Skip to content

feat(vehicle): optional HIP SCM contact-force backend (MI300X / gfx942)#755

Open
amd-pratmish wants to merge 45 commits into
projectchrono:mainfrom
amd-pratmish:feat/scm-vehicle-hip-gfx942
Open

feat(vehicle): optional HIP SCM contact-force backend (MI300X / gfx942)#755
amd-pratmish wants to merge 45 commits into
projectchrono:mainfrom
amd-pratmish:feat/scm-vehicle-hip-gfx942

Conversation

@amd-pratmish

@amd-pratmish amd-pratmish commented Jun 29, 2026

Copy link
Copy Markdown
Contributor

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)
  • Async HIP streams for pack / compute / scatter (CHRONO_SCM_GPU_ASYNC)

OpenMP → HIP split

[CPU] Update active domains
[CPU] OpenMP ray cast → hits + patch_oob
[CPU] BFS contact patches
[CPU] Pack batch (OpenMP parallel for)
[GPU] scm_compute_forces_kernel  (Bekker + Mohr-Coulomb + Janosi per hit)
[GPU] scm_reduce_body_forces_kernel (v3)
[CPU] Scatter grid / apply ChLoadBodyForce

See docs/OPENMP_TO_HIP.md — upstream OpenMP→HIP porting guide (SCM reference implementation).

Documentation

Motivation

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

  • No GPU ray cast, no FEA/ChLoadableUV, no spatial RegisterSoilParametersCallback
  • No bulldozing / erosion on GPU
  • scm_gpu built separately (same pattern as optional FSI HIP toolchain)

Files changed

  • src/chrono_vehicle/CMakeLists.txt
  • src/chrono_vehicle/ChConfigVehicle.h.in
  • cmake/FindScmGpu.cmake
  • src/chrono_vehicle/terrain/SCMGpu.h
  • src/chrono_vehicle/terrain/SCMGpuTypes.h
  • src/chrono_vehicle/terrain/SCMTerrain.h
  • src/chrono_vehicle/terrain/SCMTerrain.cpp
  • src/chrono_vehicle/terrain/SCMTerrainGpu.h
  • src/chrono_vehicle/terrain/SCMTerrainGpu.cpp
  • doxygen/documentation/manuals/vehicle/vehicle_terrain.md
  • doxygen/documentation/manuals/vehicle/vehicle_terrain_scm_gpu.md
  • docs/OPENMP_TO_HIP.md
  • docs/SCM_GPU_EXTERNAL.md

Test plan

  • Default CMake (CH_ENABLE_VEHICLE_SCM_GPU=OFF) builds with no HIP dependency
  • scm_parity_test --n-hits 65536 PASS (rtol 1e-5; per-hit + body-force parity)
  • scm_scaling_benchmark --n-hits 65536 --steps 200 kernel speedup ≥ 2×
  • CHRONO_SCM_GPU=1 wheel --load — contact timer improved vs CPU-only path
  • CHRONO_SCM_GPU_PROFILE=1 — steady-state pack/gpu/scatter logged; scatter_ms ≤ CPU-only E2E baseline

