Environment
NanoVDB master (bb597a21), CUDA 12.x nvcc, any build with -DNANOVDB_USE_CUDA=ON.
Problem
Building anything that instantiates the GPU stats kernels emits 12 copies of nvcc #20054-D:
nanovdb/tools/cuda/GridStats.cuh(68): warning #20054-D: dynamic initialization is not supported for a function-scope static __shared__ variable within a __device__/__global__ function
nanovdb/tools/cuda/GridStats.cuh(135): warning #20054-D: ...
nanovdb/tools/cuda/GridStats.cuh(136): warning #20054-D: ...
The warnings were introduced by #2250, which rewrote the stats pass into cooperative kernels using shared-memory reductions:
// processLeaf
__shared__ StatsT sStats[WarpsPerBlock * 32]; // line 68
// processInternal
__shared__ StatsT sStats[Threads]; // line 135
__shared__ CoordBBox sBBox[Threads]; // line 136
Stats/Extrema (and CoordBBox) have user-provided default constructors, so these arrays require dynamic initialization — which CUDA does not support for __shared__ variables. nvcc skips the constructors and warns. The 12 instances are the per-BuildT/StatsT kernel instantiations.
Correctness is unaffected: both kernels fully assign every shared slot before reading it (sWarp[lane] = stats / sStats[tID] = stats ahead of the reduction loops), and the early-return paths exit at whole-warp/whole-block granularity, so no partially-written slice enters a reduction. The skipped initialization never matters.
The cost is noise: GridStats.cuh is a header, so the warnings propagate into every downstream TU that instantiates these kernels — including end-user builds.
Suggested fix
Use cub::Uninitialized<T> (in <cub/util_type.cuh>, same pattern CUB uses for its own TempStorage; CUB is already in NanoVDB's include graph). Its backing store is a trivially-constructible word array, so no dynamic init is required, and Alias() reinterprets it as T:
// processLeaf
__shared__ cub::Uninitialized<StatsT[WarpsPerBlock * 32]> sStatsRaw;
StatsT (&sStats)[WarpsPerBlock * 32] = sStatsRaw.Alias();
// processInternal
__shared__ cub::Uninitialized<StatsT[Threads]> sStatsRaw;
__shared__ cub::Uninitialized<CoordBBox[Threads]> sBBoxRaw;
StatsT (&sStats)[Threads] = sStatsRaw.Alias();
CoordBBox (&sBBox)[Threads] = sBBoxRaw.Alias();
All downstream code (the assignments and reduction loops) is unchanged. A dependency-free alternative is an alignas(StatsT) unsigned char array reinterpreted as StatsT*, but cub::Uninitialized self-documents the intent and picks the widest legal access word for the backing storage.
Environment
NanoVDB master (
bb597a21), CUDA 12.x nvcc, any build with-DNANOVDB_USE_CUDA=ON.Problem
Building anything that instantiates the GPU stats kernels emits 12 copies of nvcc
#20054-D:The warnings were introduced by #2250, which rewrote the stats pass into cooperative kernels using shared-memory reductions:
Stats/Extrema(andCoordBBox) have user-provided default constructors, so these arrays require dynamic initialization — which CUDA does not support for__shared__variables. nvcc skips the constructors and warns. The 12 instances are the per-BuildT/StatsTkernel instantiations.Correctness is unaffected: both kernels fully assign every shared slot before reading it (
sWarp[lane] = stats/sStats[tID] = statsahead of the reduction loops), and the early-return paths exit at whole-warp/whole-block granularity, so no partially-written slice enters a reduction. The skipped initialization never matters.The cost is noise:
GridStats.cuhis a header, so the warnings propagate into every downstream TU that instantiates these kernels — including end-user builds.Suggested fix
Use
cub::Uninitialized<T>(in<cub/util_type.cuh>, same pattern CUB uses for its ownTempStorage; CUB is already in NanoVDB's include graph). Its backing store is a trivially-constructible word array, so no dynamic init is required, andAlias()reinterprets it asT:All downstream code (the assignments and reduction loops) is unchanged. A dependency-free alternative is an
alignas(StatsT) unsigned chararray reinterpreted asStatsT*, butcub::Uninitializedself-documents the intent and picks the widest legal access word for the backing storage.