NanoVDB: fix nvcc warnings and treat NVCC diagnostics as errors by default

[harrism]

Closes #2279.

Fixes two nvcc warning classes in the NanoVDB CUDA tool headers, then turns on warnings-as-errors for NVCC so they stay fixed.

#20054-D — GridStats.cuh (12 instances)

Stats/Extrema/CoordBBox have user-provided default constructors, so the __shared__ scratch arrays in processLeaf/processInternal require dynamic initialization, which CUDA does not support — nvcc skips the constructors and warns. The skipped initialization was always harmless: both kernels fully assign every shared slot before reading it, and the early-return paths exit at whole-warp/whole-block granularity, so no partially-written slice enters a reduction.

Fix: wrap the scratch arrays in cub::Uninitialized (<cub/util_type.cuh>, the same mechanism CUB uses for its own TempStorage): its backing store is a trivially-constructible word array, so no dynamic init is required, and Alias() re-exposes it as the element type. Array-reference bindings keep the names sStats/sBBox, so the reduction code is unchanged.

#20050-D — GridChecksum.cuh (2 kernels)

inline is ignored on __global__ functions. It was doing ODR duty for non-template kernels defined in a header, so it is replaced with static (per-TU internal linkage — the same effect as the anonymous namespace GridStats.cuh uses). Both launch sites are in the same header, so nothing needs external linkage.

Note: this class is invisible in CMake builds — CMake passes the CUDA toolkit include via -isystem, and the diagnostic is attributed through the __global__ macro machinery in those system headers, so it is suppressed (verified by flag bisection: removing only -isystem .../cuda/targets/x86_64-linux/include from the exact CMake compile line makes the warnings appear). Plain-nvcc consumers of the headers do see it.

NVCC warnings-as-errors (NANOVDB_CUDA_WERROR, default ON)

New CMake option that prepends --Werror=all-warnings to CMAKE_CUDA_FLAGS, so every nvcc front-end diagnostic fails the build of the NanoVDB tests, tools and examples. Also implied by OPENVDB_CXX_STRICT. It affects only NanoVDB's own build tree (not projects that merely include the headers) and can be set to OFF if a newer CUDA toolkit introduces new diagnostics.

Host-side -Werror (the OPENVDB_CXX_STRICT flag set applied to the nano components) is not included: a measurement pass with the strict flags replicated locally found ~74 (gcc) + ~33 (clang) unique pre-existing sites — dominated by unused parameters, -Wunused-function in the C-portable headers (CNanoVDB.h/PNanoVDB.h, where file-static functions are the design), sign-compare inside Debug-only assert bodies, and clang -Wconversion in core headers. That cleanup deserves its own PR(s) and a policy decision for the C-portable headers.

Verification

• Warnings: a TU including essentially every CUDA tool header (TestNanoVDB.cu) compiles with zero warnings in both the CMake configuration and the warning-visible plain-nvcc configuration (was 12× #20054-D + 4× #20050-D).
• Werror builds: full Release and Debug builds at CMAKE_CUDA_ARCHITECTURES=80 (matching CI) with unit tests, tools and examples enabled — zero diagnostics under --Werror=all-warnings; 7/7 ctests pass on a GPU runner (2× RTX 6000 Ada, CUDA 12.x).
• Performance:
  • Per-kernel SASS comparison on sm_89 (GridStats): all 14 kernels have identical instruction counts, register usage, shared-memory bytes, and zero spills. 12 of 14 byte-identical modulo symbol-hash/data-offset relocation; the two processLeaf stats variants differ only in register coloring and code layout (4 NOPs became 4 register moves in one). The GridChecksum change is linkage-only and does not affect codegen.
  • Timing: 200 iterations of updateGridStats on an 8M-voxel grid (1024³ box), interleaved rounds per binary: StatsMode::All 12.28–12.30 ms/iter and StatsMode::MinMax 2.55 ms/iter for both; matched-round deltas ±0.1% with alternating sign, smaller than either binary's run-to-run drift.

🤖 Generated with Claude Code

[harrism]

CI is now green on all 8 jobs with NANOVDB_CUDA_WERROR defaulted ON — including the first Windows (MSVC-host nvcc) builds under the flag, which surfaced and fixed four pre-existing issue classes across three commits:

1. 418× dll-interface diagnostics (#1394/#1388) from OpenVDB core headers included by .cu files — suppressed on Windows with --diag-suppress, mirroring the unconditional /wd4251//wd4275 policy in OpenVDBCXX.cmake. (The pre-existing -Xcudafe --diag_suppress spelling demonstrably does not reach the front-end phase that emits these.)
2. A real dead-code catch: deviceCoords in DistributedPointsToGrid.cuh, orphaned by the earlier removal of the cudaMemAdvise on external coords. Deleted.
3. 76× #177 on gtest's test_info_: gtest marks it __attribute__((unused)) on GCC/Clang but has no MSVC equivalent, so tests declared inside anonymous namespaces trip the front end. Suppressed per-target for the four CUDA gtest executables on Windows only — #177 stays enforced for all library headers, tools and examples everywhere, and for everything on Linux.
4. C4005 macro redefinition (_USE_MATH_DEFINES) in the raytrace/collide examples, which define it in source while the Windows build also passes it via -D — the flag forwards /WX to the MSVC host pass, promoting these. All 11 definition sites are now #ifndef-guarded (verified by a tree-wide sweep).

Net effect: the NanoVDB build is warning-clean under --Werror=all-warnings on Linux (gcc + clang × Release + Debug), Windows (MSVC 2022 + CUDA 12.4), and the flag-free macOS/lite configurations.

🤖 Generated with Claude Code