CodeRigor

FREE LESSON

Multiresolution hash grid encoding

Instant-NGP's input encoding: L grids of geometric resolution, coarse ones stored densely and fine ones through a spatial hash, trilinear lookups concatenated for a tiny MLP — and the scatter-add backward pass that trains the tables.

NeRF (see gfx-volume-rendering) stores a whole scene in the weights of an 8-layer, 256-wide MLP and feeds it positions through a sine/cosine frequency encoding. It works, but every sample costs a big network evaluation and every gradient step nudges all the weights: training takes a day. Müller, Evans, Schied and Keller's Instant Neural Graphics Primitives (SIGGRAPH 2022) replaced the frequency encoding with a learned one: a stack of grids whose vertices hold trainable feature vectors, with the fine grids squeezed into fixed-size tables through a spatial hash. The MLP shrinks to two hidden layers of 64, and a NeRF trains in seconds. This encoding — or a direct descendant — sits inside most fast radiance-field systems (nerfstudio's Nerfacto, Zip-NeRF, neural SDFs, neural texture compression), so you will meet it whenever someone says "Instant-NGP backbone".

1. The pipeline

Instant-NGP: encode, then a tiny MLP
Each of the L levels contributes F features (trilinear lookup into its table). They are concatenated — optionally with other inputs such as the view direction's SH encoding — and fed to a small MLP.
Instant-NGP: encode, then a tiny MLP
Each of the L levels contributes F features (trilinear lookup into its table). They are concatenated — optionally with other inputs such as the view direction's SH encoding — and fed to a small MLP.

The paper's hyperparameters (Table 1) are worth memorising: L = 16 levels, F = 2 features per entry, table size T = 2^14 … 2^24 (2^19 for NeRF), coarsest resolution N_min = 16, finest N_max from 512 to 524288 depending on the scene's extent. Tables are initialised from U(−10⁻⁴, 10⁻⁴) — nearly zero, so the MLP starts by seeing almost nothing and the features grow where the loss asks for them.

2. Resolutions grow geometrically

Level l has resolution

N_l = ⌊N_min · b^l⌋, b = exp((ln N_max − ln N_min) / (L − 1)),

so the levels are spaced like octaves: equal ratios, not equal steps. With N_min = 16, N_max = 512, L = 16 the factor is b = 32^(1/15) = 2^(1/3), and every third level doubles. (The paper chooses b between 1.26 and 2 — finer spacing than octaves — so neighbouring levels overlap in what they can represent.)

N_l for N_min = 16, N_max = 512, L = 16
Equal ratios between levels: a straight line on a log axis. Levels whose (N+1)³ vertices fit in T = 2^19 entries are dense; the rest are hashed.
N_l for N_min = 16, N_max = 512, L = 16
Equal ratios between levels: a straight line on a log axis. Levels whose (N+1)³ vertices fit in T = 2^19 entries are dense; the rest are hashed.

Compute b in double and be careful with the floor: b^l is rounded, and a value like 63.99999999999999 floors to 63 — the finest level silently loses a vertex. This task's convention adds 1e-6 before flooring:

C++
#include <cmath>
#include <print>

int main() {
    const int L = 2, n_min = 8, n_max = 64;
    double b = std::exp((std::log(double(n_max)) - std::log(double(n_min))) / (L - 1));
    std::println("b = {:.17g}", b);
    std::println("naive floor: {}", int(std::floor(n_min * std::pow(b, 1))));
    std::println("with 1e-6:   {}", int(std::floor(n_min * std::pow(b, 1) + 1e-6)));
}

