NanoVDB: parallel checksum, cooperative GridStats, coalesced IndexToGrid optimizations by swahtz · Pull Request #2250 · AcademySoftwareFoundation/openvdb

swahtz · Jul 14, 2026

Stacked on #2249. This PR is based on the #2249 branch, so until #2249 merges the diff also shows its 7 commits; GitHub will narrow it to this PR's 4 commits automatically once #2249 lands. Review the last 4 commits here.

Summary

This PR replaces the serial anti-patterns (byte-serial dependence chains, <<<1,1>>> tails, thread-per-node loops, uncoalesced thread↔address maps) in three independent NanoVDB CUDA maintenance operators with parallel associative reductions and coalesced mappings. Byte-exact where ordering permits, benchmarked A/B on an RTX PRO 6000 (sm_120; 3 rounds, 30 iters / 10 warmup) against the #2249 base.

Four commits across three independent operators — the checksum lands in two commits (a parallel rewrite plus a host-precomputed combine), GridStats and IndexToGrid one each; each touches a single independent file.

Parallel CRC32 GridChecksum (`GridChecksum.cuh`)

Replaces the serial anti-patterns in blockedCRC32:

Slicing-by-4 per-block CRC: the 256-entry base LUT is staged into shared memory and three derived slice tables built in place. Therefore, the result of this optimization is that the dependent CRC chain advances 4 bytes/step with 4 shared-memory lookups instead of 1 byte per L2 lookup.
GF(2) crc32_combine fold replacing the single-thread CRC of the whole block-CRC array (megabytes for multi-GB grids): parallel per-chunk CRCs + a zlib-construction combine (zlib's crc32_combine). This associative combine is the key optimization which makes the whole block CRC parallelizable.
Host-precomputed combine operators: the two GF(2) shift operators were originally rebuilt by binary exponentiation on a single device thread every call — a fixed ~480 µs that profiling showed dominated small/medium-grid checksums (as large as the entire streaming pass). They depend only on the chunk length, so they are now built on the host in microseconds and uploaded, and the device combine is reduced to an O(chunkCount) fold. This commit alone adds +40.7% (dragon) / +15.3% (emu) / +9.8% (crawler) / +1.9% (wdas_cloud) — a fixed cost, so it helps smaller grids most.

Checksum values are bit-identical (all transforms are value-preserving); crc32TailOld path untouched.

A/B (median of 3 rounds vs the #2249 base):

dataset	baseline	this PR	saved	speedup
dragon	7.66 ms	0.70 ms	6.96 ms	10.9×
emu	34.74 ms	2.35 ms	32.39 ms	14.8×
crawler	57.76 ms	3.56 ms	54.20 ms	16.2×
wdas_cloud	147.19 ms	8.48 ms	138.71 ms	17.4×

Cooperative GridStats (`GridStats.cuh`)

GridStats builds an accumulation of statistics over the tree aggregating stats bottom-up. Before this optimization PR, each level of the tree (leaf, lower inner, upper inner, root) was run as one thread per node with that thread doing the reduction serially. This PR recasts this work as a parallel-reduction… we assign a group of threads to each node and reduce cooperatively in shared memory.

processLeaf: one warp per leaf (lanes stride the 512 slots with 32 lanes/16 voxels-per-lane with a 5-step tree merge).
processInternal: one 128-thread block per internal node with a shared-memory tree reduction (the old thread-per-node loop gave each thread up to 32,768 serial child visits). Now 128 threads each doing 256 child visits then a 7-step merge.

min/max are order-independent → stats_minmax is bit-identical. stats_all (avg/stddev) uses a Welford-correct parallel merge but is not bit-identical (float reduction order) — it stays a bench-only measurement, never a byte-exact claim.

A/B — stats_minmax (byte-exact; median of 3 rounds):

dataset	baseline	this PR	saved	speedup
dragon	2.52 ms	0.32 ms	2.20 ms	7.9×
emu	4.36 ms	0.76 ms	3.60 ms	5.7×
crawler	5.58 ms	1.09 ms	4.48 ms	5.1×
wdas_cloud	19.00 ms	3.13 ms	15.87 ms	6.1×

A/B — stats_all (bench-only, not bit-identical):

dataset	baseline	this PR	saved	speedup
dragon	4.39 ms	1.23 ms	3.16 ms	3.6×
emu	8.68 ms	4.20 ms	4.48 ms	2.1×
crawler	11.63 ms	8.12 ms	3.51 ms	1.4×
wdas_cloud	30.84 ms	20.04 ms	10.80 ms	1.5×

Coalesced IndexToGrid remap (`IndexToGrid.cuh`)

Consecutive threads process consecutive table entries / leaf values (the old mapping gave each thread a strided run across the warp), and the 4 KB upper-node value/child masks are copied cooperatively instead of member-wise by thread 0. Preserves the scoped zero-init from #2249, so output is bit-identical.

A/B (byte-exact; median of 3 rounds) — on top of the #2249 fixes:

dataset	baseline	this PR	saved	speedup
dragon	1.28 ms	0.99 ms	0.29 ms	1.3×
emu	3.93 ms	3.06 ms	0.86 ms	1.3×
crawler	5.87 ms	4.85 ms	1.02 ms	1.2×
wdas_cloud	14.55 ms	12.42 ms	2.14 ms	1.2×

Validation

Bottom line (real time): on wdas_cloud the three byte-exact maintenance ops together drop from ~180.7 ms → ~24.0 ms (checksum + stats_minmax + indextogrid), ~157 ms saved per invocation; on dragon, ~11.5 ms → ~2.0 ms.
Byte-exact for checksum, stats_minmax, indextogrid across dragon/emu/crawler/wdas_cloud.
16 A/B wins, 0 regressions (19 ops compared); transient + retained memory unchanged; floodfill and the rest NEUTRAL.
compute-sanitizer: memcheck clean; initcheck identical to the base build (pre-existing host-side readback reports, none in the new kernels).

Notes

Profiling after the combine fix shows the checksum is now bound by the per-block streaming kernel at ~31% of peak DRAM (compute is essentially free). That kernel is limited by the checksum's definition-mandated 4 KB block layout (an uncoalesced, one-block-per-thread access pattern), not by bandwidth, so the remaining headroom would require a cooperative-staging rewrite with an uncertain payoff — deliberately out of scope here.

🤖 Generated with Claude Code, revised by a human

Five correctness/determinism fixes to the NanoVDB CUDA tools: * GridChecksum: evalChecksum copied sizeof(GridData)+sizeof(TreeData) (736 B) from a device buffer holding a single uint32 CRC into the 4-byte head/tail of a stack Checksum -- an out-of-bounds device read plus host stack corruption (compute-sanitizer visible; a hard CUDA "invalid argument" on some configs). Copy sizeof(uint32_t) at all four sites, matching validateChecksum. * Stream propagation: the NN_FACE_EDGE_VERTEX leaf-dilation kernel and SignedFloodFill's root pass ran on the default stream instead of the caller's stream, silently corrupting output for callers using a non-blocking stream. Launch and synchronize on the operation's stream. * Bounded grid-name copy: the grid-name copy read a fixed MaxNameSize bytes from a possibly shorter std::string (whose .data() is never null, making the memset branch dead), over-reading up to 255 bytes of host heap into the grid and any written .nvdb file. Zero the field, then copy only min(name.size()+1, MaxNameSize) bytes. Applied to PointsToGrid, DistributedPointsToGrid, MeshToGrid. * Deterministic output initialization: PointsToGrid and IndexToGrid left alignment padding and stats fields carrying recycled pool bytes, making output nondeterministic and leaking host heap into files. Zero the whole grid buffer once after allocation. For PointsToGrid this makes the three dense background-fill passes (upper/lower value tables, and the inactive-leaf-values pass for all but the Point build type) redundant, so they are dropped. * DeviceBuffer stream-ordered lifetime: the destructor, move-assignment and clear() freed device allocations on the default stream regardless of the stream the buffer was last used on -- a latent use-after-free for non-blocking-stream callers, masked today by legacy-stream semantics. Track the allocation stream and free on it, matching the stream-ordered free TempPool now performs. Adds gtest coverage in unittest/TestNanoVDB.cu: cross-stream dilation determinism, grid-name round-trip (empty/short/max-length), PointsToGrid output determinism, and DeviceBuffer non-blocking-stream lifetime. No performance claims. Signed-off-by: Jonathan Swartz <jonathan@jswartz.info> Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…eeds The Tranche-1 determinism fix zero-initialized the *entire* IndexToGrid output buffer so that bytes the build kernels never write - stats fields absent from the source grid, and struct alignment padding - are deterministic instead of carrying recycled pool bytes (which otherwise leak into written .nvdb files). Measured against master this cost 10-22% on indextogrid across dragon/emu/ crawler/wdas_cloud: the leaf array dominates the buffer and was memset to zero and then immediately overwritten value-by-value - a redundant multi-GB pass. Scope the initialization to the bytes that are genuinely left unwritten: * getBuffer() memsets only the non-leaf region [0, node[0]) - grid, tree, root, root tiles and the internal nodes - a small fraction of the buffer. * processLeafsKernel zeroes each leaf's own gap [&mMinimum, mValues): the stats fields (skipped when the source has no stats) and any padding before the 32-aligned mValues array. Every mValues[i] is written anyway, so this reproduces the former full zero-init byte-for-byte. Output is bit-identical to the full zero-init - all indextogrid goldens pass on dragon/emu/crawler/wdas_cloud, compute-sanitizer memcheck is clean and initcheck shows no new reports - and the regression vs master is gone (NEUTRAL on all four datasets, RTX PRO 6000 / sm_120). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…ce in test Addresses PR review: - The grid-name copy could leave GridData::mGridName without a null terminator when the name length >= MaxNameSize: min(size+1, MaxNameSize) == MaxNameSize overwrote the entire zeroed field, and gridName() (a C-string accessor) would then read past it. Copy at most MaxNameSize-1 bytes and rely on the preceding memset for termination, truncating over-long names instead. Fixed in PointsToGrid, MeshToGrid and DistributedPointsToGrid. - GridName_CudaPointsToGrid gains an over-long (>= MaxNameSize) case asserting the result is truncated to MaxNameSize-1 chars and stays null-terminated. - NonBlockingStreamDilate_ValueOnIndex queried device 0 for the clock rate; query the current device via cudaGetDevice() instead. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…ream The Tranche-1 stream-correctness fix added an unconditional cudaStreamSynchronize(stream) before processRoot's synchronous default-stream cudaMemcpy of the tree. That sync is only needed when the caller's node passes ran on a non-default stream; on the default stream (0) the blocking copy already orders after prior stream-0 work, so the extra sync there was pure overhead - a measured ~4-11% regression on floodfill/dragon vs master. Guard it with `if (stream != 0)`. Correctness for non-blocking-stream callers is unchanged. floodfill is now NEUTRAL vs master on dragon/emu/crawler (0 regressions), and byte-exact validation is unaffected. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Addresses PR review: the Tranche-1 stream-ordered-free change stored a single mStream, overwritten by every DeviceBuffer::deviceUpload. A DeviceBuffer can hold an allocation on each device (mGpuData is per-device), and the multi-GPU example (ex_make_mgpu_nanovdb) uploads one handle to every device with that device's own stream. The destructor/move-assign then freed *every* device's allocation on the last device's stream - a cross-device cudaFreeAsync(ptr@devA, stream@devB), which is invalid (the free stream must belong to the allocation's device). Pre-Tranche-1 code freed on stream 0 uniformly, so this was a regression for the multi-GPU path. Track the allocation stream per device in a cudaStream_t array parallel to mGpuData, set in init()/deviceUpload() at mStreams[device], and free each mGpuData[i] on its own mStreams[i] in clear()/move-assignment. The single-GPU path is unchanged (one device, one stream); the multi-GPU path now frees each device's allocation on its own device-bound stream. Single-GPU verified (DeviceBuffer lifetime gtest + byte-exact validation); the multi-GPU free path is correct by construction but not runtime-tested here (single-GPU environment). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Copilot

Pull request overview

This PR modernizes several NanoVDB CUDA “maintenance” operators by removing serial GPU work patterns and improving stream-correctness and determinism, while adding/expanding CUDA regression tests to lock in behavior (notably byte-determinism and non-default stream safety).

Changes:

Parallelizes GridChecksum and GridStats CUDA kernels (associative reductions / cooperative processing) to remove long serial dependence chains.
Improves determinism and hygiene by zero-initializing otherwise-unwritten regions (and tightening grid-name copying to avoid host over-reads).
Fixes stream propagation/synchronization issues in CUDA tools and adds regression tests for non-blocking-stream correctness and deterministic output.

Reviewed changes

Copilot reviewed 10 out of 10 changed files in this pull request and generated 2 comments.

Show a summary per file

File	Description
nanovdb/nanovdb/unittest/TestNanoVDB.cu	Adds CUDA regression tests for non-blocking stream correctness, grid-name round-trip, deterministic output, and DeviceBuffer stream-ordered lifetime.
nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh	Ensures root processing orders correctly with a caller-provided stream and launches device work on that stream.
nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh	Adds buffer-wide zero-init for determinism, removes redundant background fill kernels, and fixes grid-name copying to avoid over-read/truncation issues.
nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh	Fixes grid-name copying by zeroing the name field and copying only the actual string length.
nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh	Improves coalescing/cooperation for node/table/value writes and adds targeted deterministic initialization for otherwise-unwritten bytes.
nanovdb/nanovdb/tools/cuda/GridStats.cuh	Reworks stats computation to be warp/block cooperative (warp-per-leaf, block-per-internal-node) and reduces serial traversal.
nanovdb/nanovdb/tools/cuda/GridChecksum.cuh	Introduces slicing-by-4 per-block CRC and a GF(2) combine fold to avoid serial CRC passes; fixes host copy sizes for head/tail CRC.
nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh	Fixes grid-name copy to avoid host over-read by zeroing and bounded copy.
nanovdb/nanovdb/tools/cuda/DilateGrid.cuh	Ensures NN_FACE_EDGE_VERTEX dilation launch uses the caller’s stream.
nanovdb/nanovdb/cuda/DeviceBuffer.h	Tracks per-device allocation streams so frees are ordered on the owning stream (stream-safe lifetime).

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

…-wait (test) Addresses PR review: the test launches streamBusyWaitKernel on the default stream to occupy it during the candidate dilation, but only synchronized the non-blocking stream afterward. The busy-wait could remain queued on stream 0 and leak into subsequent tests, causing timing-dependent flakes. Drain it with a cudaDeviceSynchronize() after the candidate run. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Copilot

Pull request overview

Copilot reviewed 10 out of 10 changed files in this pull request and generated 4 comments.

Addresses PR review: mStreams[device] was set only when the device allocation was first created. If the buffer is later re-uploaded (or downloaded) on a different stream, the destructor/move-assign/clear would still free on the old stream, which does not necessarily order after the most recent work on the new stream (notably for cudaStreamNonBlocking). Update mStreams[device] to the current stream on every host<->device copy in deviceUpload/deviceDownload, not just on first allocation. Guarded on mStreams (null for externally-managed buffers, which are never freed here). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Copilot

Pull request overview

Copilot reviewed 10 out of 10 changed files in this pull request and generated 4 comments.

…am comments Addresses PR review: - PointsToGrid/MeshToGrid/DistributedPointsToGrid use std::min in the grid-name copy but did not include <algorithm> (it compiled only via transitive includes). Add the include so each header is self-sufficient. - DeviceBuffer comments still described mStreams as the "allocation stream", but it is now updated on every deviceUpload/deviceDownload; reword them to "last-used stream" to match the actual behavior. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…2) combine) Port of Agent A's exp/14. Replaces the two serial anti-patterns in blockedCRC32 with parallel equivalents: - Per-block CRC via slicing-by-4: the 256-entry base LUT is staged into shared memory and three derived slice tables are built in place, so the dependent CRC chain advances four bytes per step with four shared-memory lookups instead of one byte per L2 lookup. (Not __constant__: divergent indices serialize there.) - The final fold (a single thread CRCing the whole block-CRC array - megabytes for multi-GB grids) becomes parallel per-chunk CRCs plus a GF(2) crc32_combine (zlib construction: one shift operator per fixed chunk length, applied per fold). Small inputs keep the single-thread path. Checksum VALUES are bit-identical (slicing-by-4 and the GF(2) combine are both value-preserving); all checksum_full goldens pass on dragon/emu/crawler/ wdas_cloud. The pre-v32.6.0 crc32TailOld path is untouched. A/B vs the Tranche-1 base (RTX PRO 6000, 3 rounds): dragon +84.7%, emu +92.0%, crawler +93.2%, wdas_cloud +94.1% (~6.5x / 12.5x / 15x / 17x); transient memory unchanged. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…per-node) Port of Agent A's exp/13. Replaces the thread-per-node serial reductions: - processLeaf: one warp per leaf - lanes stride the 512 voxel slots (mask-gated) and merge partial statistics through shared memory; lane 0 keeps the bbox update and the final store. - processInternal: one 128-thread block per internal node striding the 4,096/32,768-entry child table with a shared-memory tree reduction of Stats + bbox. The former kernel gave each thread up to 32,768 serial child visits and ran an entire lower level of a 181M-voxel grid on ~9 warps of a 188-SM GPU. min/max are order-independent, so stats_minmax output is bit-identical - all stats_minmax goldens pass on dragon/emu/crawler/wdas_cloud. avg/stddev (stats_all) use a Welford-correct parallel merge but are NOT bit-identical to the serial result (floating-point reduction order differs); stats_all stays a bench-only measurement, never a byte-exact claim. A/B vs the Tranche-1 base (RTX PRO 6000, 3 rounds): stats_minmax dragon +87.3% emu +82.5% crawler +80.3% wdas +83.5% (5-8x) stats_all dragon +71.9% emu +51.7% crawler +29.8% wdas +35.1% transient memory unchanged. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Port of Agent A's exp/20. Consecutive threads now process consecutive table entries and leaf values - the former mapping gave each thread a consecutive RUN (a 32-entry stride across the warp for internal nodes, 8 for leaves) on every iteration. The 4 KB upper-node value/child masks are copied cooperatively across the block instead of member-wise by thread 0 while the rest idled. Combines with the Tranche-1 scoped zero-init already on this file: the leaf kernel's stats/padding-gap zeroing is preserved inside the thread-0 block, and the coalesced value loop still writes every mValues[i], so output stays bit-identical - all indextogrid goldens pass on dragon/emu/crawler/wdas_cloud. A/B vs the Tranche-1 base (RTX PRO 6000, 3 rounds): dragon +23%, emu +21%, crawler +18%, wdas_cloud +14%; transient memory unchanged. The remap gain is on top of Tranche 1's scoped init, so absolute indextogrid is now faster than both master and A's exp/20. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…-node stats slot Addresses PR review on the cooperative GridStats: - processLeaf is now launched only when nodeCount[0] > 0, matching the guards on the internal-node launches (a leaf-less grid would otherwise be a <<<0,...>>> launch). - The average path now gives each node its own d_stats slot instead of reusing one of its children's. A fully-tiled internal node (active value tiles, no children) never claimed a child slot, so the old code wrote d_stats[-1] (out of bounds) and propagated an invalid slot to its parent. d_stats is now sized for all nodes (leaves, then lower, then upper) and each node writes slot leafCount + levelOffset + nodeID; the sSlot/atomicMax machinery is removed. stats_minmax is unaffected (it never touches d_stats) and stays byte-exact; stats_all output is unchanged for grids without childless internal nodes and correct for those that have them. compute-sanitizer memcheck clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

