ex_voxelBlockManager_host_cuda: replace OpenMP threading with util::forEach

Use nanovdb::util::forEach (TBB-backed) for parallel block iteration in
the CPU benchmarks, removing the OpenMP threading dependency. Keep
#pragma omp simd as a pure vectorization hint.

This eliminates the OpenMP linker/include issues on clang++ and also
improves performance: TBB work-stealing yields ~3-4x faster sanity
checks and more consistent full-decode timings vs. OMP schedule(static).

Thread-private scratch globals are replaced by lambda-local arrays;
the reduction dummy becomes std::atomic<uint32_t>. CMakeLists no longer
needs find_package(OpenMP) or OpenMP::OpenMP_CXX for this target.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

Author: sifakis

nanovdb/nanovdb/examples/CMakeLists.txt

@@ -116,15 +116,9 @@ nanovdb_example(NAME "ex_coarsen_nanovdb_cuda" OPENVDB)
 
 nanovdb_example(NAME "ex_voxelBlockManager_host_cuda")
 if(TARGET ex_voxelBlockManager_host_cuda)
-  find_package(OpenMP)
   target_compile_options(ex_voxelBlockManager_host_cuda PRIVATE
     $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-mavx2>
     $<$<COMPILE_LANGUAGE:CXX>:-mavx2>)
-  if(OpenMP_CXX_FOUND)
-    target_link_libraries(ex_voxelBlockManager_host_cuda PRIVATE OpenMP::OpenMP_CXX)
-    target_compile_options(ex_voxelBlockManager_host_cuda PRIVATE
-      $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-fopenmp>)
-  endif()
 endif()
 
 if(CUDAToolkit_FOUND)

nanovdb/nanovdb/examples/ex_voxelBlockManager_host_cuda/vbm_host_cuda_kernels.cu

@@ -16,20 +16,17 @@
 #include <nanovdb/HostBuffer.h>
 #include <nanovdb/tools/cuda/PointsToGrid.cuh>
 #include <nanovdb/tools/VoxelBlockManager.h>
+#include <nanovdb/util/ForEach.h>
 #include <nanovdb/util/Timer.h>
 #include <nanovdb/util/cuda/Timer.h>
 #include <nanovdb/util/cuda/Util.h>
 
 #include <thrust/universal_vector.h>
 
 #include <immintrin.h>
-#ifdef _OPENMP
-#include <omp.h>
-#else
-inline int omp_get_max_threads() { return 1; }
-#endif
 
 #include <algorithm>
+#include <atomic>
 #include <vector>
 #include <iostream>
 