The reference code (tiny-cuda-nn's grid.h) uses a slightly different vertex convention — scale = N_min · b^l − 1, positions offset by half a cell, and ceil(scale) + 1 vertices per axis — which shifts the lattice so no two levels align exactly. The ideas are identical; the details matter only if you need bit-exact agreement with a trained Instant-NGP checkpoint.

3. One level: dense or hashed

A level of resolution N divides the unit cube into N³ cells with (N + 1)³ vertices. Each vertex owns a feature vector of F floats in the level's table. For the coarse levels the vertices fit and the table is simply a dense 3D array:

index = x + y·(N+1) + z·(N+1)² when (N + 1)³ ≤ T.

That's ≤, and vertices, not cells: N = 15 with T = 2^12 has exactly 16³ = 4096 vertices and is dense. Fine levels have far more vertices than entries (512³ ≈ 134 M vertices for 2^19 entries), so the vertex coordinates are hashed:

h(x) = (x₀·π₁ ⊕ x₁·π₂ ⊕ x₂·π₃) mod T, π₁ = 1, π₂ = 2654435761, π₃ = 805459861.

This is the hash of Teschner et al. 2003 ("Optimized Spatial Hashing for Collision Detection of Deformable Objects"), with π₁ = 1 so that neighbours along x land in neighbouring entries — better cache coherence on the GPU for the 8 corner reads. 2654435761 is the classic Knuth multiplicative-hash constant, ⌊2^32/φ⌋ rounded to a prime.

The 8 corners of cell (0,0,0) hashed into T = 8 entries
Computed with π = (1, 2654435761, 805459861): four slots used, each shared by two corners. No keys, no probing — colliding vertices simply share features, and gradients.
The 8 corners of cell (0,0,0) hashed into T = 8 entries
Computed with π = (1, 2654435761, 805459861): four slots used, each shared by two corners. No keys, no probing — colliding vertices simply share features, and gradients.

A table this small is absurd in practice, but it shows the mechanism: with T = 8, y·2654435761 contributes only its low 3 bits (…001), so moving one step in y flips the same bit as moving one step in x. Real tables have 2^14 to 2^24 entries and the collisions are spread far more evenly — the snippet below measures it.

The hash must be unsigned

The constants exceed INT_MAX / 3, so in int the products overflow — and signed overflow is undefined behaviour: UBSan reports it, constexpr evaluation rejects it, and the optimiser may assume it can't happen. The hash wants wraparound modulo 2^32, which is exactly the defined behaviour of std::uint32_t:

C++
#include <cstdint>
#include <print>

constexpr std::uint32_t spatial_hash(std::uint32_t x, std::uint32_t y, std::uint32_t z,
                                     std::uint32_t log2_t) {
    std::uint32_t h = x ^ (y * 2654435761u) ^ (z * 805459861u);   // wraps mod 2^32
    return h & ((1u << log2_t) - 1);                               // mod T, T a power of two
}

static_assert(spatial_hash(0, 1, 0, 32 - 1) == 2654435761u % (1u << 31));

int main() {
    // How evenly does one 64^3 lattice fill a table of 2^12 entries?
    const std::uint32_t log2_t = 12, T = 1u << log2_t;
    static unsigned count[1 << 12];
    for (std::uint32_t z = 0; z < 64; ++z)
        for (std::uint32_t y = 0; y < 64; ++y)
            for (std::uint32_t x = 0; x < 64; ++x) ++count[spatial_hash(x, y, z, log2_t)];
    unsigned lo = ~0u, hi = 0;
    for (std::uint32_t i = 0; i < T; ++i) {
        lo = count[i] < lo ? count[i] : lo;
        hi = count[i] > hi ? count[i] : hi;
    }
    std::println("{} vertices, {} entries: {} .. {} vertices per entry",
                 64 * 64 * 64, T, lo, hi);
    std::println("T = 8: (1,0,0) -> {}, (0,1,0) -> {}, (0,0,1) -> {}", spatial_hash(1, 0, 0, 3),
                 spatial_hash(0, 1, 0, 3), spatial_hash(0, 0, 1, 3));
}

x mod T is x & (T − 1) when T is a power of two, which it always is here.

4. Trilinear interpolation

To encode p ∈ [0,1]³ on level l: scale, s = p · N_l; the cell's lower corner is ⌊s⌋ (clamped to N − 1 so that p = 1 lands on the last vertex with weight 1); the fraction is w = s − ⌊s⌋. The 8 corners are ⌊s⌋ + d for d ∈ {0,1}³, and corner d gets weight Π_axes (d ? w : 1 − w). The level's F features are the weighted sum of the 8 corner features; the L levels are concatenated.

One level, in 2D: the cell around p and its corner weights
N = 4, p = (0.35, 0.6): s = (1.4, 2.4), cell corner (1, 2), w = (0.4, 0.4). The near corner gets (1 − 0.4)² = 0.36, the far one 0.4² = 0.16. In 3D there are 8 corners and the weights are products of three factors.
One level, in 2D: the cell around p and its corner weights
N = 4, p = (0.35, 0.6): s = (1.4, 2.4), cell corner (1, 2), w = (0.4, 0.4). The near corner gets (1 − 0.4)² = 0.36, the far one 0.4² = 0.16. In 3D there are 8 corners and the weights are products of three factors.

Interpolation makes the encoding continuous in p (and differentiable almost everywhere), which a nearest-vertex lookup would not be — the MLP downstream, and the volume renderer upstream, need that.

C++
#include <algorithm>
#include <cmath>
#include <print>

// One dense 1D level, resolution N, F = 1: the 1D analogue of the 3D lookup.
float lookup(const float* table, int N, float p) {
    float s = p * float(N);
    float x0 = std::clamp(std::floor(s), 0.0f, float(N - 1));
    float w = s - x0;
    int i = int(x0);
    return (1.0f - w) * table[i] + w * table[i + 1];
}

int main() {
    const int N = 4;
    float table[N + 1] = {0.0f, 1.0f, 4.0f, 9.0f, 16.0f};   // vertex v holds v^2
    for (float p : {0.0f, 0.3f, 0.5f, 0.99f, 1.0f})
        std::println("p = {:.2f} -> {:.3f}", p, lookup(table, N, p));
}

5. Gradients: a sparse scatter

Training pushes the loss gradient back through volume rendering and the MLP to the encoded features, ∂L/∂f (L·F numbers per sample). The encoding is linear in the table: level l's feature f is Σ_c w_c · θ[h(c)][f]. So

∂L/∂θ[i][f] = Σ over corners c with h(c) = i of w_c · ∂L/∂f_{l,f}.

The backward pass replays the forward lookup and scatter-adds w_c · upstream into each corner's entry — the transpose of the gather. Three consequences:

  • Accumulate, never assign. Many samples in a batch, and colliding corners within one sample, hit the same entry; their gradients must add up. On the GPU tiny-cuda-nn uses atomic adds (in half precision, two features at a time).
  • Sparsity. One sample touches 8·L entries out of L·T. Adam's moment estimates for untouched entries would decay pointlessly, so Instant-NGP skips the update for entries whose gradient is exactly zero.
  • No forward state needed. The derivative depends only on p (weights and indices), not on θ — the backward recomputes the cell instead of storing it.
C++
#include <algorithm>
#include <cmath>
#include <print>
#include <vector>

// Fit f(p) = sin(2πp) on [0,1] with a 1D two-level grid (N = 4 and N = 32) by SGD.
struct Level {
    int N;
    std::vector<float> theta;
};

int main() {
    std::vector<Level> levels{{4, std::vector<float>(5)}, {32, std::vector<float>(33)}};
    auto cell = [](int N, float p, int& i, float& w) {
        float s = p * float(N);
        float x0 = std::clamp(std::floor(s), 0.0f, float(N - 1));
        i = int(x0);
        w = s - x0;
    };
    auto encode_sum = [&](float p) {   // "MLP" = sum of the level features
        float y = 0;
        for (auto& lv : levels) {
            int i; float w;
            cell(lv.N, p, i, w);
            y += (1 - w) * lv.theta[i] + w * lv.theta[i + 1];
        }
        return y;
    };
    const float pi2 = 6.2831853f;
    unsigned seed = 1;
    for (int step = 0; step <= 4000; ++step) {
        seed = seed * 1664525u + 1013904223u;
        float p = float(seed >> 8) / 16777216.0f;
        float err = encode_sum(p) - std::sin(pi2 * p);       // dL/dy for L = err^2 / 2
        for (auto& lv : levels) {                             // backward: scatter
            int i; float w;
            cell(lv.N, p, i, w);
            lv.theta[i] -= 0.2f * (1 - w) * err;
            lv.theta[i + 1] -= 0.2f * w * err;
        }
        if (step % 1000 == 0) {
            float mse = 0;
            for (int k = 0; k < 100; ++k) {
                float q = (k + 0.5f) / 100.0f;
                float e = encode_sum(q) - std::sin(pi2 * q);
                mse += e * e / 100.0f;
            }
            std::println("step {:4}: mse {:.5f}", step, mse);
        }
    }
}

6. Why collisions don't ruin it

On a hashed level many vertices share an entry and nothing resolves the collision — no keys, no probing. It works anyway:

  • Gradient averaging by importance. A shared entry gets the sum of its vertices' gradients. Vertices near visible surfaces see many rays and large losses; vertices in empty space see few and small ones. The entry ends up serving the vertex that matters.
  • Multiresolution disambiguation. The coarse levels are dense, so collision-free. Two points that collide on one fine level almost surely don't collide on the others (the hash differs per resolution), and the MLP sees all levels at once.
  • The MLP learns what to trust, being trained jointly with the tables.

The table size T is the quality/memory knob: L·T·F parameters, e.g. 16 · 2^19 · 2 ≈ 16.8 M halves = 32 MiB, independent of N_max. A dense grid at 512³ with the same features would need 512 MiB at one level alone.

Entries needed per level: dense grid vs hash table (T = 2^19)
A dense level needs (N+1)³ entries; the hash caps every level at T. The memory goes where the levels are fine.
Entries needed per level: dense grid vs hash table (T = 2^19)
A dense level needs (N+1)³ entries; the hash caps every level at T. The memory goes where the levels are fine.

7. Why it trains fast

  • Capacity moved out of the MLP. The tables hold the detail, so the MLP can be 2 × 64 instead of 8 × 256: ~100× fewer FLOPs per sample, and tiny-cuda-nn evaluates it as a fully fused kernel with the weights in shared memory and activations in registers.
  • O(1), local updates. Each sample reads and updates 8·L entries; a step doesn't disturb the rest of the scene, so different regions converge in parallel instead of fighting over shared MLP weights.
  • Cache-friendly. Coarse dense levels fit in cache; π₁ = 1 keeps the x-neighbours of a corner in adjacent entries.
  • Instant-NGP adds an occupancy grid to skip empty space when marching rays — orthogonal to the encoding, but part of why the full system is "instant".

In Python (PyTorch)

Python
import tinycudann as tcnn   # NVIDIA's bindings: fused CUDA encoding + MLP
enc = tcnn.Encoding(3, {"otype": "HashGrid", "n_levels": 16, "n_features_per_level": 2,
                        "log2_hashmap_size": 19, "base_resolution": 16,
                        "per_level_scale": 1.3819})        # b = exp(ln(2048/16)/15)
feats = enc(x)                                             # x: (B, 3) in [0, 1]

# Pure PyTorch, one hashed level (what nerfstudio's fallback does):
primes = torch.tensor([1, 2654435761, 805459861], dtype=torch.int64)
corners = (x * N).floor().long()[:, None, :] + offsets      # (B, 8, 3)
idx = ((corners * primes).unbind(-1))                        # int64: no overflow
h = (idx[0] ^ idx[1] ^ idx[2]) % T
feat = (weights[..., None] * table[h]).sum(1)                # autograd does the scatter

PyTorch has no unsigned 32-bit multiply on most ops, so pure-PyTorch versions compute in int64 and reduce mod T — the same result as uint32 wraparound when T is a power of two ≤ 2^32, because only the low bits survive.

In an interview

  • "Why hash instead of an octree?" No tree traversal or structure updates during training; O(1) lookups, fully parallel, and the task-agnostic structure works for NeRF, SDFs and images alike. The price is collisions, handled implicitly.
  • "What does T trade off?" Memory and cache footprint against collision rate on fine levels; quality saturates once T covers the scene's occupied surface.
  • "How would you backprop this on a GPU?" Recompute the corners, atomic-add w · upstream into the table gradients; half-precision atomicAdd on __half2 pairs; skip optimizer updates for untouched entries.
  • "Gotchas?" (N+1)³ vertices not N³; floor not round; clamp at p = 1; uint32 hash (signed overflow is UB); accumulate gradients.

Watch & read

Free lectures, talks and readings that cover this skill from another angle. All resources →

A quick visual tour of what the hash encoding made possible: NeRFs, SDFs and gigapixel images trained in seconds.

The Instant NeRF demo: a radiance field trained from photos in seconds — the result this encoding is responsible for.