Validate the HIP backend on Instinct MI300X/MI350X and native Windows

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.

src/DEM/dT.h

@@ -6,6 +6,7 @@
 #ifndef DEME_DT
 #define DEME_DT
 
+#include <iostream>
 #include <mutex>
 #include <vector>
 #include <thread>
@@ -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;

src/DEM/kT.h

@@ -6,6 +6,7 @@
 #ifndef DEME_KT
 #define DEME_KT
 
+#include <iostream>
 #include <mutex>
 #include <vector>
 #include <thread>
@@ -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;

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/<version>/include
+    // Helper: find clang builtin headers by scanning <base>/<version>/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");