From f3e3bdbdd026c455b4fc34d8b2fd4b19110b5bb9 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 Aug 2026 07:22:06 +0000 Subject: [PATCH 1/6] NanoVDB: use constructor-free shared-memory storage in GridStats kernels CUDA does not support dynamic initialization of function-scope __shared__ variables, so nvcc skips the StatsT/CoordBBox constructors in processLeaf and processInternal and emits warning #20054-D -- twelve copies in every TU that instantiates the kernels, propagated to downstream builds since GridStats.cuh is a header. The skipped initialization was harmless (every shared slot is fully assigned before it is read), but the noise is not. Wrap the scratch arrays in cub::Uninitialized, whose backing store is a trivially-constructible word array, and bind array references to the aliased element type. The reduction code is unchanged. Verified warning-free (0 vs 12) and performance-neutral: per-kernel SASS comparison on sm_89 shows identical instruction counts, registers, shared-memory bytes, and zero spills for all 14 kernels; timing 200 iterations of updateGridStats on an 8M-voxel grid shows matched-round deltas within +/-0.1% with alternating sign. Closes #2279 Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/GridStats.cuh | 16 +++++++++++++--- pendingchanges/nanovdb.txt | 3 ++- 2 files changed, 15 insertions(+), 4 deletions(-) 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/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index f342f5c47a..c599212090 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -9,4 +9,5 @@ NanoVDB: 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. \ No newline at end of file From c258a552db8d5508cc1be931f0ede071d70e6ca9 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 Aug 2026 09:42:40 +0000 Subject: [PATCH 2/6] 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 Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/GridChecksum.cuh | 4 ++-- pendingchanges/nanovdb.txt | 3 ++- 2 files changed, 4 insertions(+), 3 deletions(-) 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/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index c599212090..cbf6ee5dcd 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -10,4 +10,5 @@ NanoVDB: 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. - - 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. \ No newline at end of file + - 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 From 880f4c6668b7991d184f7e9a786159e245e484f2 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 Aug 2026 12:35:30 +0000 Subject: [PATCH 3/6] NanoVDB: treat NVCC diagnostics as errors by default Add the CMake option NANOVDB_CUDA_WERROR (default ON, also implied by OPENVDB_CXX_STRICT), which prepends --Werror=all-warnings to CMAKE_CUDA_FLAGS so every nvcc front-end diagnostic fails the build of the NanoVDB tests, tools and examples. This locks in the warning-free state: the #177-D, #20050-D and #20054-D classes fixed recently cannot silently reappear. The option only affects NanoVDB's own build tree, not projects that merely include the headers, and can be disabled if a newer CUDA toolkit introduces new diagnostics. Verified locally at CMAKE_CUDA_ARCHITECTURES=80 (matching CI) in both Release and Debug with unit tests, tools and examples enabled: zero diagnostics, full builds succeed, 7/7 ctests pass on a GPU. Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/CMakeLists.txt | 8 ++++++++ pendingchanges/nanovdb.txt | 1 + 2 files changed, 9 insertions(+) diff --git a/nanovdb/nanovdb/CMakeLists.txt b/nanovdb/nanovdb/CMakeLists.txt index dede8fdfbb..5202ad2d5c 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,13 @@ 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() + # workaround for win32 bug when nvcc "--keep" is used. if(WIN32) if(NANOVDB_CUDA_KEEP_PTX) diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index cbf6ee5dcd..03772668f1 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -5,6 +5,7 @@ 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: From 9d94edfb39e0fc0ef5fdbd12bbaf9449427319cc Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 Aug 2026 13:35:32 +0000 Subject: [PATCH 4/6] NanoVDB: fix Windows fallout from NVCC warnings-as-errors The first Windows CI run under --Werror=all-warnings surfaced two issues: Remove an unreferenced local in DistributedPointsToGrid's stripe setup. deviceCoords became dead when the cudaMemAdvise call on the external coords was removed; only the MSVC-host front end reports it (#177-D), but the variable is dead on every platform. Suppress the EDG dll-interface diagnostics (#1394, #1388) for CUDA sources on Windows. This is the 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 these fire on every include of the OpenVDB core headers from a .cu file (404 + 14 instances in the CI log). Uses the top-level --diag-suppress option; the -Xcudafe --diag_suppress spelling does not reach the CUDA front-end phase that emits these. Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/CMakeLists.txt | 8 ++++++++ nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh | 1 - 2 files changed, 8 insertions(+), 1 deletion(-) diff --git a/nanovdb/nanovdb/CMakeLists.txt b/nanovdb/nanovdb/CMakeLists.txt index 5202ad2d5c..8f4fa07c56 100644 --- a/nanovdb/nanovdb/CMakeLists.txt +++ b/nanovdb/nanovdb/CMakeLists.txt @@ -104,6 +104,14 @@ if(NANOVDB_USE_CUDA) 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/tools/cuda/DistributedPointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh index 7d7875b30e..7db4f866d8 100644 --- a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh @@ -636,7 +636,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; From 233b6f35b389daafc30bb78e5cf972dd98b358e0 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 Aug 2026 14:26:20 +0000 Subject: [PATCH 5/6] NanoVDB: fix remaining Windows warnings-as-errors fallout Clearing the dll-interface diagnostics unblocked the Windows projects to compile further, exposing two more pre-existing warning classes that NANOVDB_CUDA_WERROR now promotes to errors: Guard the _USE_MATH_DEFINES definition in the four examples that define it. The Windows build also passes -D_USE_MATH_DEFINES on the command line (as 1), so the bare #define (empty) triggered MSVC C4005 macro-redefinition warnings, which the flag turns into C2220 errors. Suppress EDG #177 for the CUDA gtest targets on Windows. 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 (TestBuffer.cu, TestMemoryResource.cu, TestUtilCuda.cu) the front end proves it unreferenced and reports #177 -- 76 errors, all from the macro machinery, none actionable in our code. TestNanoVDB.cu is unaffected because its tests are declared at namespace scope. The suppression is per-target and Windows-only, so #177 stays enforced for all library headers, tools and examples on every platform, and for everything on Linux. Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- .../nanovdb/examples/ex_collide_level_set/nanovdb.cu | 2 ++ .../examples/ex_raytrace_fog_volume/nanovdb.cu | 2 ++ .../examples/ex_raytrace_iso_surface/nanovdb.cu | 2 ++ .../examples/ex_raytrace_level_set/nanovdb.cu | 2 ++ nanovdb/nanovdb/unittest/CMakeLists.txt | 12 ++++++++++++ 5 files changed, 20 insertions(+) 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_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_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/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/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() # ----------------------------------------------------------------------------- From 888e446df08c50040b1c1eef6678910b854104f3 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 Aug 2026 15:11:54 +0000 Subject: [PATCH 6/6] NanoVDB: guard _USE_MATH_DEFINES in the example headers and CXX sources The previous commit guarded only the nanovdb.cu files that the first Windows log reported; each of these examples defines the macro in up to three files, and common.h is included by the .cu sources, so the C4005 redefinition warnings persisted. Every definition site in the examples is now guarded, verified by a tree-wide sweep for unguarded instances. Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/examples/ex_collide_level_set/common.h | 2 ++ nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc | 2 ++ nanovdb/nanovdb/examples/ex_raytrace_fog_volume/common.h | 2 ++ nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc | 2 ++ nanovdb/nanovdb/examples/ex_raytrace_iso_surface/common.h | 2 ++ nanovdb/nanovdb/examples/ex_raytrace_level_set/common.h | 2 ++ nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc | 2 ++ 7 files changed, 14 insertions(+) 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/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/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_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/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