@@ -38,13 +35,6 @@ namespace {
 static constexpr int Log2BlockWidth = 7;
 static constexpr int BlockWidth     = 1 << Log2BlockWidth; // 128
 
-// Thread-private scratch arrays for CPU decodeInverseMaps benchmarks.
-// Declared at namespace scope with static storage so threadprivate is valid;
-// aligned to a cache line to avoid false sharing between threads.
-alignas(64) static uint32_t g_scratchLeafIndex[BlockWidth];
-alignas(64) static uint16_t g_scratchVoxelOffset[BlockWidth];
-#pragma omp threadprivate(g_scratchLeafIndex, g_scratchVoxelOffset)
-
 using VBM    = nanovdb::tools::cuda::VoxelBlockManager<Log2BlockWidth>;
 using CPUVBM = nanovdb::tools::VoxelBlockManager<Log2BlockWidth>;
 
@@ -381,23 +371,19 @@ void runVBMCudaTest(const std::vector<nanovdb::Coord>& coords)
         static constexpr int nPerfRuns = 20;
 
         // --- Sanity check: countOn + std::fill only, no tree access ---
-        // If OpenMP parallelism is working, this should scale near-linearly.
         {
-            std::cout << "\nCPU decodeInverseMaps sanity (countOn only, "
-                      << omp_get_max_threads() << " OMP threads):\n";
+            std::cout << "\nCPU decodeInverseMaps sanity (countOn only):\n";
 
-            // Never true: nExtraLeaves <= JumpMapLength*64, so == JumpMapLength*64+1 is impossible.
-            // Used only to prevent dead code elimination without a reduction.
+            // Never true: used only to prevent dead code elimination.
             volatile uint32_t dummy = 0;
 
             for (int run = 0; run < nPerfRuns; ++run) {
                 nanovdb::util::Timer timer;
                 timer.start("");
 
-                #pragma omp parallel
-                {
-                    #pragma omp for schedule(static)
-                    for (uint32_t bID = 0; bID < nBlocks; ++bID) {
+                nanovdb::util::forEach(0, nBlocks, 1,
+                    [&](const nanovdb::util::Range1D& range) {
+                    for (auto bID = range.begin(); bID < range.end(); ++bID) {
                         int nExtraLeaves = 0;
                         for (int i = 0; i < CPUVBM::JumpMapLength; i++)
                             nExtraLeaves += nanovdb::util::countOn(
@@ -443,28 +429,28 @@ void runVBMCudaTest(const std::vector<nanovdb::Coord>& coords)
                         if (prefixCountRealigned[31][15] == 513u)
                             dummy = prefixCountRealigned[31][15];
                     }
-                }
+                });
 
                 const float ms =
                     (float)timer.elapsed<std::chrono::microseconds>() / 1000.0f;
                 std::cout << "  run " << run << ": " << ms << " ms\n";
             }
         }
 
-        uint32_t dummy = 0;
+        std::atomic<uint32_t> dummy{0};
 
         std::cout << "\nCPU decodeInverseMaps performance ("
-                  << nBlocks << " blocks, " << nVoxels << " voxels"
-                  << ", " << omp_get_max_threads() << " OMP threads):\n";
+                  << nBlocks << " blocks, " << nVoxels << " voxels):\n";
 
         for (int run = 0; run < nPerfRuns; ++run) {
             nanovdb::util::Timer timer;
             timer.start("");
 
-            #pragma omp parallel reduction(+:dummy)
-            {
-                #pragma omp for schedule(static)
-                for (uint32_t bID = 0; bID < nBlocks; ++bID) {
+            nanovdb::util::forEach(0, nBlocks, 1,
+                [&](const nanovdb::util::Range1D& range) {
+                alignas(64) uint32_t scratchLeafIndex[BlockWidth];
+                alignas(64) uint16_t scratchVoxelOffset[BlockWidth];
+                for (auto bID = range.begin(); bID < range.end(); ++bID) {
                     const uint64_t blockFirstOffset =
                         cpuHandle.firstOffset() + (uint64_t)bID * BlockWidth;
 
@@ -473,18 +459,18 @@ void runVBMCudaTest(const std::vector<nanovdb::Coord>& coords)
                         firstLeafIDPtr[bID],
                         &jumpMapPtr[(uint64_t)bID * CPUVBM::JumpMapLength],
                         blockFirstOffset,
-                        g_scratchLeafIndex,
-                        g_scratchVoxelOffset);
+                        scratchLeafIndex,
+                        scratchVoxelOffset);
 
-                    dummy += (g_scratchLeafIndex[0]  != CPUVBM::UnusedLeafIndex)   ? 1u : 0u;
-                    dummy += (g_scratchVoxelOffset[0] != CPUVBM::UnusedVoxelOffset) ? 1u : 0u;
+                    dummy += (scratchLeafIndex[0]  != CPUVBM::UnusedLeafIndex)   ? 1u : 0u;
+                    dummy += (scratchVoxelOffset[0] != CPUVBM::UnusedVoxelOffset) ? 1u : 0u;
                 }
-            }
+            });
 
             const float ms =
                 (float)timer.elapsed<std::chrono::microseconds>() / 1000.0f;
             std::cout << "  run " << run << ": " << ms << " ms\n";
         }
-        std::cout << "  (dummy=" << dummy << ")\n";
+        std::cout << "  (dummy=" << dummy.load() << ")\n";
     }
 }