Supersedes

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_GPU lives in src/chrono_vehicle/CMakeLists.txt (OpenCRG pattern) with HIP + scm_gpu_core prerequisite checks
  • CHRONO_HAS_SCM_GPU in ChConfigVehicle.h
  • In-tree SCMGpu.h / SCMGpuTypes.h; link external scm_gpu_core via FindScmGpu.cmake
  • Fixed duplicate HitRecord definition; Doxygen vehicle_terrain_scm_gpu.md + manual cross-link (paragraph after SCM intro, not mid-sentence)
  • SCMGpu.h API wrapped in extern "C" (linkage match with scm_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 headers
  • scm_parity_test --n-hits 65536 PASS (max abs err ~3.5e-10)
  • ninja Chrono_vehicle with CH_ENABLE_VEHICLE_SCM_GPU=ON
  • chrono_scm_gpu_wheel_test CPU + GPU env smoke (100 steps)

Not re-run on this reship: scm_scaling_benchmark, btest_SCM_VEHscaling, --load fine-grid E2E (≥8k hits/step).

@amd-pratmish

Copy link
Copy Markdown
Contributor Author

@DanNegrut

@rserban rserban self-assigned this Jun 30, 2026
@amd-pratmish
amd-pratmish force-pushed the feat/scm-vehicle-hip-gfx942 branch from dafbbbe to c4bca3a Compare June 30, 2026 21:14
@amd-pratmish

Copy link
Copy Markdown
Contributor Author

@rserban : Based on your feedback, merged the PR's #756, #757 into this PR #755

https://github.com/projectchrono/chrono/pull/757#issuecomment-4841147303

@rserban

rserban commented Jul 1, 2026

Copy link
Copy Markdown
Member

@amd-pratmish - thanks for the PR. This looks like a very interesting contribution!

Having said that, a few issues and comments:

  1. The CMake option CH_ENABLE_VEHICLE_SCM_GPU should be moved in src/chrono_vehicle/CMakeLists.txt (and treated similarly to CH_ENABLE_OPENCRG).

  2. It is not enough to let the user enable "SCM_GPU" via the above CMake option. There should be additional checks to ensure that all prerequisites are available. If some are missing, a message should be displayed and this option disabled. Chrono should configure without additional user involvement in all cases (e.g., even if a GPU is not present at all).

  3. Please also include a macro CHRONO_HAS_SCM_GPU in the automatically generated ChConfigVehicle.h (modify accordingly the template src/chrono_vehicle/ChConfigVehicle.h.in and create the necessary replacement variables in the Chrono::Vehicle CMakeLists.txt) -- see how this is done for CHRONO_HAS_OPENCRG.

  4. SCMTerrain.cpp includes (line 42) "scm_gpu.h". This file is missing. Also, for consistency, please consider renaming it to something like SCMGpu.h or SCMGpuUtils.h (as appropriate -- I don't really know what that file provides).

  5. SCMTerrain.h already includes <unordered_map>. No need to do that again conditional on CHRONO_VEHICLE_SCM_GPU.

  6. When CH_ENABLE_VEHICLE_SCM_GPU=OFF, struct HitRecord is redefined! Lines 1167-1171 in SCMTerrain.cpp should be deleted.

  7. At line 916 of src/chrono_vehicle/CMakeLists.txt you add -fopenmp for clang and hipcc. What about other compilers (e.g., GCC)?

  8. This new feature must be documented somewhere in the Chrono::Vehicle "manual". Please provide a file vehicle_terrain_scm_gpu.md in doxygen/documentation/manuals/vehicle/ that describes prerequisites, configuration, etc. for this feature and reference it (with a bit of context) from the "Deformable SCM (Soil Contact Model)" section of vehicle_terrain.md. The new file vehicle_terrain_scm_gpu.md should probably reference the two documents you added under docs.

  9. Please apply clang-format to files you add or modify before committing them (a config file is available in the Chrono top-level directory).


  • For now, all I can say is that the PR code does build correctly with CH_ENABLE_VEHICLE_SCM_GPU=OFF (provided the fix for [6] above is made).
  • Because of the missing scm_gpu.h ([5] above), I was not able to check the PR with CH_ENABLE_VEHICLE_SCM_GPU=ON.

Please:

  • include the above fixes.
  • check that everything configures, builds, and runs fine in all cases and combinations (with SCM_GPU enabled and disabled, without a GPU, with an NVIDIA or an AMD GPU, with or without hipcc, etc.).

Thanks!

@amd-pratmish
amd-pratmish force-pushed the feat/scm-vehicle-hip-gfx942 branch from c4bca3a to 4597f63 Compare July 1, 2026 18:32
@amd-pratmish

Copy link
Copy Markdown
Contributor Author

@amd-pratmish - thanks for the PR. This looks like a very interesting contribution!

Having said that, a few issues and comments:

  1. The CMake option CH_ENABLE_VEHICLE_SCM_GPU should be moved in src/chrono_vehicle/CMakeLists.txt (and treated similarly to CH_ENABLE_OPENCRG).
  2. It is not enough to let the user enable "SCM_GPU" via the above CMake option. There should be additional checks to ensure that all prerequisites are available. If some are missing, a message should be displayed and this option disabled. Chrono should configure without additional user involvement in all cases (e.g., even if a GPU is not present at all).
  3. Please also include a macro CHRONO_HAS_SCM_GPU in the automatically generated ChConfigVehicle.h (modify accordingly the template src/chrono_vehicle/ChConfigVehicle.h.in and create the necessary replacement variables in the Chrono::Vehicle CMakeLists.txt) -- see how this is done for CHRONO_HAS_OPENCRG.
  4. SCMTerrain.cpp includes (line 42) "scm_gpu.h". This file is missing. Also, for consistency, please consider renaming it to something like SCMGpu.h or SCMGpuUtils.h (as appropriate -- I don't really know what that file provides).
  5. SCMTerrain.h already includes <unordered_map>. No need to do that again conditional on CHRONO_VEHICLE_SCM_GPU.
  6. When CH_ENABLE_VEHICLE_SCM_GPU=OFF, struct HitRecord is redefined! Lines 1167-1171 in SCMTerrain.cpp should be deleted.
  7. At line 916 of src/chrono_vehicle/CMakeLists.txt you add -fopenmp for clang and hipcc. What about other compilers (e.g., GCC)?
  8. This new feature must be documented somewhere in the Chrono::Vehicle "manual". Please provide a file vehicle_terrain_scm_gpu.md in doxygen/documentation/manuals/vehicle/ that describes prerequisites, configuration, etc. for this feature and reference it (with a bit of context) from the "Deformable SCM (Soil Contact Model)" section of vehicle_terrain.md. The new file vehicle_terrain_scm_gpu.md should probably reference the two documents you added under docs.
  9. Please apply clang-format to files you add or modify before committing them (a config file is available in the Chrono top-level directory).
  • For now, all I can say is that the PR code does build correctly with CH_ENABLE_VEHICLE_SCM_GPU=OFF (provided the fix for [6] above is made).
  • Because of the missing scm_gpu.h ([5] above), I was not able to check the PR with CH_ENABLE_VEHICLE_SCM_GPU=ON.

Please:

  • include the above fixes.
  • check that everything configures, builds, and runs fine in all cases and combinations (with SCM_GPU enabled and disabled, without a GPU, with an NVIDIA or an AMD GPU, with or without hipcc, etc.).

Thanks!

Hi @rserban,

Thank you for your review and providing fine-granular comments.

Based on your feedback, the following changes were made and pushed:

  1. Option + prerequisite checks live in vehicle CMake; root CMakeLists.txt no longer touched
  2. Added find_package(HIP) + find_package(ScmGpu) with warnings
  3. CHRONO_SCM_GPU substitution wired in CMake
  4. Headers added; SCMTerrainGpu.cpp includes chrono_vehicle/terrain/SCMGpu.h
  5. Conditional duplicate stripped in integrate script
  6. Deleted duplicate HitRecord struct - dedupe_hitrecord() + validation
  7. Links OpenMP::OpenMP_CXX when target exists. BTW, this PR was originally meant for OpenMP to HIP translation.
  8. New page + link from vehicle_terrain.md
  9. clang format: Applied and integrate scripts on auto-formats on export.

@amd-pratmish
amd-pratmish force-pushed the feat/scm-vehicle-hip-gfx942 branch from 538b991 to 90f8a9c Compare July 1, 2026 20:51
amd-chronos ship-kit and others added 2 commits July 1, 2026 21:04
### 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.
@amd-pratmish
amd-pratmish force-pushed the feat/scm-vehicle-hip-gfx942 branch from 90f8a9c to d8714d2 Compare July 1, 2026 21:04
DanNegrut and others added 17 commits July 3, 2026 13:04
…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.
…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.
rserban and others added 19 commits July 6, 2026 08:45
…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>
…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.
@rserban

rserban commented Jul 14, 2026

Copy link
Copy Markdown
Member

@amd-pratmish - sorry for the delayed response.

I am confused by the scm_gpu_core external dependency:

  • why add the complication of an external dependency for what is meant to be Chrono code?
  • unless I miss something, I don't see how to build this scm_gpu_core library.

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)?
Is there really a need to use environment variables (in `SCMTerrainGGpu.cpp)?

  1. Could you please try to streamline this implementation?
  2. Please also merge the latest code from the main branch.
  3. Finally, is it possible to use this capability also with an Nvidia GPU? If yes, how? If no, why not?

Thanks!

rserban and others added 6 commits July 14, 2026 11:07
…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.
@amd-pratmish

amd-pratmish commented Jul 17, 2026

Copy link
Copy Markdown
Contributor Author

Hi @rserban — thank you for the follow-up. I merged the latest main and streamlined the implementation as requested.
Summary below:

Re: external scm_gpu_core

Agreed- it was a prototype, so no good reason to keep this as a separate library. External scm_gpu_core dependency — removed

  • Deleted cmake/FindScmGpu.cmake and docs/SCM_GPU_EXTERNAL.md
  • HIP code lives in-tree: terrain/gpu/SCMGpuKernels.hip.cpp, SCMGpuHost.cpp
  • Linked into libChrono_vehicle.so when -DCH_ENABLE_VEHICLE_SCM_GPU=ON and HIP/ROCm is found (same pattern as Chrono::DEM / FSI::SPH)

Build is a normal Chrono configure — no separate library, no environment variables:

cmake -S chrono -B build \
  -DCH_ENABLE_MODULE_VEHICLE=ON \
  -DCH_ENABLE_VEHICLE_SCM_GPU=ON \
  -DCHRONO_GPU_BACKEND=HIP \
  -DCHRONO_HIP_ARCHITECTURES=gfx942
ninja -C build Chrono_vehicle

If HIP/ROCm is missing, CMake warns and turns CH_ENABLE_VEHICLE_SCM_GPU off; the CPU SCM path still configures and builds.

Re: environment variables

Removed. Runtime control is via SCMTerrain API:

terrain.SetScmGpuEnabled(true);
auto cfg = terrain.GetScmGpuConfig();
cfg.min_hits = 4096;
terrain.SetScmGpuConfig(cfg);

Validation (MI300X / gfx942, ROCm 7.2)

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.

@amd-pratmish
amd-pratmish force-pushed the feat/scm-vehicle-hip-gfx942 branch from 3992f54 to 17050b8 Compare July 17, 2026 18:41
Resolve ChConfigVehicle.h.in conflict: keep both CHRONO_CRM and
CHRONO_HAS_SCM_GPU substitution macros.
@amd-pratmish
amd-pratmish force-pushed the feat/scm-vehicle-hip-gfx942 branch from 96adb3b to 5955da9 Compare July 17, 2026 19:12
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants