AcademySoftwareFoundation · swahtz · Aug 18, 2026 · Aug 12, 2026 · Aug 12, 2026 · Aug 12, 2026
@@ -31,6 +31,7 @@ option(NANOVDB_BUILD_PYTHON_MODULE "Build the nanovdb Python module" OFF)
 option(NANOVDB_USE_INTRINSICS "Build with hardware intrinsics support" OFF)
 option(NANOVDB_USE_CUDA "Build with CUDA support" OFF)
 option(NANOVDB_CUDA_KEEP_PTX "Keep CUDA PTX" OFF)
+option(NANOVDB_CUDA_WERROR "Treat all NVCC diagnostics as errors" ON)
 option(NANOVDB_USE_OLD_ACCESSOR "Use the old (buggy) ReadAccessor instead of the new fixed one" OFF)
 
 option(NANOVDB_USE_OPENVDB "Build with OpenVDB support" OFF)
@@ -96,6 +97,21 @@ if(NANOVDB_USE_CUDA)
 
   set(CMAKE_CUDA_FLAGS "${NANOVDB_CUDA_EXTENDED_LAMBDA} -use_fast_math -lineinfo ${CMAKE_CUDA_FLAGS}")
 
+  if(NANOVDB_CUDA_WERROR OR OPENVDB_CXX_STRICT)
+    # Device-side analog of the host -Werror configured in OpenVDBCXX.cmake:
+    # every nvcc front-end diagnostic becomes an error. NANOVDB_CUDA_WERROR
+    # enables it independently of the host-side strict flags.
+    set(CMAKE_CUDA_FLAGS "--Werror=all-warnings ${CMAKE_CUDA_FLAGS}")
+  endif()
+
+  if(WIN32)
+    # NVCC analog of the /wd4251 and /wd4275 suppressions in OpenVDBCXX.cmake:
+    # it's not possible to use STL types in DLL interfaces in a portable and
+    # reliable way, so the dll-interface diagnostics (#1394 and #1388) fire on
+    # every include of the OpenVDB core headers from a CUDA source.
+    set(CMAKE_CUDA_FLAGS "--diag-suppress=1388,1394 ${CMAKE_CUDA_FLAGS}")
+  endif()
+
   # workaround for win32 bug when nvcc "--keep" is used.
   if(WIN32)
     if(NANOVDB_CUDA_KEEP_PTX)

@@ -3,7 +3,9 @@
 
 #pragma once
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 #include <fstream>

@@ -1,7 +1,9 @@
 // Copyright Contributors to the OpenVDB Project
 // SPDX-License-Identifier: Apache-2.0
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 

@@ -3,7 +3,9 @@
 
 #if defined(NANOVDB_USE_OPENVDB)
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 

@@ -3,7 +3,9 @@
 
 #pragma once
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 #include <fstream>

@@ -1,7 +1,9 @@
 // Copyright Contributors to the OpenVDB Project
 // SPDX-License-Identifier: Apache-2.0
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 

@@ -3,7 +3,9 @@
 
 #if defined(NANOVDB_USE_OPENVDB)
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 

@@ -3,7 +3,9 @@
 
 #pragma once
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 #include <fstream>

@@ -1,7 +1,9 @@
 // Copyright Contributors to the OpenVDB Project
 // SPDX-License-Identifier: Apache-2.0
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 

@@ -3,7 +3,9 @@
 
 #pragma once
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 #include <fstream>

@@ -1,7 +1,9 @@
 // Copyright Contributors to the OpenVDB Project
 // SPDX-License-Identifier: Apache-2.0
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 

@@ -3,7 +3,9 @@
 
 #if defined(NANOVDB_USE_OPENVDB)
 
+#ifndef _USE_MATH_DEFINES
 #define _USE_MATH_DEFINES
+#endif
 #include <cmath>
 #include <chrono>
 

