feat(vehicle): optional HIP SCM contact-force backend (MI300X / gfx942)#755
feat(vehicle): optional HIP SCM contact-force backend (MI300X / gfx942)#755amd-pratmish wants to merge 45 commits into
Conversation
dafbbbe to
c4bca3a
Compare
|
@rserban : Based on your feedback, merged the PR's #756, #757 into this PR #755 https://github.com/projectchrono/chrono/pull/757#issuecomment-4841147303 |
|
@amd-pratmish - thanks for the PR. This looks like a very interesting contribution! Having said that, a few issues and comments:
Please:
Thanks! |
c4bca3a to
4597f63
Compare
Hi @rserban, Thank you for your review and providing fine-granular comments. Based on your feedback, the following changes were made and pushed:
|
538b991 to
90f8a9c
Compare
### Summary Single PR combining the SCM terrain HIP port (foundation + E2E perf + device body reduction), per maintainer request. Adds an **opt-in HIP path** for the SCM Bekker / Mohr-Coulomb / Janosi contact-force loop in `SCMLoader::ComputeInternalForces()`. Ray casting and contact-patch BFS stay on CPU. **Foundation** - CMake: `CH_ENABLE_VEHICLE_SCM_GPU=ON` in `src/chrono_vehicle/CMakeLists.txt` (OpenCRG pattern) + `CHRONO_SCM_GPU_LIB_DIR` - In-tree API headers `SCMGpu.h` / `SCMGpuTypes.h`; external HIP impl in `scm_gpu_core` (`FindScmGpu.cmake`) - Runtime: `CHRONO_SCM_GPU=1` (uniform soil, rigid `ChBody` contactables) - Auto-fallback to CPU when `hits.size() < CHRONO_SCM_GPU_MIN_HITS` (default **8192**) - Host compiler stays **g++**; HIP device code via CMake HIP language. Set `-DCMAKE_HIP_ARCHITECTURES=gfx942` (MI300X) or `gfx90a` (MI210). **E2E perf** - Pass ray-cast `hits` map directly to `ComputeContactForcesGpu` (no intermediate vector copy) - Call `scm_gpu::PrimeBuffers()` from all `SCMTerrain::Initialize` overloads - Dense per-body force accumulation in `SCMTerrainGpu.cpp` **Device body reduction + async bridge** - Grid-only host scatter; body forces from device reduction (`scm_reduce_body_forces_kernel`) Ship kit: amd-chronos contrib/upstream_ready/phase2c/combined
Insert HIP backend reference as its own paragraph after the SCM intro (per review feedback); repair broken mid-sentence insertion.
90f8a9c to
d8714d2
Compare
…ojectchrono#758) * Allow building CRM terrain demos without Chrono::VSG The vehicle terrain demo directory returned early unless CH_ENABLE_MODULE_VSG was enabled, so the CRM terrain demos could not be built headless even though demo_VEH_CRMTerrain_WheeledVehicle and demo_VEH_CRMTerrain_TrackedVehicle have complete no-VSG code paths. Build those two whenever CH_ENABLE_MODULE_FSI_SPH is enabled. demo_VEH_CRMTerrain_MovingPatch includes VSG headers unconditionally and remains gated on VSG. With VSG enabled, the demo set is unchanged. * Guard visualization calls in CRM terrain demos for headless runs Both demo_VEH_CRMTerrain_WheeledVehicle and _TrackedVehicle dereferenced the null visualization system in their simulation loops when built and run without a run-time visualization module. Guard the Synchronize and Advance calls with if(vis), and include ChVehicleVisualSystem.h in the tracked demo (previously pulled in only transitively through the VSG headers). Verified headless on Windows 11 with an AMD Radeon 8060S (gfx1151, HIP backend); the null dereference is platform-independent.
… after initialization (WIP)
…dule. - eliminate the redundant RoboSimian RS_Driver class
…rono#759) * Link Ws2_32 on Windows regardless of compiler The Winsock library was linked into Chrono_core only when CH_COMPILER is COMPILER_MSVC or COMPILER_MSVC_X64. Building with clang on Windows (required for the HIP backend, since CMake mandates a single compiler family across CXX and HIP) left the Winsock symbols unresolved at link time. * Treat unsupported hipMemAdvise as benign in the HIP runtime wrapper hipMemAdvise is not implemented for managed memory on some HIP platforms (native Windows, APUs) and returns hipErrorInvalidValue, which the DEM error-check macro escalated to a fatal abort during solver initialization. Memory advice is a performance hint; tolerate the unsupported-advice error codes and clear the sticky error state. The CUDA path is unchanged. * Document native Windows (HIP SDK) build in the AMD GPU guide Add a Windows section to docs/README_AMD_GPU.md covering the AMD HIP SDK toolchain: the single-compiler-family requirement (ROCm clang for C, CXX, and HIP), CH_ENABLE_OPENMP=OFF (no libomp in the Windows SDK), a reference CMake configuration, run notes, and Windows-specific troubleshooting. Validated on Windows 11 with a Radeon 8060S (gfx1151): FSI-SPH, CRM terrain, and DEM demos. Also list gfx1151 among the ISA examples and add a pointer from the install guide's HIP subsection.
…d scripts (projectchrono#760) * Compile the FSI VSG plugin as HIP when using the HIP backend ChSphVisualizationVSG.cpp includes the FSI data-manager headers, which reference rocThrust/rocPRIM device types that do not compile as plain C++ (error: use of undeclared identifier '__builtin_amdgcn_wavefrontsize'). Compile it with the HIP compiler under CHRONO_HIP_FOUND, the same treatment applied to the FSISPH_GPU_SOURCE_FILES of the main library. This makes CH_ENABLE_MODULE_VSG usable together with CHRONO_GPU_BACKEND=HIP. * Fix glslang configuration in the VSG build scripts Building glslang from a plain clone fails to configure with its default ENABLE_OPT=1 (SPIRV-Tools sources are not present), and the scripts do not check for errors, so VSG silently builds without runtime shader compilation. Every Chrono::VSG application then aborts at startup with 'FATAL: VulkanSceneGraph not compiled with GLSLang, unable to compile shaders.' Configure glslang with ENABLE_OPT=0 (the optimizer is not needed for VSG's GLSL to SPIR-V compilation) and document the Python 3 requirement of the glslang build.
(cherry picked from commit c61d66b)
…rames on a per-object basis
…ctchrono#761) * Use portable compile definitions for WNT in the Python wrapper targets target_compile_definitions() expects definition names, but five SWIG Python targets (cascade, vehicle, robot, parsers, ros) passed the MSVC flag spelling "/DWNT". MSVC's cl tolerated the resulting malformed flag; building the Python module with the ROCm clang toolchain (the HIP-on-Windows configuration of PR projectchrono#759) receives -D/DWNT and fails with "macro name must be an identifier". Pass the plain definition name "WNT" instead; identical behavior under MSVC. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * Update demo_ROBOT_Viper_CRM.py to the renamed EnableGPUErrorCheck The Python Viper-on-CRM demo still calls EnableCudaErrorCheck, which was renamed EnableGPUErrorCheck when the FSI/SPH error-check API was generalized beyond CUDA; the demo crashes with an AttributeError on current main. Only occurrence of the stale name in the repository. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…ed from a URDF file
…materials that are referenced in the associated OBJ file
- Allow testing a custom wheel assembly (which must be derived from ChWheelTestRig::WheelAssembly). - Allow testing a wheel assembly consisting of a Chrono::Vehicle wheel and tire (same as previous implementation). - Rename `ChTireTestRig` to `ChWheelTestRig`. - Rename associated demos for consistency.
|
@amd-pratmish - sorry for the delayed response. I am confused by the
Is there a particular reason why all of this cannot be done transparently and automatically when building the Chrono::Vehicle library (relying only on the already existing support for finding ROCm and HIP)?
Thanks! |
…a ChBodyGeometry object at a given body-relative frame
…ic return mapping poisons p_c In the MCC branch of TauEulerStep, the return-mapping quadratic coefficient b = -K*c_v^2 - 3*G*c_q^2 reaches |b| > sqrt(FLT_MAX) = 1.84e19 once a particle carries a trial deviatoric stress of roughly 390 kPa at clamped moduli, so square(b) overflows to inf in single-precision builds (the default). The inf propagates to delta_lambda, and the hardening update then multiplies p_c by (1 + inf): p_c becomes permanently infinite (the fmax floor does not catch inf), the yield function evaluates to -inf ever after, the particle never yields again, its stress grows without bound, and the simulation later aborts with the generic density-NaN check. Bearing-capacity-scale problems (plate loading at hundreds of kPa) hit this reliably; low-stress terramechanics scenarios never reach the threshold. Fix: (1) perform the quadratic solve in double precision (a few double ops per yielding particle; double-precision builds are numerically unaffected); (2) guard the hardening update against a non-finite plastic strain increment so no single bad step can poison p_c. Validated on plate-loading runs that previously crashed 100% reproducibly (now complete full stroke; an instrumented build confirms zero non-finite events) and against a mu(I) regression run whose output is unchanged. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…n-mapping-fp32-overflow Fix MCC CRM abort (A particle density is NaN): fp32 overflow in plast…
Merge upstream main and streamline the SCM GPU backend per review: - Build HIP kernels directly into Chrono_vehicle (DEM/FSI pattern) - Remove external scm_gpu_core dependency and FindScmGpu.cmake - Replace runtime env vars with SCMTerrain::SetScmGpuConfig API - Document NVIDIA limitation (HIP/ROCm only; CUDA port out of scope)
Move SCMGpu.h include outside chrono::vehicle namespace to avoid nested scm_gpu symbols. Align SCMLoader/SCMTerrain implementation guards with CHRONO_HAS_SCM_GPU from ChConfigVehicle.h.
|
Hi @rserban — thank you for the follow-up. I merged the latest main and streamlined the implementation as requested. Re: external
|
| Case | Result |
|---|---|
CH_ENABLE_VEHICLE_SCM_GPU=OFF — configure + build |
PASS |
CH_ENABLE_VEHICLE_SCM_GPU=ON — configure + HIP build (791/791) |
PASS |
Wheel smoke, CPU path (SetScmGpuEnabled(false)) |
PASS |
Wheel smoke, GPU path (SetScmGpuEnabled(true)) |
PASS |
Implementation is HIP/ROCm only (validated on gfx942 and gfx90a). Chrono DEM/FSI already use separate CUDA and HIP backends; an SCM CUDA port would follow that same pattern and supporting other GPU vendors is out of scope for this PR.
I hope this helps the PR being merged with main.
3992f54 to
17050b8
Compare
Resolve ChConfigVehicle.h.in conflict: keep both CHRONO_CRM and CHRONO_HAS_SCM_GPU substitution macros.
96adb3b to
5955da9
Compare
Summary
Single PR combining the SCM terrain HIP port (foundation + E2E perf + device body reduction), per maintainer request.
Adds an opt-in HIP path for the SCM Bekker / Mohr-Coulomb / Janosi contact-force loop in
SCMLoader::ComputeInternalForces(). Ray casting and contact-patch BFS stay on CPU.Foundation
CH_ENABLE_VEHICLE_SCM_GPU=ONinsrc/chrono_vehicle/CMakeLists.txt(OpenCRG pattern) +CHRONO_SCM_GPU_LIB_DIRSCMGpu.h/SCMGpuTypes.h; external HIP impl inscm_gpu_core(FindScmGpu.cmake)CHRONO_SCM_GPU=1(uniform soil, rigidChBodycontactables)hits.size() < CHRONO_SCM_GPU_MIN_HITS(default 8192)-DCMAKE_HIP_ARCHITECTURES=gfx942(MI300X) orgfx90a(MI210).E2E perf
hitsmap directly toComputeContactForcesGpu(no intermediate vector copy)scm_gpu::PrimeBuffers()from allSCMTerrain::InitializeoverloadsSCMTerrainGpu.cppDevice body reduction + async bridge
scm_reduce_body_forces_kernel)CHRONO_SCM_GPU_ASYNC)OpenMP → HIP split
See
docs/OPENMP_TO_HIP.md— upstream OpenMP→HIP porting guide (SCM reference implementation).Documentation
docs/OPENMP_TO_HIP.md— porting playbookdocs/SCM_GPU_EXTERNAL.md— build/configure externalscm_gpulibraryMotivation
Large SCM batches (multi-vehicle / fine soil grids) spend significant time in the per-hit force loop. Validated on MI300X (gfx942, ROCm 7.x): ~2.5× kernel speedup at 65k hits/step; E2E contact timer improves with device body reduce + async streams.
Scope / non-goals
ChLoadableUV, no spatialRegisterSoilParametersCallbackscm_gpubuilt separately (same pattern as optional FSI HIP toolchain)Files changed
src/chrono_vehicle/CMakeLists.txtsrc/chrono_vehicle/ChConfigVehicle.h.incmake/FindScmGpu.cmakesrc/chrono_vehicle/terrain/SCMGpu.hsrc/chrono_vehicle/terrain/SCMGpuTypes.hsrc/chrono_vehicle/terrain/SCMTerrain.hsrc/chrono_vehicle/terrain/SCMTerrain.cppsrc/chrono_vehicle/terrain/SCMTerrainGpu.hsrc/chrono_vehicle/terrain/SCMTerrainGpu.cppdoxygen/documentation/manuals/vehicle/vehicle_terrain.mddoxygen/documentation/manuals/vehicle/vehicle_terrain_scm_gpu.mddocs/OPENMP_TO_HIP.mddocs/SCM_GPU_EXTERNAL.mdTest plan
CH_ENABLE_VEHICLE_SCM_GPU=OFF) builds with no HIP dependencyscm_parity_test --n-hits 65536PASS (rtol 1e-5; per-hit + body-force parity)scm_scaling_benchmark --n-hits 65536 --steps 200kernel speedup ≥ 2×CHRONO_SCM_GPU=1wheel--load— contact timer improved vs CPU-only pathCHRONO_SCM_GPU_PROFILE=1— steady-state pack/gpu/scatter logged; scatter_ms ≤ CPU-only E2E baselineSupersedes
Previously opened as stacked PRs #756 (E2E) and #757 (body reduce) — closed in favor of this single PR per review feedback.
Review updates (@rserban)
CH_ENABLE_VEHICLE_SCM_GPUlives insrc/chrono_vehicle/CMakeLists.txt(OpenCRG pattern) with HIP +scm_gpu_coreprerequisite checksCHRONO_HAS_SCM_GPUinChConfigVehicle.hSCMGpu.h/SCMGpuTypes.h; link externalscm_gpu_coreviaFindScmGpu.cmakeHitRecorddefinition; Doxygenvehicle_terrain_scm_gpu.md+ manual cross-link (paragraph after SCM intro, not mid-sentence)SCMGpu.hAPI wrapped inextern "C"(linkage match withscm_gpu_core)AMD validation (Radha MI300X / gfx942)
Pre-push gate on
rad-mi300x-2(ROCm 7.2, job 42847):integrate_scm_gpu_upstream.py— vehicle CMake,CHRONO_HAS_SCM_GPU, in-tree headersscm_parity_test --n-hits 65536PASS (max abs err ~3.5e-10)ninja Chrono_vehiclewithCH_ENABLE_VEHICLE_SCM_GPU=ONchrono_scm_gpu_wheel_testCPU + GPU env smoke (100 steps)Not re-run on this reship:
scm_scaling_benchmark,btest_SCM_VEHscaling,--loadfine-grid E2E (≥8k hits/step).