From c4e228dabebaa4f9201ce588fdf689af09e58942 Mon Sep 17 00:00:00 2001 From: auc7us Date: Fri, 4 Sep 2026 12:01:10 -0500 Subject: [PATCH 1/5] Make the SCM GPU kernels a single source compiled as either CUDA or HIP Renames the two kernel translation units from .hip.cpp to .cu and removes their dependence on the HIP runtime, so the same source can be compiled by nvcc or by hipcc. This is the arrangement Chrono::DEM and Chrono::FSI::SPH already use, and it reuses their machinery: chrono_set_gpu_source_language() relabels the sources LANGUAGE HIP when the HIP backend is selected, and .cu is CUDA by default. The device code needed no changes -- __global__, __shared__, threadIdx and the <<<>>> launch are spelled identically in both. Only two runtime names appeared in these files, hipStream_t and hipGetLastError, and they are now selected by a short conditional at the top of each file. That conditional is not optional: nvcc implicitly includes for a .cu, but HIP-clang provides nothing, so a .cu with no includes compiles under nvcc and fails under hipcc with threadIdx, blockIdx and blockDim undeclared. The host bridges are untouched and remain HIP-only, so CH_ENABLE_VEHICLE_SCM_GPU still REQUIRES HIP and every existing build behaves exactly as before. They use around thirty runtime symbols and will be duplicated per backend rather than routed through a compatibility layer, for which Chrono has no precedent. Verified on both toolchains: SCMRaycastGpuKernels.cu nvcc -c PASS SCMRaycastGpuKernels.cu hipcc -fsyntax-only --offload-arch=gfx942 PASS SCMGpuKernels.cu nvcc -c PASS SCMGpuKernels.cu hipcc -fsyntax-only --offload-arch=gfx942 PASS and the full HIP build of Chrono_vehicle links unchanged. --- PR_MESSAGE.md | 37 ++++++++++++++++ src/chrono_vehicle/CMakeLists.txt | 19 +++++--- ...SCMGpuKernels.hip.cpp => SCMGpuKernels.cu} | 43 ++++++++++++++++--- ...ernels.hip.cpp => SCMRaycastGpuKernels.cu} | 41 +++++++++++++++--- 4 files changed, 120 insertions(+), 20 deletions(-) create mode 100644 PR_MESSAGE.md rename src/chrono_vehicle/terrain/gpu/{SCMGpuKernels.hip.cpp => SCMGpuKernels.cu} (73%) rename src/chrono_vehicle/terrain/gpu/{SCMRaycastGpuKernels.hip.cpp => SCMRaycastGpuKernels.cu} (85%) diff --git a/PR_MESSAGE.md b/PR_MESSAGE.md new file mode 100644 index 0000000000..0db3af76ec --- /dev/null +++ b/PR_MESSAGE.md @@ -0,0 +1,37 @@ +# [UPDATE] Improve SCM memory use and long-run performance + +**Summary** + +This PR improves SCM performance and memory use, especially for large terrain patches and long simulations. + +- Store SCM height and node data in single precision. This cuts the memory used by these fields roughly in half and allows larger or finer terrain grids. +- Use wider counters when creating the SCM visualization mesh. This prevents integer overflow on very large grids and reports a clear error when the mesh index limit is exceeded. +- Store deformed nodes in 16 x 16 tiles. This keeps terrain lookups fast as the vehicle covers more ground and prevents long runs from gradually slowing down. +- Add active-domain and deformed-node counters, along with a way to remove active domains. This makes it easier to monitor SCM work and stop processing bodies that no longer need terrain interaction. +- Limit rendering in affected demos to 60 frames per simulated second. This avoids unnecessary rendering work and makes reported simulation performance more meaningful. + +**Related Issue(s)** + +None. + +**Author(s)** + +Keshav Sharan + +**Licensing** + +By submitting this pull request, I agree that my contribution will be included in Chrono and redistributed under the BSD-3-Clause License. + +**Backward Compatibility** + +There are no input or public API breaks. Public SCM values remain double precision. Small numerical differences are possible because internal terrain storage now uses single precision. + +**Implementation Notes** + +The changes were tested with the Curiosity SCM demo in the `chrono-orb` container. + +**Post Submission Checklist** + +- [x] The changes are complete +- [x] The changes build with CMake +- [x] The SCM demo was tested diff --git a/src/chrono_vehicle/CMakeLists.txt b/src/chrono_vehicle/CMakeLists.txt index a96ab3f35c..92e9601fb2 100644 --- a/src/chrono_vehicle/CMakeLists.txt +++ b/src/chrono_vehicle/CMakeLists.txt @@ -38,8 +38,9 @@ set(CH_USE_SCM_GPU OFF) # The GPU requirement is declared for this FEATURE, not for Chrono::Vehicle: # the module builds fine with no GPU backend at all, and only the optional SCM -# GPU path needs one. The kernels are written in HIP and have no CUDA variant, -# hence REQUIRES HIP. Both HIP platforms are accepted: nothing here depends on +# GPU path needs one. The kernels themselves are backend-neutral (single-source .cu, +# compiled as CUDA or HIP), but the host bridges are still HIP-only, hence REQUIRES HIP +# for now; this becomes CUDA_OR_HIP once the CUDA host bridges land. Both HIP platforms are accepted: nothing here depends on # the ROCm-only libraries (hipCUB/rocPRIM, rocThrust) that restrict # Chrono::DEM and Chrono::FSI::SPH to the amd platform. So "HIP + ROCm on AMD" # and "HIP + nvcc on NVIDIA" both give the GPU path with no further @@ -318,11 +319,15 @@ endif() source_group("terrain" FILES ${CV_TERRAIN_FILES}) if(CH_USE_SCM_GPU) - set(CV_SCM_GPU_HIP_FILES - terrain/gpu/SCMGpuKernels.hip.cpp - terrain/gpu/SCMRaycastGpuKernels.hip.cpp + # Single source for both backends: .cu is CUDA by default and + # chrono_set_gpu_source_language() relabels it LANGUAGE HIP when the HIP backend is + # selected, exactly as Chrono::DEM and Chrono::FSI::SPH do. The host bridges below are + # NOT shared -- they are almost entirely runtime API calls and are duplicated per backend. + set(CV_SCM_GPU_KERNEL_FILES + terrain/gpu/SCMGpuKernels.cu + terrain/gpu/SCMRaycastGpuKernels.cu ) - chrono_set_gpu_source_language(${CHRONO_VEHICLE_SCM_BACKEND} ${CV_SCM_GPU_HIP_FILES}) + chrono_set_gpu_source_language(${CHRONO_VEHICLE_SCM_BACKEND} ${CV_SCM_GPU_KERNEL_FILES}) list(APPEND CV_TERRAIN_FILES terrain/SCMGpu.h terrain/SCMGpuTypes.h @@ -333,7 +338,7 @@ if(CH_USE_SCM_GPU) terrain/SCMRaycastGpuTypes.h terrain/SCMTerrainRaycastGpu.cpp terrain/gpu/SCMRaycastGpuHost.cpp - ${CV_SCM_GPU_HIP_FILES} + ${CV_SCM_GPU_KERNEL_FILES} ) endif() diff --git a/src/chrono_vehicle/terrain/gpu/SCMGpuKernels.hip.cpp b/src/chrono_vehicle/terrain/gpu/SCMGpuKernels.cu similarity index 73% rename from src/chrono_vehicle/terrain/gpu/SCMGpuKernels.hip.cpp rename to src/chrono_vehicle/terrain/gpu/SCMGpuKernels.cu index f6f6657492..c62e02cde5 100644 --- a/src/chrono_vehicle/terrain/gpu/SCMGpuKernels.hip.cpp +++ b/src/chrono_vehicle/terrain/gpu/SCMGpuKernels.cu @@ -1,6 +1,35 @@ -// SCMGpuKernels.hip.cpp — HIP kernels for SCM contact forces. - -#include +// SCMGpuKernels.cu — SCM contact-force kernels, compiled as CUDA or HIP. + +// Backend-neutral source: this file is compiled by nvcc for the CUDA backend and by hipcc for the HIP +// one, selected by chrono_set_gpu_source_language() -- the same single-source arrangement Chrono::DEM +// and Chrono::FSI::SPH use. +// +// The conditional below is not optional. nvcc implicitly includes for a .cu, but +// HIP-clang provides nothing: a .cu with no includes compiles under nvcc and fails under +// `hipcc --offload-arch=gfx942` with threadIdx/blockIdx/blockDim undeclared. +// +// gpuStream_t and gpuGetLastError are the ONLY two runtime names this file touches; everything else +// here is device syntax, which CUDA and HIP spell identically. The host bridges, which use around +// thirty runtime symbols, are duplicated per backend instead -- they can be, and Chrono has no +// precedent for a compatibility layer. This file cannot be duplicated: it is the shared source. +// Which branch each build takes, verified rather than assumed: +// CUDA backend nvcc -> CUDA branch (nvcc implicitly includes +// for a .cu) +// HIP on NVIDIA nvcc -D__HIP_PLATFORM_NVIDIA__ -> CUDA branch. CMake's HIP language calls +// nvcc directly here and defines neither +// __HIPCC__ nor __HIP_PLATFORM_AMD__; that +// is correct, because on this platform +// hipStream_t IS cudaStream_t. +// HIP on AMD hipcc / HIP-clang -> HIP branch, the only one that needs the +// header. +#if defined(__HIPCC__) || defined(__HIP_PLATFORM_AMD__) + #include +using gpuStream_t = hipStream_t; + #define gpuGetLastError hipGetLastError +#else +using gpuStream_t = cudaStream_t; + #define gpuGetLastError cudaGetLastError +#endif #include #include @@ -181,7 +210,7 @@ extern "C" int scm_launch_compute_forces(const void* soil_host, const void* in_dev, void* out_dev, int n, - hipStream_t stream) { + gpuStream_t stream) { if (n <= 0) return 0; @@ -194,7 +223,7 @@ extern "C" int scm_launch_compute_forces(const void* soil_host, static_cast(in_dev), static_cast(out_dev), n); - return hipGetLastError(); + return gpuGetLastError(); } extern "C" int scm_launch_reduce_body_forces(const void* in_dev, @@ -202,7 +231,7 @@ extern "C" int scm_launch_reduce_body_forces(const void* in_dev, void* body_forces_dev, int n, int n_bodies, - hipStream_t stream) { + gpuStream_t stream) { if (n <= 0 || n_bodies <= 0) return 0; @@ -213,5 +242,5 @@ extern "C" int scm_launch_reduce_body_forces(const void* in_dev, static_cast(body_forces_dev), n, n_bodies); - return hipGetLastError(); + return gpuGetLastError(); } diff --git a/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.hip.cpp b/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.cu similarity index 85% rename from src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.hip.cpp rename to src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.cu index df0666335f..78c5f36af7 100644 --- a/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.hip.cpp +++ b/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.cu @@ -1,4 +1,4 @@ -// SCMRaycastGpuKernels.hip.cpp — HIP kernel for the SCM GPU ray-cast backend. +// SCMRaycastGpuKernels.cu — ray-cast kernel for the SCM GPU backend, compiled as CUDA or HIP. // // One thread BLOCK per query (SCM grid node ray), not one thread. Threads within a block split the // triangle scan (each thread handles a strided subset), transform each triangle from local to world @@ -24,7 +24,36 @@ // Includes the empirically-determined margin-correction sign; an exact match to Bullet's hit set is // not the goal. -#include +// Backend-neutral source: this file is compiled by nvcc for the CUDA backend and by hipcc for the HIP +// one, selected by chrono_set_gpu_source_language() -- the same single-source arrangement Chrono::DEM +// and Chrono::FSI::SPH use. +// +// The conditional below is not optional. nvcc implicitly includes for a .cu, but +// HIP-clang provides nothing: a .cu with no includes compiles under nvcc and fails under +// `hipcc --offload-arch=gfx942` with threadIdx/blockIdx/blockDim undeclared. +// +// gpuStream_t and gpuGetLastError are the ONLY two runtime names this file touches; everything else +// here is device syntax, which CUDA and HIP spell identically. The host bridges, which use around +// thirty runtime symbols, are duplicated per backend instead -- they can be, and Chrono has no +// precedent for a compatibility layer. This file cannot be duplicated: it is the shared source. +// Which branch each build takes, verified rather than assumed: +// CUDA backend nvcc -> CUDA branch (nvcc implicitly includes +// for a .cu) +// HIP on NVIDIA nvcc -D__HIP_PLATFORM_NVIDIA__ -> CUDA branch. CMake's HIP language calls +// nvcc directly here and defines neither +// __HIPCC__ nor __HIP_PLATFORM_AMD__; that +// is correct, because on this platform +// hipStream_t IS cudaStream_t. +// HIP on AMD hipcc / HIP-clang -> HIP branch, the only one that needs the +// header. +#if defined(__HIPCC__) || defined(__HIP_PLATFORM_AMD__) + #include +using gpuStream_t = hipStream_t; + #define gpuGetLastError hipGetLastError +#else +using gpuStream_t = cudaStream_t; + #define gpuGetLastError cudaGetLastError +#endif #include @@ -309,7 +338,7 @@ int launch(const void* queries_dev, const void* margins_dev, int n_bodies, void* results_dev, - hipStream_t stream) { + gpuStream_t stream) { if (n_queries <= 0) return 0; @@ -322,7 +351,7 @@ int launch(const void* queries_dev, static_cast*>(margins_dev), n_bodies, static_cast*>(results_dev)); - return hipGetLastError(); + return gpuGetLastError(); } } // namespace @@ -336,7 +365,7 @@ extern "C" int scm_launch_raycast_fp64(const void* queries_dev, const void* margins_dev, int n_bodies, void* results_dev, - hipStream_t stream) { + gpuStream_t stream) { return launch(queries_dev, n_queries, verts_dev, faces_dev, n_faces, xforms_dev, margins_dev, n_bodies, results_dev, stream); } @@ -350,7 +379,7 @@ extern "C" int scm_launch_raycast_fp32(const void* queries_dev, const void* margins_dev, int n_bodies, void* results_dev, - hipStream_t stream) { + gpuStream_t stream) { return launch(queries_dev, n_queries, verts_dev, faces_dev, n_faces, xforms_dev, margins_dev, n_bodies, results_dev, stream); } From aaf29b1202866e228411b57efe04fd808ac1b265 Mon Sep 17 00:00:00 2001 From: auc7us Date: Fri, 4 Sep 2026 12:23:37 -0500 Subject: [PATCH 2/5] Add a CUDA host bridge for the SCM GPU backend alongside the HIP one The kernels are already a single source compiled by either toolchain. The host bridges cannot be: they are almost entirely GPU runtime API calls, about thirty symbols of allocation, copies, streams and events. Those are duplicated rather than routed through a compatibility layer, for which Chrono has no precedent. terrain/gpu/hip/SCMGpuHost.cpp moved, unchanged terrain/gpu/hip/SCMRaycastGpuHost.cpp moved, unchanged terrain/gpu/cuda/SCMGpuHost.cpp new terrain/gpu/cuda/SCMRaycastGpuHost.cpp new Exactly one directory is compiled per build, selected from CHRONO_VEHICLE_SCM_BACKEND. Each file carries a banner naming its counterpart, because nothing enforces that the two stay in step -- only one is ever compiled, so a change made to one alone produces no error anywhere. The translation is not name-for-name everywhere. hipHostMalloc(p, n) is cudaHostAlloc(p, n, flags): the flags argument that HIP defaults is mandatory in CUDA, so a mechanical rename produces a call that fails to compile. Three sites in the contact-force bridge are affected, all in the pinned staging buffers of the async double-buffered pipeline. Also removes a stale #include from SCMTerrainRaycastGpu.cpp, which uses no HIP symbol at all. It went unnoticed while the only backend was HIP. CH_ENABLE_VEHICLE_SCM_GPU now REQUIRES CUDA_OR_HIP, which gives the intended default on each vendor with no further configuration: an NVIDIA machine resolves AUTO to CUDA (both backends are candidates and CUDA is first), and an AMD machine resolves it to HIP on ROCm (the only candidate). HIP over CUDA remains available on NVIDIA for anyone who wants it, via CHRONO_VEHICLE_SCM_GPU_BACKEND=HIP. The one consequence to be aware of is that an existing NVIDIA build directory that was using HIP moves to CUDA on reconfigure, which is the point. Verified both ways on one machine (RTX 4080, CUDA 13.2, ROCm 7.2 with HIP platform nvidia): AUTO -> CUDA CHRONO_VEHICLE_SCM_GPU_BACKEND=HIP gpu/hip bridges, HIP_COMPILER, links HIP CHRONO_VEHICLE_SCM_GPU_BACKEND=CUDA gpu/cuda bridges, CUDA_COMPILER, links CXX Both produce a Chrono_vehicle carrying the SCM GPU entry points. The AMD leg (gfx942) is unbuilt here: this container has no ROCm device bitcode, only the front end. --- src/chrono_vehicle/CMakeLists.txt | 26 +- .../terrain/SCMTerrainRaycastGpu.cpp | 2 - .../terrain/gpu/cuda/SCMGpuHost.cpp | 471 ++++++++++++++++++ .../terrain/gpu/cuda/SCMRaycastGpuHost.cpp | 357 +++++++++++++ .../terrain/gpu/{ => hip}/SCMGpuHost.cpp | 8 + .../gpu/{ => hip}/SCMRaycastGpuHost.cpp | 10 +- 6 files changed, 864 insertions(+), 10 deletions(-) create mode 100644 src/chrono_vehicle/terrain/gpu/cuda/SCMGpuHost.cpp create mode 100644 src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp rename src/chrono_vehicle/terrain/gpu/{ => hip}/SCMGpuHost.cpp (96%) rename src/chrono_vehicle/terrain/gpu/{ => hip}/SCMRaycastGpuHost.cpp (95%) diff --git a/src/chrono_vehicle/CMakeLists.txt b/src/chrono_vehicle/CMakeLists.txt index 92e9601fb2..95aee0fd7f 100644 --- a/src/chrono_vehicle/CMakeLists.txt +++ b/src/chrono_vehicle/CMakeLists.txt @@ -38,9 +38,11 @@ set(CH_USE_SCM_GPU OFF) # The GPU requirement is declared for this FEATURE, not for Chrono::Vehicle: # the module builds fine with no GPU backend at all, and only the optional SCM -# GPU path needs one. The kernels themselves are backend-neutral (single-source .cu, -# compiled as CUDA or HIP), but the host bridges are still HIP-only, hence REQUIRES HIP -# for now; this becomes CUDA_OR_HIP once the CUDA host bridges land. Both HIP platforms are accepted: nothing here depends on +# GPU path needs one. The kernels are backend-neutral (single-source .cu compiled as +# CUDA or HIP) and the host bridges are duplicated per backend under terrain/gpu/cuda +# and terrain/gpu/hip, so either toolchain is sufficient. Both HIP platforms remain +# acceptable: nothing here uses the ROCm-only libraries that restrict Chrono::DEM and +# Chrono::FSI::SPH to the amd platform. Both HIP platforms are accepted: nothing here depends on # the ROCm-only libraries (hipCUB/rocPRIM, rocThrust) that restrict # Chrono::DEM and Chrono::FSI::SPH to the amd platform. So "HIP + ROCm on AMD" # and "HIP + nvcc on NVIDIA" both give the GPU path with no further @@ -50,7 +52,7 @@ if(CH_ENABLE_VEHICLE_SCM_GPU) chrono_select_gpu_backend(CHRONO_VEHICLE_SCM_BACKEND FEATURE "Chrono::Vehicle SCM GPU" NAME VEHICLE_SCM - REQUIRES HIP + REQUIRES CUDA_OR_HIP HIP_PLATFORMS amd nvidia) if(CHRONO_VEHICLE_SCM_BACKEND STREQUAL "NONE") @@ -66,7 +68,7 @@ if(CH_ENABLE_VEHICLE_SCM_GPU) # silently build the CPU path. Leaving the cache alone lets availability be # re-evaluated every configure. chrono_gpu_feature_unavailable(FEATURE "Chrono::Vehicle SCM GPU" - REQUIRES "HIP" + REQUIRES "CUDA or HIP" CLASS IMPLIED) else() @@ -328,16 +330,26 @@ if(CH_USE_SCM_GPU) terrain/gpu/SCMRaycastGpuKernels.cu ) chrono_set_gpu_source_language(${CHRONO_VEHICLE_SCM_BACKEND} ${CV_SCM_GPU_KERNEL_FILES}) + # Host bridges are duplicated, not shared: they are almost entirely runtime API calls + # (about thirty symbols), and Chrono has no compatibility layer for those. Exactly one + # directory is compiled per build; the two files in each pair must be kept in step by + # hand, which the banner at the top of each states. + if(CHRONO_VEHICLE_SCM_BACKEND STREQUAL "CUDA") + set(CV_SCM_GPU_HOST_DIR terrain/gpu/cuda) + else() + set(CV_SCM_GPU_HOST_DIR terrain/gpu/hip) + endif() + list(APPEND CV_TERRAIN_FILES terrain/SCMGpu.h terrain/SCMGpuTypes.h terrain/SCMTerrainGpu.h terrain/SCMTerrainGpu.cpp - terrain/gpu/SCMGpuHost.cpp + ${CV_SCM_GPU_HOST_DIR}/SCMGpuHost.cpp terrain/SCMRaycastGpu.h terrain/SCMRaycastGpuTypes.h terrain/SCMTerrainRaycastGpu.cpp - terrain/gpu/SCMRaycastGpuHost.cpp + ${CV_SCM_GPU_HOST_DIR}/SCMRaycastGpuHost.cpp ${CV_SCM_GPU_KERNEL_FILES} ) endif() diff --git a/src/chrono_vehicle/terrain/SCMTerrainRaycastGpu.cpp b/src/chrono_vehicle/terrain/SCMTerrainRaycastGpu.cpp index 0781aac965..90e0ce256e 100644 --- a/src/chrono_vehicle/terrain/SCMTerrainRaycastGpu.cpp +++ b/src/chrono_vehicle/terrain/SCMTerrainRaycastGpu.cpp @@ -19,8 +19,6 @@ #include #include - #include - #include "chrono/physics/ChBody.h" #include "chrono/collision/ChCollisionShapeTriangleMesh.h" #include "chrono/geometry/ChTriangleMeshConnected.h" diff --git a/src/chrono_vehicle/terrain/gpu/cuda/SCMGpuHost.cpp b/src/chrono_vehicle/terrain/gpu/cuda/SCMGpuHost.cpp new file mode 100644 index 0000000000..27d2d68d03 --- /dev/null +++ b/src/chrono_vehicle/terrain/gpu/cuda/SCMGpuHost.cpp @@ -0,0 +1,471 @@ +// SCMGpuHost.cpp — CUDA host bridge: pinned staging, async copy/compute streams, body reduce. +// NOTE: this file is DUPLICATED per GPU backend. Its counterpart is +// terrain/gpu/hip/SCMGpuHost.cpp, and the two differ only in the runtime API names +// (about thirty symbols). Any change here must be made there too -- nothing enforces it, +// because only one of the two is compiled in a given build. Chrono has no compatibility +// layer for this and deliberately does not grow one; the kernels, which CAN be shared, +// are single-source .cu files instead. +// +// Not a name-for-name mapping: hipHostMalloc(p, n) is cudaHostAlloc(p, n, flags). + +#include "chrono_vehicle/terrain/SCMGpu.h" + +#include + +#include +#include +#include +#include + +namespace { + +using chrono::vehicle::scm::gpu::BodyForceAccum; +using chrono::vehicle::scm::gpu::HitInput; +using chrono::vehicle::scm::gpu::HitOutput; +using chrono::vehicle::scm::gpu::SoilParams; + +extern "C" int scm_launch_compute_forces(const void* soil_host, + const void* in_dev, + void* out_dev, + int n, + cudaStream_t stream); + +extern "C" int scm_launch_reduce_body_forces(const void* in_dev, + const void* out_dev, + void* body_forces_dev, + int n, + int n_bodies, + cudaStream_t stream); + +struct BufferSlot { + HitInput* h_in = nullptr; + HitOutput* h_out = nullptr; + HitInput* d_in = nullptr; + HitOutput* d_out = nullptr; +}; + +struct ScmGpuContextImpl { + int device = 0; + cudaStream_t stream_copy = nullptr; + cudaStream_t stream_compute = nullptr; + cudaEvent_t event_h2d_done = nullptr; + cudaEvent_t event_compute_done = nullptr; + + BufferSlot slot; + bool in_flight = false; + + BodyForceAccum* h_body = nullptr; + double* d_body = nullptr; + std::size_t hit_capacity = 0; + std::size_t body_capacity = 0; + bool warmed_up = false; +}; + +chrono::vehicle::scm_gpu::Config& MutableConfig() { + static chrono::vehicle::scm_gpu::Config cfg; + return cfg; +} + +void die_cuda(const char* msg, cudaError_t err) { + fprintf(stderr, "SCM GPU FATAL: %s — %s\n", msg, cudaGetErrorString(err)); + std::abort(); +} + +void ensure_device(int device) { + int current = -1; + cudaGetDevice(¤t); + if (current != device) + cudaSetDevice(device); +} + +void free_slot(BufferSlot& slot) { + if (slot.d_in) + (void)cudaFree(slot.d_in); + if (slot.d_out) + (void)cudaFree(slot.d_out); + if (slot.h_in) + (void)cudaFreeHost(slot.h_in); + if (slot.h_out) + (void)cudaFreeHost(slot.h_out); + slot = {}; +} + +void ensure_hit_capacity(ScmGpuContextImpl* ctx, std::size_t n) { + if (n <= ctx->hit_capacity) + return; + + free_slot(ctx->slot); + + const std::size_t bytes_in = n * sizeof(HitInput); + const std::size_t bytes_out = n * sizeof(HitOutput); + + cudaError_t e1 = cudaMalloc(&ctx->slot.d_in, bytes_in); + if (e1 != cudaSuccess) + die_cuda("cudaMalloc d_in", e1); + cudaError_t e2 = cudaMalloc(&ctx->slot.d_out, bytes_out); + if (e2 != cudaSuccess) + die_cuda("cudaMalloc d_out", e2); + cudaError_t e3 = cudaHostAlloc(&ctx->slot.h_in, bytes_in, cudaHostAllocDefault); + if (e3 != cudaSuccess) + die_cuda("cudaHostAlloc h_in", e3); + cudaError_t e4 = cudaHostAlloc(&ctx->slot.h_out, bytes_out, cudaHostAllocDefault); + if (e4 != cudaSuccess) + die_cuda("cudaHostAlloc h_out", e4); + + ctx->hit_capacity = n; +} + +void ensure_body_capacity(ScmGpuContextImpl* ctx, std::size_t n_bodies) { + if (n_bodies <= ctx->body_capacity) + return; + + if (ctx->d_body) + (void)cudaFree(ctx->d_body); + if (ctx->h_body) + (void)cudaFreeHost(ctx->h_body); + + const std::size_t bytes = n_bodies * 6 * sizeof(double); + cudaError_t e1 = cudaMalloc(&ctx->d_body, bytes); + if (e1 != cudaSuccess) + die_cuda("cudaMalloc d_body", e1); + cudaError_t e2 = cudaHostAlloc(reinterpret_cast(&ctx->h_body), bytes, cudaHostAllocDefault); + if (e2 != cudaSuccess) + die_cuda("cudaHostAlloc h_body", e2); + + ctx->body_capacity = n_bodies; +} + +BufferSlot& current_slot(ScmGpuContextImpl* impl) { + return impl->slot; +} + +void sync_impl(ScmGpuContextImpl* impl) { + if (!impl->in_flight) + return; + cudaError_t e1 = cudaStreamSynchronize(impl->stream_copy); + if (e1 != cudaSuccess) + die_cuda("cudaStreamSynchronize copy", e1); + cudaError_t e2 = cudaStreamSynchronize(impl->stream_compute); + if (e2 != cudaSuccess) + die_cuda("cudaStreamSynchronize compute", e2); + impl->in_flight = false; +} + +int launch_pipelined(ScmGpuContextImpl* impl, + const SoilParams& soil, + std::size_t n_hits, + std::size_t n_bodies) { + BufferSlot& slot = current_slot(impl); + const std::size_t bytes_in = n_hits * sizeof(HitInput); + const std::size_t bytes_out = n_hits * sizeof(HitOutput); + const bool reduce_bodies = n_bodies > 0; + + if (reduce_bodies) + ensure_body_capacity(impl, n_bodies); + + sync_impl(impl); + + cudaError_t e_h2d = + cudaMemcpyAsync(slot.d_in, slot.h_in, bytes_in, cudaMemcpyHostToDevice, impl->stream_copy); + if (e_h2d != cudaSuccess) + return static_cast(e_h2d); + + cudaError_t e_rec_h2d = cudaEventRecord(impl->event_h2d_done, impl->stream_copy); + if (e_rec_h2d != cudaSuccess) + return static_cast(e_rec_h2d); + + cudaError_t e_wait_h2d = cudaStreamWaitEvent(impl->stream_compute, impl->event_h2d_done, 0); + if (e_wait_h2d != cudaSuccess) + return static_cast(e_wait_h2d); + + const int launch_err = scm_launch_compute_forces(&soil, + slot.d_in, + slot.d_out, + static_cast(n_hits), + impl->stream_compute); + if (launch_err != cudaSuccess) + return launch_err; + + if (reduce_bodies) { + const std::size_t body_bytes = n_bodies * 6 * sizeof(double); + cudaError_t e_zero = + cudaMemsetAsync(impl->d_body, 0, body_bytes, impl->stream_compute); + if (e_zero != cudaSuccess) + return static_cast(e_zero); + + const int reduce_err = scm_launch_reduce_body_forces(slot.d_in, + slot.d_out, + impl->d_body, + static_cast(n_hits), + static_cast(n_bodies), + impl->stream_compute); + if (reduce_err != cudaSuccess) + return reduce_err; + } + + cudaError_t e_rec_compute = cudaEventRecord(impl->event_compute_done, impl->stream_compute); + if (e_rec_compute != cudaSuccess) + return static_cast(e_rec_compute); + + cudaError_t e_wait_compute = cudaStreamWaitEvent(impl->stream_copy, impl->event_compute_done, 0); + if (e_wait_compute != cudaSuccess) + return static_cast(e_wait_compute); + + cudaError_t e_d2h_out = + cudaMemcpyAsync(slot.h_out, slot.d_out, bytes_out, cudaMemcpyDeviceToHost, impl->stream_copy); + if (e_d2h_out != cudaSuccess) + return static_cast(e_d2h_out); + + if (reduce_bodies) { + const std::size_t body_bytes = n_bodies * 6 * sizeof(double); + cudaError_t e_d2h_body = cudaMemcpyAsync(impl->h_body, + impl->d_body, + body_bytes, + cudaMemcpyDeviceToHost, + impl->stream_copy); + if (e_d2h_body != cudaSuccess) + return static_cast(e_d2h_body); + } + + impl->in_flight = true; + sync_impl(impl); + return static_cast(cudaSuccess); +} + +int launch_simple(ScmGpuContextImpl* impl, const SoilParams& soil, std::size_t n_hits, std::size_t n_bodies) { + BufferSlot& slot = current_slot(impl); + const std::size_t bytes_in = n_hits * sizeof(HitInput); + const std::size_t bytes_out = n_hits * sizeof(HitOutput); + const bool reduce_bodies = n_bodies > 0; + + if (reduce_bodies) + ensure_body_capacity(impl, n_bodies); + + cudaError_t e1 = cudaMemcpy(slot.d_in, slot.h_in, bytes_in, cudaMemcpyHostToDevice); + if (e1 != cudaSuccess) + return static_cast(e1); + + const int launch_err = scm_launch_compute_forces(&soil, + slot.d_in, + slot.d_out, + static_cast(n_hits), + impl->stream_compute); + if (launch_err != cudaSuccess) + return launch_err; + + if (reduce_bodies) { + const std::size_t body_bytes = n_bodies * 6 * sizeof(double); + cudaError_t e_zero = cudaMemset(impl->d_body, 0, body_bytes); + if (e_zero != cudaSuccess) + return static_cast(e_zero); + + const int reduce_err = scm_launch_reduce_body_forces(slot.d_in, + slot.d_out, + impl->d_body, + static_cast(n_hits), + static_cast(n_bodies), + impl->stream_compute); + if (reduce_err != cudaSuccess) + return reduce_err; + + cudaError_t e_body = cudaMemcpy(impl->h_body, impl->d_body, body_bytes, cudaMemcpyDeviceToHost); + if (e_body != cudaSuccess) + return static_cast(e_body); + } + + cudaError_t e2 = cudaMemcpy(slot.h_out, slot.d_out, bytes_out, cudaMemcpyDeviceToHost); + if (e2 != cudaSuccess) + return static_cast(e2); + + return static_cast(cudaSuccess); +} + +} // namespace + +namespace chrono { +namespace vehicle { +namespace scm_gpu { + +void SetConfig(const Config& config) { + MutableConfig() = config; +} + +Config GetConfig() { + return MutableConfig(); +} + +} // namespace scm_gpu +} // namespace vehicle +} // namespace chrono + +extern "C" std::size_t scm_gpu_min_hits(void) { + return chrono::vehicle::scm_gpu::GetConfig().min_hits; +} + +extern "C" std::size_t scm_gpu_reserve_hits(void) { + return chrono::vehicle::scm_gpu::GetConfig().reserve_hits; +} + +extern "C" int scm_gpu_async_enabled(void) { + return chrono::vehicle::scm_gpu::GetConfig().async ? 1 : 0; +} + +extern "C" ScmGpuContext* scm_gpu_create(int device_id) { + auto* impl = new ScmGpuContextImpl(); + impl->device = device_id; + ensure_device(device_id); + + cudaError_t e1 = cudaStreamCreateWithFlags(&impl->stream_copy, cudaStreamNonBlocking); + if (e1 != cudaSuccess) + die_cuda("cudaStreamCreate copy", e1); + cudaError_t e2 = cudaStreamCreateWithFlags(&impl->stream_compute, cudaStreamNonBlocking); + if (e2 != cudaSuccess) + die_cuda("cudaStreamCreate compute", e2); + cudaError_t e3 = cudaEventCreateWithFlags(&impl->event_h2d_done, cudaEventDisableTiming); + if (e3 != cudaSuccess) + die_cuda("cudaEventCreate h2d", e3); + cudaError_t e4 = cudaEventCreateWithFlags(&impl->event_compute_done, cudaEventDisableTiming); + if (e4 != cudaSuccess) + die_cuda("cudaEventCreate compute", e4); + + return reinterpret_cast(impl); +} + +extern "C" void scm_gpu_destroy(ScmGpuContext* ctx) { + if (!ctx) + return; + auto* impl = reinterpret_cast(ctx); + ensure_device(impl->device); + sync_impl(impl); + free_slot(impl->slot); + if (impl->d_body) + cudaFree(impl->d_body); + if (impl->h_body) + cudaFreeHost(impl->h_body); + if (impl->event_h2d_done) + cudaEventDestroy(impl->event_h2d_done); + if (impl->event_compute_done) + cudaEventDestroy(impl->event_compute_done); + if (impl->stream_copy) + cudaStreamDestroy(impl->stream_copy); + if (impl->stream_compute) + cudaStreamDestroy(impl->stream_compute); + delete impl; +} + +extern "C" void scm_gpu_reserve(ScmGpuContext* ctx, std::size_t n_hits) { + if (!ctx || n_hits == 0) + return; + auto* impl = reinterpret_cast(ctx); + ensure_device(impl->device); + ensure_hit_capacity(impl, n_hits); +} + +extern "C" void scm_gpu_warmup(ScmGpuContext* ctx) { + if (!ctx) + return; + auto* impl = reinterpret_cast(ctx); + if (impl->warmed_up) + return; + ensure_device(impl->device); + const std::size_t reserve_n = std::max(scm_gpu_reserve_hits(), 1); + ensure_hit_capacity(impl, reserve_n); + ensure_body_capacity(impl, 1); + current_slot(impl).h_in[0] = {}; + current_slot(impl).h_in[0].active = 1; + SoilParams soil{}; + soil.elastic_k = 1.0; + soil.area = 1.0; + soil.dt = 1e-3; + impl->warmed_up = true; + if (scm_gpu_async_enabled()) + (void)launch_pipelined(impl, soil, 1, 1); + else + (void)launch_simple(impl, soil, 1, 1); +} + +extern "C" HitInput* scm_gpu_prepare_input(ScmGpuContext* ctx, std::size_t n_hits) { + if (!ctx || n_hits == 0) + return nullptr; + auto* impl = reinterpret_cast(ctx); + ensure_device(impl->device); + ensure_hit_capacity(impl, n_hits); + return current_slot(impl).h_in; +} + +extern "C" HitOutput* scm_gpu_prepare_output(ScmGpuContext* ctx, std::size_t n_hits) { + if (!ctx || n_hits == 0) + return nullptr; + auto* impl = reinterpret_cast(ctx); + ensure_device(impl->device); + ensure_hit_capacity(impl, n_hits); + return current_slot(impl).h_out; +} + +extern "C" BodyForceAccum* scm_gpu_prepare_body_forces(ScmGpuContext* ctx, std::size_t n_bodies) { + if (!ctx || n_bodies == 0) + return nullptr; + auto* impl = reinterpret_cast(ctx); + ensure_device(impl->device); + ensure_body_capacity(impl, n_bodies); + return impl->h_body; +} + +extern "C" void scm_gpu_sync(ScmGpuContext* ctx) { + if (!ctx) + return; + auto* impl = reinterpret_cast(ctx); + ensure_device(impl->device); + sync_impl(impl); +} + +extern "C" int scm_gpu_compute_forces_staged(ScmGpuContext* ctx, + const SoilParams& soil, + std::size_t n_hits, + std::size_t n_bodies) { + if (!ctx) + return -1; + if (n_hits == 0) + return 0; + + auto* impl = reinterpret_cast(ctx); + ensure_device(impl->device); + if (!impl->warmed_up) + scm_gpu_warmup(ctx); + ensure_hit_capacity(impl, n_hits); + + if (scm_gpu_async_enabled()) + return launch_pipelined(impl, soil, n_hits, n_bodies); + return launch_simple(impl, soil, n_hits, n_bodies); +} + +extern "C" int scm_gpu_compute_forces(ScmGpuContext* ctx, + const SoilParams& soil, + const HitInput* in, + HitOutput* out_host, + std::size_t n_hits, + std::size_t n_bodies) { + if (!ctx || !in || !out_host) + return -1; + if (n_hits == 0) + return 0; + + auto* impl = reinterpret_cast(ctx); + ensure_device(impl->device); + if (!impl->warmed_up) + scm_gpu_warmup(ctx); + ensure_hit_capacity(impl, n_hits); + + BufferSlot& slot = current_slot(impl); + const std::size_t bytes_in = n_hits * sizeof(HitInput); + const std::size_t bytes_out = n_hits * sizeof(HitOutput); + if (in != slot.h_in) + std::memcpy(slot.h_in, in, bytes_in); + + const int err = scm_gpu_compute_forces_staged(ctx, soil, n_hits, n_bodies); + if (err != cudaSuccess) + return err; + if (out_host != slot.h_out) + std::memcpy(out_host, slot.h_out, bytes_out); + return 0; +} diff --git a/src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp b/src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp new file mode 100644 index 0000000000..979a1c20ab --- /dev/null +++ b/src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp @@ -0,0 +1,357 @@ +// SCMRaycastGpuHost.cpp — CUDA host bridge for the SCM ray-cast backend: device buffer management, +// NOTE: this file is DUPLICATED per GPU backend. Its counterpart is +// terrain/gpu/hip/SCMRaycastGpuHost.cpp, and the two differ only in the runtime API names +// (about thirty symbols). Any change here must be made there too -- nothing enforces it, +// because only one of the two is compiled in a given build. Chrono has no compatibility +// layer for this and deliberately does not grow one; the kernels, which CAN be shared, +// are single-source .cu files instead. +// +// Not a name-for-name mapping: hipHostMalloc(p, n) is cudaHostAlloc(p, n, flags). +// synchronous upload/run (v1 -- see SCMRaycastGpu.h for why this isn't pipelined yet). +// +// Supports two kernel precisions (ScmRaycastGpuPrecision): FP64, the validated default, and FP32, added for GPUs with weak double-precision throughput -- +// notably consumer NVIDIA cards (e.g. RTX 4080/5090), unlike this project's AMD MI300X target, a proper +// datacenter part with strong FP64. The public API types (SCMRaycastGpuTypes.h) stay double-precision +// throughout -- Chrono itself is double internally -- this file downcasts to float on upload and +// upconverts results back to double when precision == kFP32, so callers (SCMTerrainRaycastGpu.cpp) +// don't need to know or care which precision is active. +// +// Mesh geometry and per-body transforms are uploaded through SEPARATE calls: mesh (vertices/faces/ +// margins) only changes when the candidate body set changes (rare -- never, in the common +// single-vehicle case), while transforms change every step. The caller (SCMTerrainRaycastGpu.cpp) +// caches the candidate set and only calls scm_raycast_gpu_upload_mesh when it actually changes. + +#include "chrono_vehicle/terrain/SCMRaycastGpu.h" + +#include + +#include +#include +#include +#include + +namespace { + +using chrono::vehicle::scm::gpu::RaycastBodyMargin; +using chrono::vehicle::scm::gpu::RaycastBodyTransform; +using chrono::vehicle::scm::gpu::RaycastFace; +using chrono::vehicle::scm::gpu::RaycastQuery; +using chrono::vehicle::scm::gpu::RaycastResult; +using chrono::vehicle::scm::gpu::RaycastVertex; + +// Float mirrors of the double-precision public types. Only the memory layout needs to match the +// SCMRaycastGpuKernels.cu kernel's own VertexDevT/TransformDevT/MarginDevT/QueryDevT/ +// ResultDevT -- not the C++ type identity -- since the two are compiled by different compilers +// (g++ here, nvcc there) and only ever communicate via raw device pointers. +struct VertexF { + float x, y, z; +}; +struct TransformF { + float px, py, pz; + float r00, r01, r02; + float r10, r11, r12; + float r20, r21, r22; + float bx0, by0, bz0; + float bx1, by1, bz1; +}; +struct MarginF { + float margin; + int32_t face_begin; + int32_t face_end; +}; +struct QueryF { + float from_x, from_y, from_z; + float to_x, to_y, to_z; +}; +struct ResultF { + int32_t hit; + int32_t body_slot; + float hit_x, hit_y, hit_z; +}; + +extern "C" int scm_launch_raycast_fp64(const void* queries_dev, + int n_queries, + const void* verts_dev, + const void* faces_dev, + int n_faces, + const void* xforms_dev, + const void* margins_dev, + int n_bodies, + void* results_dev, + cudaStream_t stream); +extern "C" int scm_launch_raycast_fp32(const void* queries_dev, + int n_queries, + const void* verts_dev, + const void* faces_dev, + int n_faces, + const void* xforms_dev, + const void* margins_dev, + int n_bodies, + void* results_dev, + cudaStream_t stream); + +void die_cuda(const char* msg, cudaError_t err) { + fprintf(stderr, "SCM RAYCAST GPU FATAL: %s -- %s\n", msg, cudaGetErrorString(err)); + std::abort(); +} + +void ensure_device(int device) { + int current = -1; + cudaGetDevice(¤t); + if (current != device) + cudaSetDevice(device); +} + +template +void ensure_capacity(T** d_ptr, std::size_t& capacity, std::size_t n) { + if (n <= capacity) + return; + if (*d_ptr) + (void)cudaFree(*d_ptr); + cudaError_t e = cudaMalloc(d_ptr, n * sizeof(T)); + if (e != cudaSuccess) + die_cuda("cudaMalloc", e); + capacity = n; +} + +} // namespace + +struct ScmRaycastGpuContext { + int device = 0; + ScmRaycastGpuPrecision precision = ScmRaycastGpuPrecision::kFP32; + + // FP64 device buffers (used when precision == kFP64). + RaycastVertex* d_verts = nullptr; + RaycastBodyTransform* d_xforms = nullptr; + RaycastBodyMargin* d_margins = nullptr; + RaycastQuery* d_queries = nullptr; + RaycastResult* d_results = nullptr; + std::size_t cap_verts = 0; + std::size_t cap_bodies = 0; + std::size_t cap_margins = 0; + std::size_t cap_queries = 0; + std::size_t cap_results = 0; + + // FP32 device buffers (used when precision == kFP32). + VertexF* d_verts_f = nullptr; + TransformF* d_xforms_f = nullptr; + MarginF* d_margins_f = nullptr; + QueryF* d_queries_f = nullptr; + ResultF* d_results_f = nullptr; + std::size_t cap_verts_f = 0; + std::size_t cap_bodies_f = 0; + std::size_t cap_margins_f = 0; + std::size_t cap_queries_f = 0; + std::size_t cap_results_f = 0; + + // Faces (indices + body_slot) are precision-independent -- always int32_t, one shared buffer. + RaycastFace* d_faces = nullptr; + std::size_t cap_faces = 0; + + int n_faces_current = 0; + int n_bodies_current = 0; +}; + +extern "C" ScmRaycastGpuContext* scm_raycast_gpu_create(int device_id, ScmRaycastGpuPrecision precision) { + auto* ctx = new ScmRaycastGpuContext(); + ctx->device = device_id; + ctx->precision = precision; + ensure_device(device_id); + return ctx; +} + +extern "C" void scm_raycast_gpu_destroy(ScmRaycastGpuContext* ctx) { + if (!ctx) + return; + ensure_device(ctx->device); + if (ctx->d_verts) + cudaFree(ctx->d_verts); + if (ctx->d_faces) + cudaFree(ctx->d_faces); + if (ctx->d_xforms) + cudaFree(ctx->d_xforms); + if (ctx->d_margins) + cudaFree(ctx->d_margins); + if (ctx->d_queries) + cudaFree(ctx->d_queries); + if (ctx->d_results) + cudaFree(ctx->d_results); + if (ctx->d_verts_f) + cudaFree(ctx->d_verts_f); + if (ctx->d_xforms_f) + cudaFree(ctx->d_xforms_f); + if (ctx->d_margins_f) + cudaFree(ctx->d_margins_f); + if (ctx->d_queries_f) + cudaFree(ctx->d_queries_f); + if (ctx->d_results_f) + cudaFree(ctx->d_results_f); + delete ctx; +} + +extern "C" int scm_raycast_gpu_upload_mesh(ScmRaycastGpuContext* ctx, + const RaycastVertex* verts, + int n_verts, + const RaycastFace* faces, + int n_faces, + const RaycastBodyMargin* margins, + int n_bodies) { + if (!ctx) + return -1; + ensure_device(ctx->device); + + if (n_verts > 0) { + if (ctx->precision == ScmRaycastGpuPrecision::kFP32) { + std::vector tmp(n_verts); + for (int i = 0; i < n_verts; ++i) + tmp[i] = {static_cast(verts[i].x), static_cast(verts[i].y), + static_cast(verts[i].z)}; + ensure_capacity(&ctx->d_verts_f, ctx->cap_verts_f, static_cast(n_verts)); + cudaError_t e = cudaMemcpy(ctx->d_verts_f, tmp.data(), n_verts * sizeof(VertexF), cudaMemcpyHostToDevice); + if (e != cudaSuccess) + return static_cast(e); + } else { + ensure_capacity(&ctx->d_verts, ctx->cap_verts, static_cast(n_verts)); + cudaError_t e = cudaMemcpy(ctx->d_verts, verts, n_verts * sizeof(RaycastVertex), cudaMemcpyHostToDevice); + if (e != cudaSuccess) + return static_cast(e); + } + } + if (n_faces > 0) { + ensure_capacity(&ctx->d_faces, ctx->cap_faces, static_cast(n_faces)); + cudaError_t e = cudaMemcpy(ctx->d_faces, faces, n_faces * sizeof(RaycastFace), cudaMemcpyHostToDevice); + if (e != cudaSuccess) + return static_cast(e); + } + if (n_bodies > 0) { + if (ctx->precision == ScmRaycastGpuPrecision::kFP32) { + std::vector tmp(n_bodies); + for (int i = 0; i < n_bodies; ++i) + tmp[i] = {static_cast(margins[i].margin), margins[i].face_begin, margins[i].face_end}; + ensure_capacity(&ctx->d_margins_f, ctx->cap_margins_f, static_cast(n_bodies)); + cudaError_t e = + cudaMemcpy(ctx->d_margins_f, tmp.data(), n_bodies * sizeof(MarginF), cudaMemcpyHostToDevice); + if (e != cudaSuccess) + return static_cast(e); + } else { + ensure_capacity(&ctx->d_margins, ctx->cap_margins, static_cast(n_bodies)); + cudaError_t e = + cudaMemcpy(ctx->d_margins, margins, n_bodies * sizeof(RaycastBodyMargin), cudaMemcpyHostToDevice); + if (e != cudaSuccess) + return static_cast(e); + } + } + + ctx->n_faces_current = n_faces; + ctx->n_bodies_current = n_bodies; + return static_cast(cudaSuccess); +} + +extern "C" int scm_raycast_gpu_upload_transforms(ScmRaycastGpuContext* ctx, + const RaycastBodyTransform* xforms, + int n_bodies) { + if (!ctx) + return -1; + if (n_bodies <= 0) + return 0; + ensure_device(ctx->device); + + if (ctx->precision == ScmRaycastGpuPrecision::kFP32) { + std::vector tmp(n_bodies); + for (int i = 0; i < n_bodies; ++i) { + const RaycastBodyTransform& t = xforms[i]; + tmp[i] = {static_cast(t.px), static_cast(t.py), static_cast(t.pz), + static_cast(t.r00), static_cast(t.r01), static_cast(t.r02), + static_cast(t.r10), static_cast(t.r11), static_cast(t.r12), + static_cast(t.r20), static_cast(t.r21), static_cast(t.r22), + static_cast(t.bx0), static_cast(t.by0), static_cast(t.bz0), + static_cast(t.bx1), static_cast(t.by1), static_cast(t.bz1)}; + } + ensure_capacity(&ctx->d_xforms_f, ctx->cap_bodies_f, static_cast(n_bodies)); + cudaError_t e = cudaMemcpy(ctx->d_xforms_f, tmp.data(), n_bodies * sizeof(TransformF), cudaMemcpyHostToDevice); + if (e != cudaSuccess) + return static_cast(e); + } else { + ensure_capacity(&ctx->d_xforms, ctx->cap_bodies, static_cast(n_bodies)); + cudaError_t e = + cudaMemcpy(ctx->d_xforms, xforms, n_bodies * sizeof(RaycastBodyTransform), cudaMemcpyHostToDevice); + if (e != cudaSuccess) + return static_cast(e); + } + + return static_cast(cudaSuccess); +} + +extern "C" int scm_raycast_gpu_run(ScmRaycastGpuContext* ctx, + const RaycastQuery* queries, + RaycastResult* out_results, + int n_queries) { + if (!ctx || !queries || !out_results) + return -1; + if (n_queries <= 0) + return 0; + + ensure_device(ctx->device); + + if (ctx->precision == ScmRaycastGpuPrecision::kFP32) { + std::vector q_tmp(n_queries); + for (int i = 0; i < n_queries; ++i) + q_tmp[i] = {static_cast(queries[i].from_x), static_cast(queries[i].from_y), + static_cast(queries[i].from_z), static_cast(queries[i].to_x), + static_cast(queries[i].to_y), static_cast(queries[i].to_z)}; + + ensure_capacity(&ctx->d_queries_f, ctx->cap_queries_f, static_cast(n_queries)); + ensure_capacity(&ctx->d_results_f, ctx->cap_results_f, static_cast(n_queries)); + + cudaError_t e1 = cudaMemcpy(ctx->d_queries_f, q_tmp.data(), n_queries * sizeof(QueryF), cudaMemcpyHostToDevice); + if (e1 != cudaSuccess) + return static_cast(e1); + + int launch_err = scm_launch_raycast_fp32(ctx->d_queries_f, n_queries, ctx->d_verts_f, ctx->d_faces, + ctx->n_faces_current, ctx->d_xforms_f, ctx->d_margins_f, + ctx->n_bodies_current, ctx->d_results_f, nullptr); + if (launch_err != cudaSuccess) + return launch_err; + + cudaError_t e2 = cudaDeviceSynchronize(); + if (e2 != cudaSuccess) + return static_cast(e2); + std::vector r_tmp(n_queries); + cudaError_t e3 = + cudaMemcpy(r_tmp.data(), ctx->d_results_f, n_queries * sizeof(ResultF), cudaMemcpyDeviceToHost); + if (e3 != cudaSuccess) + return static_cast(e3); + + for (int i = 0; i < n_queries; ++i) { + out_results[i].hit = r_tmp[i].hit; + out_results[i].body_slot = r_tmp[i].body_slot; + out_results[i].hit_x = r_tmp[i].hit_x; + out_results[i].hit_y = r_tmp[i].hit_y; + out_results[i].hit_z = r_tmp[i].hit_z; + } + return static_cast(cudaSuccess); + } + + ensure_capacity(&ctx->d_queries, ctx->cap_queries, static_cast(n_queries)); + ensure_capacity(&ctx->d_results, ctx->cap_results, static_cast(n_queries)); + + cudaError_t e1 = cudaMemcpy(ctx->d_queries, queries, n_queries * sizeof(RaycastQuery), cudaMemcpyHostToDevice); + if (e1 != cudaSuccess) + return static_cast(e1); + + int launch_err = scm_launch_raycast_fp64(ctx->d_queries, n_queries, ctx->d_verts, ctx->d_faces, + ctx->n_faces_current, ctx->d_xforms, ctx->d_margins, + ctx->n_bodies_current, ctx->d_results, nullptr); + if (launch_err != cudaSuccess) + return launch_err; + + cudaError_t e2 = cudaDeviceSynchronize(); + if (e2 != cudaSuccess) + return static_cast(e2); + + cudaError_t e3 = cudaMemcpy(out_results, ctx->d_results, n_queries * sizeof(RaycastResult), cudaMemcpyDeviceToHost); + if (e3 != cudaSuccess) + return static_cast(e3); + + return static_cast(cudaSuccess); +} diff --git a/src/chrono_vehicle/terrain/gpu/SCMGpuHost.cpp b/src/chrono_vehicle/terrain/gpu/hip/SCMGpuHost.cpp similarity index 96% rename from src/chrono_vehicle/terrain/gpu/SCMGpuHost.cpp rename to src/chrono_vehicle/terrain/gpu/hip/SCMGpuHost.cpp index 0c83ef60fb..39bfa865e3 100644 --- a/src/chrono_vehicle/terrain/gpu/SCMGpuHost.cpp +++ b/src/chrono_vehicle/terrain/gpu/hip/SCMGpuHost.cpp @@ -1,4 +1,12 @@ // SCMGpuHost.cpp — HIP host bridge: pinned staging, async copy/compute streams, body reduce. +// NOTE: this file is DUPLICATED per GPU backend. Its counterpart is +// terrain/gpu/cuda/SCMGpuHost.cpp, and the two differ only in the runtime API names +// (about thirty symbols). Any change here must be made there too -- nothing enforces it, +// because only one of the two is compiled in a given build. Chrono has no compatibility +// layer for this and deliberately does not grow one; the kernels, which CAN be shared, +// are single-source .cu files instead. +// +// Not a name-for-name mapping: hipHostMalloc(p, n) is cudaHostAlloc(p, n, flags). #include "chrono_vehicle/terrain/SCMGpu.h" diff --git a/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuHost.cpp b/src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp similarity index 95% rename from src/chrono_vehicle/terrain/gpu/SCMRaycastGpuHost.cpp rename to src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp index 9304cea101..9dfeeddbff 100644 --- a/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuHost.cpp +++ b/src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp @@ -1,4 +1,12 @@ // SCMRaycastGpuHost.cpp — HIP host bridge for the SCM ray-cast backend: device buffer management, +// NOTE: this file is DUPLICATED per GPU backend. Its counterpart is +// terrain/gpu/cuda/SCMRaycastGpuHost.cpp, and the two differ only in the runtime API names +// (about thirty symbols). Any change here must be made there too -- nothing enforces it, +// because only one of the two is compiled in a given build. Chrono has no compatibility +// layer for this and deliberately does not grow one; the kernels, which CAN be shared, +// are single-source .cu files instead. +// +// Not a name-for-name mapping: hipHostMalloc(p, n) is cudaHostAlloc(p, n, flags). // synchronous upload/run (v1 -- see SCMRaycastGpu.h for why this isn't pipelined yet). // // Supports two kernel precisions (ScmRaycastGpuPrecision): FP64, the validated default, and FP32, added for GPUs with weak double-precision throughput -- @@ -32,7 +40,7 @@ using chrono::vehicle::scm::gpu::RaycastResult; using chrono::vehicle::scm::gpu::RaycastVertex; // Float mirrors of the double-precision public types. Only the memory layout needs to match the -// .hip.cpp kernel's own VertexDevT/TransformDevT/MarginDevT/QueryDevT/ +// SCMRaycastGpuKernels.cu kernel's own VertexDevT/TransformDevT/MarginDevT/QueryDevT/ // ResultDevT -- not the C++ type identity -- since the two are compiled by different compilers // (g++ here, hipcc there) and only ever communicate via raw device pointers. struct VertexF { From 190ee87aef983586752a9fa80230fe29a6363ee4 Mon Sep 17 00:00:00 2001 From: auc7us Date: Tue, 8 Sep 2026 10:55:34 -0500 Subject: [PATCH 3/5] Drop a duplicated sentence from the SCM GPU backend comment The comment rewrite that came with REQUIRES CUDA_OR_HIP left two phrasings of the same point about HIP platforms. Keep the one that names the ROCm-only libraries. Co-Authored-By: Claude Opus 5 (1M context) --- src/chrono_vehicle/CMakeLists.txt | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/src/chrono_vehicle/CMakeLists.txt b/src/chrono_vehicle/CMakeLists.txt index 95aee0fd7f..d8bd8ad9cf 100644 --- a/src/chrono_vehicle/CMakeLists.txt +++ b/src/chrono_vehicle/CMakeLists.txt @@ -40,13 +40,11 @@ set(CH_USE_SCM_GPU OFF) # the module builds fine with no GPU backend at all, and only the optional SCM # GPU path needs one. The kernels are backend-neutral (single-source .cu compiled as # CUDA or HIP) and the host bridges are duplicated per backend under terrain/gpu/cuda -# and terrain/gpu/hip, so either toolchain is sufficient. Both HIP platforms remain -# acceptable: nothing here uses the ROCm-only libraries that restrict Chrono::DEM and -# Chrono::FSI::SPH to the amd platform. Both HIP platforms are accepted: nothing here depends on -# the ROCm-only libraries (hipCUB/rocPRIM, rocThrust) that restrict -# Chrono::DEM and Chrono::FSI::SPH to the amd platform. So "HIP + ROCm on AMD" -# and "HIP + nvcc on NVIDIA" both give the GPU path with no further -# configuration. +# and terrain/gpu/hip, so either toolchain is sufficient. Both HIP platforms are +# accepted: nothing here depends on the ROCm-only libraries (hipCUB/rocPRIM, +# rocThrust) that restrict Chrono::DEM and Chrono::FSI::SPH to the amd +# platform. So "HIP + ROCm on AMD" and "HIP + nvcc on NVIDIA" both give the GPU +# path with no further configuration. if(CH_ENABLE_VEHICLE_SCM_GPU) chrono_select_gpu_backend(CHRONO_VEHICLE_SCM_BACKEND From 2eb07aaeefabcef62809f2fdb7143768e9d83c17 Mon Sep 17 00:00:00 2001 From: auc7us Date: Tue, 8 Sep 2026 11:01:47 -0500 Subject: [PATCH 4/5] Remove a stray PR draft file PR_MESSAGE.md is a scratch note for a different branch that was swept in by a git add -A. It is not part of the SCM CUDA backend. --- PR_MESSAGE.md | 37 ------------------------------------- 1 file changed, 37 deletions(-) delete mode 100644 PR_MESSAGE.md diff --git a/PR_MESSAGE.md b/PR_MESSAGE.md deleted file mode 100644 index 0db3af76ec..0000000000 --- a/PR_MESSAGE.md +++ /dev/null @@ -1,37 +0,0 @@ -# [UPDATE] Improve SCM memory use and long-run performance - -**Summary** - -This PR improves SCM performance and memory use, especially for large terrain patches and long simulations. - -- Store SCM height and node data in single precision. This cuts the memory used by these fields roughly in half and allows larger or finer terrain grids. -- Use wider counters when creating the SCM visualization mesh. This prevents integer overflow on very large grids and reports a clear error when the mesh index limit is exceeded. -- Store deformed nodes in 16 x 16 tiles. This keeps terrain lookups fast as the vehicle covers more ground and prevents long runs from gradually slowing down. -- Add active-domain and deformed-node counters, along with a way to remove active domains. This makes it easier to monitor SCM work and stop processing bodies that no longer need terrain interaction. -- Limit rendering in affected demos to 60 frames per simulated second. This avoids unnecessary rendering work and makes reported simulation performance more meaningful. - -**Related Issue(s)** - -None. - -**Author(s)** - -Keshav Sharan - -**Licensing** - -By submitting this pull request, I agree that my contribution will be included in Chrono and redistributed under the BSD-3-Clause License. - -**Backward Compatibility** - -There are no input or public API breaks. Public SCM values remain double precision. Small numerical differences are possible because internal terrain storage now uses single precision. - -**Implementation Notes** - -The changes were tested with the Curiosity SCM demo in the `chrono-orb` container. - -**Post Submission Checklist** - -- [x] The changes are complete -- [x] The changes build with CMake -- [x] The SCM demo was tested From 55b1f64ca061dc67279b36c03ddc8ca3f69fccbc Mon Sep 17 00:00:00 2001 From: auc7us Date: Tue, 8 Sep 2026 11:07:26 -0500 Subject: [PATCH 5/5] Correct the ray-cast precision comments Three files claimed FP64 was the validated default, chosen by which HIP platform the build targets. Neither is true: DesiredRaycastGpuPrecision returns FP32 unconditionally, overridable only by SCM_RAYCAST_GPU_PRECISION, and it does so deliberately so the same model does not diverge between machines. Point at that function instead of restating its reasoning, and drop the AMD-only framing now that the kernels also build as CUDA. Co-Authored-By: Claude Opus 5 (1M context) --- .../terrain/gpu/SCMRaycastGpuKernels.cu | 17 ++++++++++------- .../terrain/gpu/cuda/SCMRaycastGpuHost.cpp | 7 ++++--- .../terrain/gpu/hip/SCMRaycastGpuHost.cpp | 7 ++++--- 3 files changed, 18 insertions(+), 13 deletions(-) diff --git a/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.cu b/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.cu index 78c5f36af7..5ec25a65a2 100644 --- a/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.cu +++ b/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.cu @@ -13,13 +13,16 @@ // ~2400-3400 x kThreadsPerRay threads, without changing the total amount of work // (queries x triangles) or the algorithm itself. // -// Templated on Real (double or float) so the same kernel logic runs at either precision. double is the -// validated default for this project's AMD MI300X target. float is offered for GPUs with weak double-precision throughput -- notably consumer NVIDIA -// cards (RTX 4080/5090-class), where FP64 is deliberately throttled relative to FP32 (unlike MI300X, a -// proper datacenter part) -- so a straight double-precision port would be correct there but far slower -// than it needs to be. Both precisions are exported (scm_launch_raycast_fp64 / _fp32); the host bridge -// (SCMRaycastGpuHost.cpp) selects one per-context based on ScmRaycastGpuPrecision, which -// SCMTerrainRaycastGpu.cpp defaults by which HIP platform (AMD vs NVIDIA) this build targets. +// Templated on Real (double or float) so the same kernel logic runs at either precision. FP32 is the +// default, on every backend and every target: making it depend on the hardware would let the same +// model diverge between machines, which is worse than the precision itself. See the comment on +// DesiredRaycastGpuPrecision in SCMTerrainRaycastGpu.cpp for the validation behind that choice. +// FP32 also matters on GPUs with weak double-precision throughput -- notably consumer NVIDIA cards +// (RTX 4080/5090-class), where FP64 is deliberately throttled relative to FP32, unlike a datacenter +// part such as the MI300X -- so a double-only port would be correct there but far slower than it +// needs to be. Both precisions are exported (scm_launch_raycast_fp64 / _fp32) and the host bridge +// (SCMRaycastGpuHost.cpp) selects one per context from ScmRaycastGpuPrecision; env +// SCM_RAYCAST_GPU_PRECISION=fp32|fp64 overrides the default at run time. // // Includes the empirically-determined margin-correction sign; an exact match to Bullet's hit set is // not the goal. diff --git a/src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp b/src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp index 979a1c20ab..41db86c4fb 100644 --- a/src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp +++ b/src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp @@ -9,9 +9,10 @@ // Not a name-for-name mapping: hipHostMalloc(p, n) is cudaHostAlloc(p, n, flags). // synchronous upload/run (v1 -- see SCMRaycastGpu.h for why this isn't pipelined yet). // -// Supports two kernel precisions (ScmRaycastGpuPrecision): FP64, the validated default, and FP32, added for GPUs with weak double-precision throughput -- -// notably consumer NVIDIA cards (e.g. RTX 4080/5090), unlike this project's AMD MI300X target, a proper -// datacenter part with strong FP64. The public API types (SCMRaycastGpuTypes.h) stay double-precision +// Supports two kernel precisions (ScmRaycastGpuPrecision). FP32 is the default on every backend and +// target -- see DesiredRaycastGpuPrecision in SCMTerrainRaycastGpu.cpp -- and it is also what makes +// this usable on GPUs with weak double-precision throughput, notably consumer NVIDIA cards (e.g. RTX +// 4080/5090) as opposed to a datacenter part with strong FP64 such as the MI300X. The public API types (SCMRaycastGpuTypes.h) stay double-precision // throughout -- Chrono itself is double internally -- this file downcasts to float on upload and // upconverts results back to double when precision == kFP32, so callers (SCMTerrainRaycastGpu.cpp) // don't need to know or care which precision is active. diff --git a/src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp b/src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp index 9dfeeddbff..ee593b10ce 100644 --- a/src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp +++ b/src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp @@ -9,9 +9,10 @@ // Not a name-for-name mapping: hipHostMalloc(p, n) is cudaHostAlloc(p, n, flags). // synchronous upload/run (v1 -- see SCMRaycastGpu.h for why this isn't pipelined yet). // -// Supports two kernel precisions (ScmRaycastGpuPrecision): FP64, the validated default, and FP32, added for GPUs with weak double-precision throughput -- -// notably consumer NVIDIA cards (e.g. RTX 4080/5090), unlike this project's AMD MI300X target, a proper -// datacenter part with strong FP64. The public API types (SCMRaycastGpuTypes.h) stay double-precision +// Supports two kernel precisions (ScmRaycastGpuPrecision). FP32 is the default on every backend and +// target -- see DesiredRaycastGpuPrecision in SCMTerrainRaycastGpu.cpp -- and it is also what makes +// this usable on GPUs with weak double-precision throughput, notably consumer NVIDIA cards (e.g. RTX +// 4080/5090) as opposed to a datacenter part with strong FP64 such as the MI300X. The public API types (SCMRaycastGpuTypes.h) stay double-precision // throughout -- Chrono itself is double internally -- this file downcasts to float on upload and // upconverts results back to double when precision == kFP32, so callers (SCMTerrainRaycastGpu.cpp) // don't need to know or care which precision is active.