diff --git a/src/chrono_vehicle/CMakeLists.txt b/src/chrono_vehicle/CMakeLists.txt index a96ab3f35c..d8bd8ad9cf 100644 --- a/src/chrono_vehicle/CMakeLists.txt +++ b/src/chrono_vehicle/CMakeLists.txt @@ -38,18 +38,19 @@ 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 -# 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. +# 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 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 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") @@ -65,7 +66,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() @@ -318,22 +319,36 @@ 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}) + # 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_HIP_FILES} + ${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/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 80% rename from src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.hip.cpp rename to src/chrono_vehicle/terrain/gpu/SCMRaycastGpuKernels.cu index df0666335f..5ec25a65a2 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 @@ -13,18 +13,50 @@ // ~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. -#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 +341,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 +354,7 @@ int launch(const void* queries_dev, static_cast*>(margins_dev), n_bodies, static_cast*>(results_dev)); - return hipGetLastError(); + return gpuGetLastError(); } } // namespace @@ -336,7 +368,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 +382,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); } 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..41db86c4fb --- /dev/null +++ b/src/chrono_vehicle/terrain/gpu/cuda/SCMRaycastGpuHost.cpp @@ -0,0 +1,358 @@ +// 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). 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. +// +// 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 92% rename from src/chrono_vehicle/terrain/gpu/SCMRaycastGpuHost.cpp rename to src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp index 9304cea101..ee593b10ce 100644 --- a/src/chrono_vehicle/terrain/gpu/SCMRaycastGpuHost.cpp +++ b/src/chrono_vehicle/terrain/gpu/hip/SCMRaycastGpuHost.cpp @@ -1,9 +1,18 @@ // 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 -- -// 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. @@ -32,7 +41,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 {