From 6e6d82b27c931530dd1757856ea06da588a76a9c Mon Sep 17 00:00:00 2001 From: Dan Negrut Date: Sat, 29 Aug 2026 12:39:48 -0500 Subject: [PATCH 1/2] Fix HIP JIT include paths on Windows; report worker-thread exceptions Two fixes found while validating the HIP backend on a Windows machine with the AMD HIP SDK 7.1 (Radeon 8060S, gfx1151): 1. JitHelper: hiprtc on Windows has no implicit ROCm include path, so every runtime kernel compilation failed with 'hip/hip_runtime.h file not found'. The include-path lookup now accepts HIP_PATH (the Windows HIP SDK convention) in addition to ROCM_PATH, and the clang builtin header scan also covers the Windows SDK layout (lib/clang/) in addition to the Linux one (lib/llvm/lib/clang/). 2. kT/dT worker threads: an exception escaping a std::thread calls std::terminate, which on Windows prints nothing, so JIT compilation errors were silently swallowed (exit 0xC0000409 with no message). The worker-thread entry points now catch, report to stderr, and rethrow. Co-Authored-By: Claude Fable 5 --- src/DEM/dT.h | 14 ++++++++++++-- src/DEM/kT.h | 14 ++++++++++++-- src/core/utils/JitHelper.cpp | 25 ++++++++++++++++--------- 3 files changed, 40 insertions(+), 13 deletions(-) diff --git a/src/DEM/dT.h b/src/DEM/dT.h index 689c1a75..c5fa9851 100644 --- a/src/DEM/dT.h +++ b/src/DEM/dT.h @@ -6,6 +6,7 @@ #ifndef DEME_DT #define DEME_DT +#include #include #include #include @@ -317,8 +318,17 @@ class DEMDynamicThread { // spawning the child thread. DEME_GPU_CALL(cudaStreamCreate(&streamInfo.stream)); - // Launch a worker thread bound to this instance - th = std::move(std::thread([this]() { this->workerThread(); })); + // Launch a worker thread bound to this instance. An exception escaping a std::thread calls + // std::terminate; on some platforms (notably Windows) that prints nothing, so report it here + // before dying, otherwise runtime errors like JIT compilation failures are silently swallowed. + th = std::move(std::thread([this]() { + try { + this->workerThread(); + } catch (const std::exception& e) { + std::cerr << "Fatal error in dT worker thread: " << e.what() << std::endl; + throw; + } + })); } ~DEMDynamicThread() { // std::cout << "Dynamic thread closing..." << std::endl; diff --git a/src/DEM/kT.h b/src/DEM/kT.h index 40ba1b04..295c1d67 100644 --- a/src/DEM/kT.h +++ b/src/DEM/kT.h @@ -6,6 +6,7 @@ #ifndef DEME_KT #define DEME_KT +#include #include #include #include @@ -225,8 +226,17 @@ class DEMKinematicThread { // spawning the child thread. DEME_GPU_CALL(cudaStreamCreate(&streamInfo.stream)); - // Launch a worker thread bound to this instance - th = std::move(std::thread([this]() { this->workerThread(); })); + // Launch a worker thread bound to this instance. An exception escaping a std::thread calls + // std::terminate; on some platforms (notably Windows) that prints nothing, so report it here + // before dying, otherwise runtime errors like JIT compilation failures are silently swallowed. + th = std::move(std::thread([this]() { + try { + this->workerThread(); + } catch (const std::exception& e) { + std::cerr << "Fatal error in kT worker thread: " << e.what() << std::endl; + throw; + } + })); } ~DEMKinematicThread() { // std::cout << "Kinematic thread closing..." << std::endl; diff --git a/src/core/utils/JitHelper.cpp b/src/core/utils/JitHelper.cpp index a3620da3..d10695cd 100644 --- a/src/core/utils/JitHelper.cpp +++ b/src/core/utils/JitHelper.cpp @@ -86,20 +86,27 @@ deme::jit::Program JitHelper::buildProgram(const std::string& name, // Common fallbacks #if defined(USE_HIP) - // Helper: find clang builtin headers by scanning lib/llvm/lib/clang//include + // Helper: find clang builtin headers by scanning //include. + // Linux ROCm keeps them under lib/llvm/lib/clang; the Windows HIP SDK under lib/clang. auto add_clang_builtins = [&](const std::filesystem::path& rocm_root) { - std::filesystem::path clang_base = rocm_root / "lib" / "llvm" / "lib" / "clang"; - std::error_code scan_ec; - for (const auto& entry : std::filesystem::directory_iterator(clang_base, scan_ec)) { - if (entry.is_directory()) { - add_inc(entry.path() / "include"); - break; // use the first (and typically only) version dir + for (const auto& clang_base : {rocm_root / "lib" / "llvm" / "lib" / "clang", rocm_root / "lib" / "clang"}) { + std::error_code scan_ec; + for (const auto& entry : std::filesystem::directory_iterator(clang_base, scan_ec)) { + if (entry.is_directory()) { + add_inc(entry.path() / "include"); + break; // use the first (and typically only) version dir + } } } }; - // ROCm include paths for hipRTC - if (const char* rocm_path = std::getenv("ROCM_PATH")) { + // ROCm include paths for hipRTC. ROCM_PATH is the Linux convention; the Windows HIP SDK + // sets HIP_PATH instead, so accept either (hiprtc on Windows has no implicit include path, + // which makes this lookup load-bearing there, not just a fallback). + const char* rocm_path = std::getenv("ROCM_PATH"); + if (rocm_path == nullptr || rocm_path[0] == '\0') + rocm_path = std::getenv("HIP_PATH"); + if (rocm_path != nullptr && rocm_path[0] != '\0') { add_inc(std::filesystem::path(rocm_path) / "include"); add_inc(std::filesystem::path(rocm_path) / "include" / "hipcub"); add_inc(std::filesystem::path(rocm_path) / "include" / "rocprim"); From a8eb50a38405fe23a6f6ce37a4bb392bd7c1dbb6 Mon Sep 17 00:00:00 2001 From: Dan Negrut Date: Sat, 29 Aug 2026 15:00:35 -0500 Subject: [PATCH 2/2] Document validated AMD architectures and the native Windows HIP build Validation performed on the AMD University Program AI & HPC Cluster (this work used computing resources made available through the AMD University Program (AUP) AI & HPC Cluster) and on a Windows 11 machine: - MI300X (gfx942) and MI350X (gfx950), ROCm 7.2.0, Linux: build with -DUSE_HIP=ON and a multi-arch CMAKE_HIP_ARCHITECTURES list; the SingleSphereCollide and Mixer demos run to completion on both, with output files free of NaN/Inf and physically sensible. - Radeon 8060S (gfx1151), HIP SDK 7.1, native Windows: builds with the SDK's clang++ for both CXX and HIP under the Ninja generator (CMake rejects mixing MSVC cl for CXX with clang for HIP); smoke demos pass. Co-Authored-By: Claude Fable 5 --- README.md | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/README.md b/README.md index aef51a09..27140199 100644 --- a/README.md +++ b/README.md @@ -202,6 +202,24 @@ ninja If `CMAKE_HIP_ARCHITECTURES` is left unset it defaults to `gfx90a`. You can list the architecture of an installed device with `rocminfo`. If the ROCm install is not on CMake's default search path, point `-DCMAKE_PREFIX_PATH` at it (e.g. `/opt/rocm`) so `find_package` can locate hip and hipCUB. The demos are then run exactly as in the **Numerical examples** section. +This backend has also been validated on AMD Instinct data-center GPUs: MI300X (`gfx942`) and MI350X (`gfx950`) nodes running ROCm 7.2.0, configured with `-DCMAKE_HIP_ARCHITECTURES="gfx90a;gfx942;gfx950"`. Note that `CMAKE_HIP_ARCHITECTURES` only governs ahead-of-time compiled code; the simulation kernels are compiled at run time by `hiprtc` for whatever device is present, so one multi-architecture build can serve nodes with different GPUs. + +#### AMD GPUs on native Windows + +_DEME_ also builds on Windows with the [AMD HIP SDK](https://rocm.docs.amd.com/projects/install-on-windows/en/latest/) (validated with HIP SDK 7.1 on a Radeon 8060S, `gfx1151`). Two Windows-specific points: + +- CMake does not support mixing MSVC for C++ with clang for HIP, so use the HIP SDK's own clang for both languages. From a VS 2022 Developer shell: + +``` +cmake -G Ninja -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1151 ^ + -DCMAKE_PREFIX_PATH="C:/Program Files/AMD/ROCm/7.1" ^ + -DCMAKE_CXX_COMPILER="C:/Program Files/AMD/ROCm/7.1/bin/clang++.exe" ^ + -DCMAKE_HIP_COMPILER="C:/Program Files/AMD/ROCm/7.1/bin/clang++.exe" .. +ninja +``` + +- At run time, the runtime kernel compilation locates the ROCm headers through the `ROCM_PATH` or `HIP_PATH` environment variable; the HIP SDK installer sets `HIP_PATH`, so no extra setup is normally needed. + ### Windows The process is similar to [the installation of Chrono](https://api.projectchrono.org/tutorial_install_chrono.html), which you can use as reference. The steps depend on your choice of tools, and what listed here are our recommendation.