NanoVDB: make GridChecksum's header-defined CRC32 kernels static

nvcc ignores the inline qualifier on __global__ functions and emits
warning #20050-D for crc32SlicedKernel and crc32CombineKernel. The
inline was doing ODR duty for a non-template kernel defined in a
header, so removing it outright would invite multiple-definition link
errors; static provides internal linkage per TU instead, matching the
effect of the anonymous namespace used by GridStats.cuh. Both launch
sites live in the same header, so nothing needs external linkage.

Note these warnings are invisible in CMake builds: CMake passes the
CUDA toolkit include directory via -isystem, and the diagnostic is
attributed through the __global__ macro machinery in those system
headers, so it is suppressed. Plain nvcc consumers of the headers see
it. With this change, TestNanoVDB.cu -- which includes essentially
every CUDA tool header -- compiles warning-free even in the
warning-visible (non -isystem) configuration.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

Author: harrism

nanovdb/nanovdb/tools/cuda/GridChecksum.cuh

@@ -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) {

pendingchanges/nanovdb.txt

@@ -10,4 +10,5 @@ NanoVDB:
   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.
-  - 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 #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.