[svt] add sparse volume traversal benchmark - #325

Open

zjin-lcf wants to merge 3 commits into

zjin-lcf commented Aug 22, 2026

Collaborator

Summary

add a self-contained sparse volume traversal benchmark with a CPU reference
provide CUDA, HIP, SYCL, and OpenMP target-offload implementations
register SVT in CMake and document it in the benchmark catalog

Test plan

Build and run HIP on AMD MI100 (PASS)
Build and run SYCL on AMD MI100 (PASS)
Build and run OpenMP target offload on AMD MI100 (PASS)
Build and run CUDA on an NVIDIA GPU

Made with Cursor


          [svt] add sparse volume traversal benchmark

c558a16

Add portable CUDA, HIP, SYCL, and OpenMP implementations with reference validation so hierarchical sparse-grid traversal can be compared across programming models.

Co-authored-by: Cursor <cursoragent@cursor.com>

zjin-lcf requested a review from Geekdude

August 24, 2026 14:21

Geekdude requested a balanced review from Copilot

August 26, 2026 01:25

Copilot started reviewing on behalf of Geekdude

August 26, 2026 01:25

View session

Copilot AI reviewed

View reviewed changes

Copilot AI left a comment

Pull request overview

Adds a sparse-volume traversal benchmark with a shared CPU reference across CUDA, HIP, SYCL, and OpenMP offload.

Changes:

Implements two-level DDA TSDF ray casting and verification.
Adds backend-specific build and test registration.
Registers and documents SVT as a simulation benchmark.

Reviewed changes

Copilot reviewed 17 out of 17 changed files in this pull request and generated 11 comments.

Show a summary per file

File	Description
`src/svt-cuda/main.cu`	CUDA implementation
`src/svt-cuda/reference.h`	Shared CPU reference and volume generation
`src/svt-cuda/Makefile`	CUDA standalone build
`src/svt-cuda/CMakeLists.txt`	CUDA CMake registration
`src/svt-hip/main.cu`	HIP implementation
`src/svt-hip/Makefile`	HIP standalone build
`src/svt-hip/CMakeLists.txt`	HIP CMake registration
`src/svt-sycl/main.cpp`	SYCL implementation
`src/svt-sycl/Makefile`	SYCL standalone build
`src/svt-sycl/CMakeLists.txt`	SYCL CMake registration
`src/svt-omp/main.cpp`	OpenMP offload implementation
`src/svt-omp/Makefile`	Intel OpenMP build
`src/svt-omp/Makefile.aomp`	AOMP build
`src/svt-omp/Makefile.nvc`	NVIDIA HPC SDK build
`src/svt-omp/CMakeLists.txt`	OpenMP CMake registration
`src/CMakeLists.txt`	Registers the SVT benchmark
`README.md`	Adds SVT to the benchmark catalog

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

src/svt-cuda/main.cu Outdated

Comment on lines +174 to +177

+                      if (vi[0] < 0 || vi[0] >= LEAF_DIM ||
+                          vi[1] < 0 || vi[1] >= LEAF_DIM ||
+                          vi[2] < 0 || vi[2] >= LEAF_DIM ||
+                          cur_t > t_next) break;

zjin-lcf Aug 26, 2026

Collaborator Author

address the issue in the updated PR

src/svt-hip/main.cu Outdated

Comment on lines +159 to +162

+                      if (vi[0] < 0 || vi[0] >= LEAF_DIM ||
+                          vi[1] < 0 || vi[1] >= LEAF_DIM ||
+                          vi[2] < 0 || vi[2] >= LEAF_DIM ||
+                          cur_t > t_next) break;

zjin-lcf Aug 26, 2026

Collaborator Author

address the issue in the updated PR

src/svt-sycl/main.cpp Outdated

Comment on lines +142 to +145

+                      if (vi[0] < 0 || vi[0] >= LEAF_DIM ||
+                          vi[1] < 0 || vi[1] >= LEAF_DIM ||
+                          vi[2] < 0 || vi[2] >= LEAF_DIM ||
+                          cur_t > t_next) break;

zjin-lcf Aug 26, 2026

Collaborator Author

address the issue in the updated PR

src/svt-omp/main.cpp Outdated

Comment on lines +156 to +159

+                      if (vi[0] < 0 || vi[0] >= LEAF_DIM ||
+                          vi[1] < 0 || vi[1] >= LEAF_DIM ||
+                          vi[2] < 0 || vi[2] >= LEAF_DIM ||
+                          cur_t > t_next) break;

zjin-lcf Aug 26, 2026

Collaborator Author

address the issue in the updated PR

src/svt-cuda/reference.h Outdated

Comment on lines +356 to +359

+                      if (vi[0] < 0 || vi[0] >= LEAF_DIM ||
+                          vi[1] < 0 || vi[1] >= LEAF_DIM ||
+                          vi[2] < 0 || vi[2] >= LEAF_DIM ||
+                          cur_t > t_next) break;

zjin-lcf Aug 26, 2026

Collaborator Author

address the issue in the updated PR

src/svt-omp/main.cpp Outdated

Comment on lines +298 to +299

		#pragma omp target teams distribute parallel for collapse(2) \
		num_teams(num_teams) thread_limit(thread_limit) nowait

zjin-lcf Aug 26, 2026

Collaborator Author

nowait is disabled in the updated PR

src/svt-cuda/main.cu

+                CHECK(cudaFree(d_leaf));
+                CHECK(cudaFree(d_image));
+                return 0;

src/svt-hip/main.cu

+                CHECK(hipFree(d_leaf));
+                CHECK(hipFree(d_image));
+                return 0;

zjin-lcf Aug 26, 2026

Collaborator Author

the issue may be addressed across all hecbench programs

src/svt-sycl/main.cpp

+                sycl::free(d_leaf,  q);
+                sycl::free(d_image, q);
+                return 0;

zjin-lcf Aug 26, 2026

Collaborator Author

the issue may be addressed across all hecbench programs

src/svt-omp/main.cpp

+                  pass = pass_result;
+                }
+                return 0;

zjin-lcf Aug 26, 2026

Collaborator Author

the issue may be addressed across all hecbench programs

zjin-lcf added 2 commits

August 26, 2026 09:10


          [svt] increase the runs from 100 to 1K in Makefiles; address the gap:…

bcce20c

… The ray casting is a two-level DDA (digital differential analyzer): an outer loop steps across coarse blocks in the sparse grid, and for each active block, an inner loop steps voxel-by-voxel through that block leaf looking for a sign change in the TSDF (truncated signed distance field) value — a positive-to-non-positive transition marks a surface crossing. When the ray exits one leaf and the outer loop advances to the next block, nothing carries over the last voxel sample that was examined in the leaf just left. So if the final voxel of leaf A had a positive TSDF (still outside the surface) and the first voxel of the neighboring leaf B is non-positive (now inside the surface), that is zero-crossing — a real surface — sitting right at the seam between two leaves. But because prev for leaf B is initialized from that same first voxel of B rather than from the last voxel of A, the comparison prev > 0 && cur <= 0 is never made across that specific pair. The surface at that boundary is silently skipped, i.e. a hole appears at every point where a surface happens to fall exactly on a leaf boundary


          [svt] disable nowait in the OMP program

50e3a62

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

Labels

None yet