@@ -636,7 +636,6 @@ void DistributedPointsToGrid<BuildT>::countNodes(const PtrT coords, size_t coord
         deviceStripeCounts[deviceId] = deviceStripeCount;
 
         if (deviceStripeCount) {
-            nanovdb::Coord* deviceCoords = coords + deviceStripeOffset;
             uint64_t* deviceInputKeys = mKeys + deviceStripeOffset;
             uint32_t* deviceInputIndices = mIndices + deviceStripeOffset;
             uint64_t* deviceOutputKeys = mData->d_keys + deviceStripeOffset;

@@ -115,7 +115,7 @@ inline unique_ptr<uint32_t> createCrc32Lut(size_t extra = 0, cudaStream_t stream
 /// step hit shared memory and the dependent update chain advances four bytes
 /// per step instead of one. Bit-identical to the byte-serial crc32().
 /// The final block absorbs any remainder of @c totalSize.
-__global__ inline void crc32SlicedKernel(const void *d_data, uint32_t* d_blockCRC, uint64_t blockCount, uint32_t log2BlockSize, uint64_t totalSize, const uint32_t *d_lut)
+static __global__ void crc32SlicedKernel(const void *d_data, uint32_t* d_blockCRC, uint64_t blockCount, uint32_t log2BlockSize, uint64_t totalSize, const uint32_t *d_lut)
 {
     __shared__ uint32_t sLut[4][256];
     for (uint32_t i = threadIdx.x; i < 256; i += blockDim.x) sLut[0][i] = d_lut[i];
@@ -182,7 +182,7 @@ __host__ __device__ inline void crc32BuildShiftOp(uint32_t *dst, uint64_t bits)
 /// possibly shorter, chunk - are precomputed on the host, so this is just an
 /// O(chunkCount) fold. Bit-identical to a serial crc32 over the concatenated
 /// stream. (@c d_accLast equals @c d_acc when the last chunk is full.)
-__global__ inline void crc32CombineKernel(const uint32_t *d_chunkCRC, uint64_t chunkCount, const uint32_t *d_acc, const uint32_t *d_accLast, uint32_t *d_crc)
+static __global__ void crc32CombineKernel(const uint32_t *d_chunkCRC, uint64_t chunkCount, const uint32_t *d_acc, const uint32_t *d_accLast, uint32_t *d_crc)
 {
     uint32_t crc = d_chunkCRC[0];
     for (uint64_t i = 1; i < chunkCount; ++i) {

@@ -18,6 +18,8 @@
 #include <nanovdb/NanoVDB.h>
 #include <nanovdb/tools/GridStats.h>
 
+#include <cub/util_type.cuh>// for cub::Uninitialized
+
 namespace nanovdb {
 
 namespace tools::cuda {
@@ -65,7 +67,11 @@ template<typename BuildT, typename StatsT>
 __global__ void processLeaf(NodeManager<BuildT> *d_nodeMgr, StatsT *d_stats)
 {
     constexpr uint32_t WarpsPerBlock = 4;
-    __shared__ StatsT sStats[WarpsPerBlock * 32];
+    // CUDA does not run constructors for __shared__ variables, so wrap the scratch
+    // array in cub::Uninitialized (raw, constructor-free storage). Safe here because
+    // every slot is fully assigned before it is read.
+    __shared__ cub::Uninitialized<StatsT[WarpsPerBlock * 32]> sStatsStorage;
+    StatsT (&sStats)[WarpsPerBlock * 32] = sStatsStorage.Alias();
 
     const uint32_t warpID = threadIdx.x >> 5, lane = threadIdx.x & 31u;
     const uint32_t tid = blockIdx.x * WarpsPerBlock + warpID;// leaf index
@@ -132,8 +138,12 @@ __global__ void processInternal(NodeManager<BuildT> *d_nodeMgr, StatsT *d_stats)
     using ChildT = typename NanoNode<BuildT,LEVEL-1>::type;
     using NodeT = typename NanoNode<BuildT,LEVEL>::type;
     constexpr uint32_t Threads = 128;
-    __shared__ StatsT sStats[Threads];
-    __shared__ CoordBBox sBBox[Threads];
+    // Constructor-free shared storage, as in processLeaf: every slot is assigned
+    // before it is read.
+    __shared__ cub::Uninitialized<StatsT[Threads]> sStatsStorage;
+    __shared__ cub::Uninitialized<CoordBBox[Threads]> sBBoxStorage;
+    StatsT (&sStats)[Threads] = sStatsStorage.Alias();
+    CoordBBox (&sBBox)[Threads] = sBBoxStorage.Alias();
 
     const uint32_t nodeID = blockIdx.x;
     const uint32_t tID = threadIdx.x;

@@ -79,6 +79,18 @@ if(NANOVDB_USE_CUDA)
   target_link_libraries(nanovdb_test_cuda_buffer PRIVATE nanovdb GTest::GTest GTest::Main)
   set_target_properties(nanovdb_test_cuda_buffer PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
   add_test(nanovdb_cuda_buffer_unit_test nanovdb_test_cuda_buffer)
+
+  if(WIN32)
+    # gtest's TEST macro declares each test's static test_info_ member with an
+    # unused-attribute on GCC/Clang but not under MSVC, so for tests defined
+    # inside an anonymous namespace the CUDA front end can prove it
+    # unreferenced and reports #177, which NANOVDB_CUDA_WERROR promotes to an
+    # error. Suppress #177 for the CUDA test sources that use that style.
+    foreach(_cuda_gtest nanovdb_test_cuda_buffer nanovdb_test_cuda_memory_resource
+            nanovdb_test_cuda_util nanovdb_test_cuda_util_sync)
+      target_compile_options(${_cuda_gtest} PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:--diag-suppress=177>")
+    endforeach()
+  endif()
 endif()
 
 # -----------------------------------------------------------------------------

diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt
@@ -5,8 +5,11 @@ NanoVDB:
   - Added nanovdb::cuda::Buffer and nanovdb::cuda::BufferView (CUDA): a typed, resource-aware, stream-ordered container that allocates from an injectable memory resource and frees on its retained stream, and a non-owning view over externally managed memory that a GridHandle can wrap without copying. Member names follow cuda::buffer (destroy, set_stream, swap). Also added the synchronous resource concept nanovdb::cuda::is_resource alongside is_async_resource.
 
   Improvements:
+  - Added the CMake option NANOVDB_CUDA_WERROR (default ON, also implied by OPENVDB_CXX_STRICT) which passes --Werror=all-warnings to NVCC so that every device-side diagnostic becomes an error when building the NanoVDB tests, tools and examples. Disable it if a newer CUDA toolkit introduces diagnostics that block your build.
   - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU).
 
   Fixes:
   - Fixed a caching bug in nanovdb::ReadAccessor<LEVEL0, LEVEL1, LEVEL2>, e.g ReadAccessor<0,1,2> and ReadAccessor<0,1>: the accessors now cache nodes at all requested levels (matching the behaviour of openvdb::ValueAccessor) rather than caching only leaf nodes. This is the new default behaviour. The previous (buggy) behaviour can be restored by setting the CMake option -DNANOVDB_USE_OLD_ACCESSOR=ON, or by defining the macro NANOVDB_USE_OLD_ACCESSOR before including NanoVDB.h.
-  - Deprecated nanovdb::math::ZeroCrossingNode and renamed it nanovdb::math::zeroCrossingNode to follow our naming convention.
+  - Deprecated nanovdb::math::ZeroCrossingNode and renamed it nanovdb::math::zeroCrossingNode to follow our naming convention.
+  - Fixed the nvcc #20054-D warnings emitted by the nanovdb::tools::cuda::GridStats kernels: the __shared__ statistics arrays now use constructor-free storage (cub::Uninitialized), since CUDA does not support dynamic initialization of __shared__ variables.
+  - Fixed the nvcc #20050-D warnings emitted by the nanovdb::tools::cuda::GridChecksum CRC32 kernels: the inline qualifier is ignored for __global__ functions, so the header-defined kernels are now static instead.