kmuseth

Impressive optimizations, though it's hard to follow the details of the tricks, but they appear to be primarily related to better use of shared memory and asynchronous processing with streams. The use of raw intrinsics produces cryptic-looking code, so please add comments. The same is true for some of the more intricate algorithmic optimization tricks. Happy to see lots of performance and unit tests.

kmuseth · Jul 16, 2026

-    if (d_leaf.updateBBox()) {// updates active bounding box (also updates data->mFlags) and return true if non-empty
+    bool nonEmpty = false;
+    if (lane == 0) nonEmpty = d_leaf.updateBBox();// updates active bounding box (also updates data->mFlags)
+    nonEmpty = __shfl_sync(0xffffffffu, nonEmpty, 0);


please add a comment - looks cryptical

True, and the subsequent code was also poorly commented. I added commentary here to explain plus in the logic below and correspondingly in processInternal

kmuseth · Jul 16, 2026

-                *reinterpret_cast<uint32_t*>(&d_leaf.mMinimum) = tid;
-            } else {
-                stats.setStats(d_leaf);
+            const auto &mask = d_leaf.valueMask();


it looks like a cool optimization but I'd like to understand the underlying principle. Just a comment

Sure, I've added some more colour to what is happening in this optimization to the PR summary and put some more explanation in the comments to what is going on at each stage. Instead of a thread processing a whole leaf node's stats, we map a 32-thread warp to each leaf node. This cooperative reduction lets us increase occupancy and coalesce memory reads done by the warp (instead of each thread in a warp reading stats from different nodes). The core inspiration was chapter 10 from the Programming Massively Parallel Processors book.

kmuseth

Oh, one mandatory change: please update the pendingchanges/nanovdb.txt file to reflect these improvements.

crc32CombineKernel rebuilt its two GF(2) shift operators by binary exponentiation on a single device thread every call - a fixed ~480 us that dominated the checksum for small and medium grids (profiling showed it as large as the entire tail-streaming pass). The operators depend only on the chunk length, so they are now built on the host in microseconds and uploaded, and the device kernel is reduced to the O(chunkCount) fold. Bit-identical: all checksum goldens pass on dragon/emu/crawler/wdas_cloud. A/B vs the parallel-checksum base (RTX PRO 6000, 3 rounds): dragon +40.7% (1.69x) emu +15.3% crawler +9.8% wdas_cloud +1.9% The gain shrinks with grid size because it removes a fixed cost; transient memory unchanged. Signed-off-by: Jonathan Swartz <jonathan@jswartz.info> Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

swahtz and others added 7 commits July 14, 2026 02:42

Fix up comments

42afea5

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

comment fix

fb1c692

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

swahtz requested a review from kmuseth as a code owner July 14, 2026 05:25

swahtz changed the title ~~NanoVDB CUDA: Tranche 2 — parallel checksum, cooperative GridStats, coalesced IndexToGrid~~ NanoVDB: parallel checksum, cooperative GridStats, coalesced IndexToGrid Jul 14, 2026

swahtz added the nanovdb label Jul 14, 2026

swahtz changed the title ~~NanoVDB: parallel checksum, cooperative GridStats, coalesced IndexToGrid~~ NanoVDB CUDA: parallel checksum, cooperative GridStats, coalesced IndexToGrid Jul 14, 2026

swahtz requested a review from Copilot July 14, 2026 05:35

Copilot started reviewing on behalf of swahtz July 14, 2026 05:35 View session

Copilot AI reviewed Jul 14, 2026

View reviewed changes

Comment thread nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh Outdated

Comment thread nanovdb/nanovdb/unittest/TestNanoVDB.cu

swahtz force-pushed the tranche-2-maintenance-monoids branch from 90ef727 to f7220d6 Compare July 14, 2026 22:22

swahtz requested a review from Copilot July 14, 2026 22:24

swahtz changed the title ~~NanoVDB CUDA: parallel checksum, cooperative GridStats, coalesced IndexToGrid~~ NanoVDB: parallel checksum, cooperative GridStats, coalesced IndexToGrid Jul 14, 2026

Copilot started reviewing on behalf of swahtz July 14, 2026 22:25 View session

Copilot AI reviewed Jul 14, 2026

View reviewed changes

Comment thread nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh Outdated

Comment thread nanovdb/nanovdb/tools/cuda/GridStats.cuh Outdated

Comment thread nanovdb/nanovdb/tools/cuda/GridStats.cuh Outdated

Comment thread nanovdb/nanovdb/cuda/DeviceBuffer.h

swahtz force-pushed the tranche-2-maintenance-monoids branch from f7220d6 to 9f85d9f Compare July 14, 2026 22:57

swahtz requested a review from Copilot July 14, 2026 23:11

Copilot started reviewing on behalf of swahtz July 14, 2026 23:12 View session

Copilot AI reviewed Jul 14, 2026

View reviewed changes

Comment thread nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh

Comment thread nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh

Comment thread nanovdb/nanovdb/cuda/DeviceBuffer.h Outdated

Comment thread nanovdb/nanovdb/cuda/DeviceBuffer.h Outdated

swahtz and others added 5 commits July 15, 2026 00:02

swahtz force-pushed the tranche-2-maintenance-monoids branch from 9f85d9f to 72a4bfe Compare July 15, 2026 01:10

kmuseth approved these changes Jul 16, 2026

View reviewed changes

kmuseth requested changes Jul 16, 2026

View reviewed changes

swahtz changed the title ~~NanoVDB: parallel checksum, cooperative GridStats, coalesced IndexToGrid~~ NanoVDB: parallel checksum, cooperative GridStats, coalesced IndexToGrid optimizations Jul 16, 2026

swahtz requested review from Idclip, apradhana, danrbailey, jmlait and richhones as code owners July 19, 2026 01:57

swahtz added 4 commits July 19, 2026 03:53

Improve comments for processInternal/Leaf

06ad21b

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Comment cleanup and swap some constants for semantically defined values

ad25134

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Update pendingchanges

65c9d48

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Merge branch 'master' into tranche-2-maintenance-monoids

b0b99d6

Search code, repositories, users, issues, pull requests...

Uh oh!

Conversation

swahtz commented Jul 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Parallel CRC32 GridChecksum (GridChecksum.cuh)

Cooperative GridStats (GridStats.cuh)

Coalesced IndexToGrid remap (IndexToGrid.cuh)

Validation

Notes

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

kmuseth left a comment

Choose a reason for hiding this comment

Uh oh!

kmuseth Jul 16, 2026

Choose a reason for hiding this comment

Uh oh!

swahtz Jul 19, 2026

Choose a reason for hiding this comment

Uh oh!

kmuseth Jul 16, 2026

Choose a reason for hiding this comment

Uh oh!

swahtz Jul 19, 2026

Choose a reason for hiding this comment

Uh oh!

kmuseth left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

swahtz commented Jul 14, 2026 •

edited

Loading

Parallel CRC32 GridChecksum (`GridChecksum.cuh`)

Cooperative GridStats (`GridStats.cuh`)

Coalesced IndexToGrid remap (`IndexToGrid.cuh`)

kmuseth left a comment •

edited

Loading