diff --git a/nanovdb/nanovdb/CMakeLists.txt b/nanovdb/nanovdb/CMakeLists.txt index dede8fdfbb..8f4fa07c56 100644 --- a/nanovdb/nanovdb/CMakeLists.txt +++ b/nanovdb/nanovdb/CMakeLists.txt @@ -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) diff --git a/nanovdb/nanovdb/examples/ex_collide_level_set/common.h b/nanovdb/nanovdb/examples/ex_collide_level_set/common.h index 0b476588fd..f01ab0f2c9 100644 --- a/nanovdb/nanovdb/examples/ex_collide_level_set/common.h +++ b/nanovdb/nanovdb/examples/ex_collide_level_set/common.h @@ -3,7 +3,9 @@ #pragma once +#ifndef _USE_MATH_DEFINES #define _USE_MATH_DEFINES +#endif #include #include #include diff --git a/nanovdb/nanovdb/examples/ex_collide_level_set/nanovdb.cu b/nanovdb/nanovdb/examples/ex_collide_level_set/nanovdb.cu index 98f9e5128f..312e93975a 100644 --- a/nanovdb/nanovdb/examples/ex_collide_level_set/nanovdb.cu +++ b/nanovdb/nanovdb/examples/ex_collide_level_set/nanovdb.cu @@ -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 #include diff --git a/nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc b/nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc index 7c71392d32..3a4c456376 100644 --- a/nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc +++ b/nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc @@ -3,7 +3,9 @@ #if defined(NANOVDB_USE_OPENVDB) +#ifndef _USE_MATH_DEFINES #define _USE_MATH_DEFINES +#endif #include #include diff --git a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/common.h b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/common.h index dc0e601f26..477fd97c6d 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/common.h +++ b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/common.h @@ -3,7 +3,9 @@ #pragma once +#ifndef _USE_MATH_DEFINES #define _USE_MATH_DEFINES +#endif #include #include #include diff --git a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/nanovdb.cu b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/nanovdb.cu index a7b9c94a5f..e0e7577610 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/nanovdb.cu +++ b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/nanovdb.cu @@ -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 #include diff --git a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc index f622c4da0e..16360157fd 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc +++ b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc @@ -3,7 +3,9 @@ #if defined(NANOVDB_USE_OPENVDB) +#ifndef _USE_MATH_DEFINES #define _USE_MATH_DEFINES +#endif #include #include diff --git a/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/common.h b/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/common.h index 8b3d6bb447..551e702b4b 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/common.h +++ b/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/common.h @@ -3,7 +3,9 @@ #pragma once +#ifndef _USE_MATH_DEFINES #define _USE_MATH_DEFINES +#endif #include #include #include diff --git a/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu b/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu index 1e53c65a9c..3eafd79ce8 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu +++ b/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu @@ -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 #include diff --git a/nanovdb/nanovdb/examples/ex_raytrace_level_set/common.h b/nanovdb/nanovdb/examples/ex_raytrace_level_set/common.h index dc0e601f26..477fd97c6d 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_level_set/common.h +++ b/nanovdb/nanovdb/examples/ex_raytrace_level_set/common.h @@ -3,7 +3,9 @@ #pragma once +#ifndef _USE_MATH_DEFINES #define _USE_MATH_DEFINES +#endif #include #include #include diff --git a/nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu b/nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu index 7b92541b9d..1b5e9084d3 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu +++ b/nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu @@ -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 #include diff --git a/nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc b/nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc index b8ca62c838..75ee170acc 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc +++ b/nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc @@ -3,7 +3,9 @@ #if defined(NANOVDB_USE_OPENVDB) +#ifndef _USE_MATH_DEFINES #define _USE_MATH_DEFINES +#endif #include #include diff --git a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh index da098c98f7..5a0ca7a762 100644 --- a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh @@ -617,7 +617,6 @@ void DistributedPointsToGrid::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; diff --git a/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh b/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh index f5585e8b96..100a405416 100644 --- a/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh +++ b/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh @@ -115,7 +115,7 @@ inline unique_ptr 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) { diff --git a/nanovdb/nanovdb/tools/cuda/GridStats.cuh b/nanovdb/nanovdb/tools/cuda/GridStats.cuh index f2c85f89d5..1705c629f2 100644 --- a/nanovdb/nanovdb/tools/cuda/GridStats.cuh +++ b/nanovdb/nanovdb/tools/cuda/GridStats.cuh @@ -18,6 +18,8 @@ #include #include +#include // for cub::Uninitialized + namespace nanovdb { namespace tools::cuda { @@ -65,7 +67,11 @@ template __global__ void processLeaf(NodeManager *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 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 *d_nodeMgr, StatsT *d_stats) using ChildT = typename NanoNode::type; using NodeT = typename NanoNode::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 sStatsStorage; + __shared__ cub::Uninitialized sBBoxStorage; + StatsT (&sStats)[Threads] = sStatsStorage.Alias(); + CoordBBox (&sBBox)[Threads] = sBBoxStorage.Alias(); const uint32_t nodeID = blockIdx.x; const uint32_t tID = threadIdx.x; diff --git a/nanovdb/nanovdb/unittest/CMakeLists.txt b/nanovdb/nanovdb/unittest/CMakeLists.txt index 5538a3d04e..439f05610a 100644 --- a/nanovdb/nanovdb/unittest/CMakeLists.txt +++ b/nanovdb/nanovdb/unittest/CMakeLists.txt @@ -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 "$<$:--diag-suppress=177>") + endforeach() + endif() endif() # ----------------------------------------------------------------------------- diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index 906e6fdaeb..cb3a42d13c 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -6,8 +6,11 @@ NanoVDB: Improvements: - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. PointsToGrid's device scratch and intermediate arrays are now owned by cuda::Buffer as well, replacing all of its hand-paired allocate/free calls. + - 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, 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. \ No newline at end of file + - 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. \ No newline at end of file