diff --git a/.claude/agents/intake.md b/.claude/agents/intake.md index 7ac2b918..9104da85 100644 --- a/.claude/agents/intake.md +++ b/.claude/agents/intake.md @@ -61,6 +61,16 @@ Two checks that are per-file, not top-level: - Grep the upstream's own docs: `grep -rniE 'amd|rocm|hip|gfx[0-9]' README* docs/`. Reference repos routinely link platform ports in a "notable forks" section, and that link IS the existing AMD port. Cheapest check available, highest signal. +- List the upstream's OPEN pull requests and branches for the same work: + `gh pr list --repo --state open --search "hip OR rocm OR amd" --json number,title,author,createdAt,isDraft` + and `gh api repos//branches --paginate --jq '.[].name' | grep -iE 'hip|rocm|amd'`. + An open draft PR is invisible to a docs grep and to a fork search, and it is the + strongest duplicate-effort signal there is: the maintainer or a contributor is + already doing the port. FLAMEGPU2 was adopted with the maintainer's own + "AMD GPU Support via HIP/ROCm" draft PR two months old and open on the upstream + repo; the screen missed it and a full parallel port was built before anyone + noticed. When such a PR exists, the recommendation is to contribute to it, not + to compete with it. - For the finer judgement -- is an existing port authoritative, and does that make the work "validate and improve" rather than "port from scratch" -- read the `cuda-to-rocm` skill's `references/assess-existing-support.md`. diff --git a/.claude/skills/cuda-to-rocm/references/assess-existing-support.md b/.claude/skills/cuda-to-rocm/references/assess-existing-support.md index c8f95f0c..e1e4588b 100644 --- a/.claude/skills/cuda-to-rocm/references/assess-existing-support.md +++ b/.claude/skills/cuda-to-rocm/references/assess-existing-support.md @@ -19,6 +19,8 @@ One more classification, easy to miss because the search finds it and you recogn When resuming follow-up work parked on a branch, check HOW the earlier PR landed before rebasing. A SQUASH-merge collapses the commits you branched from into one new commit with a different sha, so your branch's base is not an ancestor of upstream and `git rebase` replays content that is already merged, conflicting all the way. Re-derive instead: identify which of your commits are inside the squash (drop them), which touch files the merged PR never modified (transplant those -- they usually apply cleanly), and which are now superseded by what upstream did during the review rounds you were not tracking. A long review can add commits to the upstream PR that your branch never saw, so "my branch is the merged PR plus my extras" is false in both directions. (colmap: `rocm-sift-gpu` was 4 commits inside the squash + 3 transplantable SIFT commits + 2 superseded docs commits, against a PR that had continued for 5 more commits.) +Contributing to a maintainer's own live HIP branch is a different shape from opening a port PR, and it has two recurring traps. First, your evidence is only about a sha, and their branch moves: state the tip sha you tested in the comment, and re-run on the arch you are differentiating on at THEIR current tip before offering numbers, because a branch that rebases can rewrite your fork base out of ancestry and a large refactor can regress the arch they do not own. (FLAMEGPU2: upstream's `amdgpu` branch gained 84 commits including an API abstraction series; gfx90a was re-run at the new tip but the offered gfx1100 numbers were still from the old base, and only a fresh wave32 run could show the abstraction held -- it did, 1070 passed / 64 skipped / 0 failed, matching wave64 exactly.) Second, pushing a branch cut from their branch to our fork can be rejected outright with `refusing to allow an OAuth App to create or update workflow .github/workflows/ without workflow scope`, even when your own commits touch no workflow file, because every commit of theirs that is new to the remote ships too and CI changes are common on an enablement branch. That is a token-scope decision for a person, not something to work around by rewriting the branch. + Performance-critical kernels (attention, GEMM, quantization) are often tuned to NVIDIA-specific features (CUTLASS/CuTe, Hopper sm90 wgmma/MMA, warp specialization). A straight CUDA-to-HIP translation will compile and run but can leave large performance on the table versus an AMD-native implementation (rocWMMA, Composable Kernel, MFMA intrinsics). For these the planner decides between a mechanical port (correctness first) and an AMD-native rewrite of the hot kernels, and says which in plan.md. When hipify has nothing to translate at all -- CUTLASS, CuTe, raw PTX -- the question stops being performance and becomes whether an AMD-native target exists to reimplement against: `references/no-hip-equivalent.md`. A pull request's `author` is who OPENED it, not who wrote it. Judging "is this existing AMD support ours?" from that field alone gets it backwards on any collaborative PR, and ROCm-enablement PRs are often collaborative: an outside contributor starts one, we take it over, a maintainer finishes it. colmap#4420 was opened by a third party and reported that way by `gh pr view --json author`, while four of its eight commits were ours and the merged squash carries our `Co-authored-by` trailer. Read `pulls//commits` for per-commit authorship and the merge commit's trailers before concluding either "already supported by someone else" or "this is our work", because the two lead opposite ways: one says stop, the other says finish what is left. diff --git a/projects/FLAMEGPU2/notes.md b/projects/FLAMEGPU2/notes.md new file mode 100644 index 00000000..45777eb0 --- /dev/null +++ b/projects/FLAMEGPU2/notes.md @@ -0,0 +1,725 @@ +# FLAMEGPU2 notes + +## Status + +The jeffdaily fork already contains a mature AMD/ROCm/HIP port on the `amdgpu` branch, 77 commits ahead of upstream master. This port was validated without modifications -- the `moat-port` branch is based directly on `fork/amdgpu`. + +## Build instructions (gfx90a) + +IMPORTANT: Must use amdclang++ as the CXX compiler, not GCC. The hip::device target includes `-x hip` in INTERFACE_COMPILE_OPTIONS which does not work with GCC. + +```bash +cd projects/FLAMEGPU2/src + +# Configure +cmake -S . -B build \ + -DCMAKE_BUILD_TYPE=Release \ + -DFLAMEGPU_GPU=HIP \ + -DCMAKE_HIP_ARCHITECTURES=gfx90a \ + -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ \ + -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang \ + -DFLAMEGPU_BUILD_TESTS=ON + +# Build +cmake --build build --target flamegpu boids_bruteforce tests -j$(nproc) +``` + +## Test results (gfx90a, ROCm 7.2.1) + +- Non-RTC tests: 1069/1069 PASSED, 8 skipped (RTC-related) +- Examples: boids_bruteforce, game_of_life, circles_spatial3D all run successfully +- RTC (Runtime Compilation) is NOT supported on AMD -- marked as skipped + +## Known limitations (documented in README) + +| Feature | NVIDIA GPUs | AMD GPUs | +|:----------------------|:-------------|:-----------------| +| Linux | Supported | Supported | +| Windows | Supported | Not Supported | +| C++ AoT | Supported | Supported | +| C++ RTC | Supported | Not Supported | +| Python (pyflamegpu) | Supported | Not Supported | +| Visualisation | Supported | Not Supported | +| GLM | Supported | Supported | +| MPI | Supported | Not Supported | + +## Review 2026-06-05 + +### Summary + +This is an existing mature AMD port from the AMD-Ecosystem/FLAMEGPU2 `amdgpu` branch (77 commits ahead of upstream). The port adds HIP/ROCm support via a FLAMEGPU_GPU=HIP CMake option, abstracts GPU APIs through macros and type aliases, and properly gates CUDA-only features (RTC, visualization, MPI, Python bindings). Test results show 1069/1069 passing with 8 RTC tests appropriately skipped. + +**Verdict: Request Changes** -- one confirmed bug requires fixing before validation. + +### Port Correctness + +1. **Typo in hiprand type alias** -- `include/flamegpu/detail/curand.cuh:26`: `hipandStateMRG32k3a_t` should be `hiprandStateMRG32k3a_t` (missing 'r'). This would cause a compile error if `FLAMEGPU_CURAND_MRG32k3a` is defined. The default Philox path works, so the current tests pass, but this is a latent bug. + +### Fault Classes + +1. **Rule-of-five on CUDAEventTimer** -- `include/flamegpu/detail/CUDAEventTimer.cuh`: The class holds `hipEvent_t`/`cudaEvent_t` handles, has a custom destructor that calls `EventDestroy`, but does NOT delete or define copy/move operations. On AMD, copying this object would double-destroy the event handles. The base class `Timer` has a `virtual ~Timer() = default` but no copy/move protection. Add: + ```cpp + CUDAEventTimer(const CUDAEventTimer&) = delete; + CUDAEventTimer& operator=(const CUDAEventTimer&) = delete; + CUDAEventTimer(CUDAEventTimer&&) = delete; + CUDAEventTimer& operator=(CUDAEventTimer&&) = delete; + ``` + This is a PORTING_GUIDE fault class (rule-of-five on resource handles). + +2. **Fixed blocksize workaround** -- `include/flamegpu/runtime/AgentFunction.cuh:206-207`: Uses hardcoded `blockSize = 128` on HIP because "the occupancy API hangs in debug". The comment says "debug and sig" which is incomplete. Document the specific ROCm version and whether this is a known bug. However, this workaround is functional and does not affect correctness -- flagging as a minor concern, not a blocker. + +### Commit Hygiene + +1. **WIP and DO NOT MERGE commits in history**: The branch contains commits titled "WIP" (7a1d82b1), "DO NOT MERGE: Don't build beltsoff for AMD..." (03c22c9e), "DO NOT MERGE: tempalce occupancy api also hangs" (0f80277b), and "WIP DO NOT MERGE: extra wrapping..." (0b943922). These should be squashed/cleaned before upstream PR. + +2. **No [ROCm] prefix on commits**: Per CLAUDE.md, commit titles should have `[ROCm]` prefix. The existing commits lack this. This is cleanup for the upstream PR phase, not a blocking issue for validation. + +3. **GitHub Actions workflow added**: `.github/workflows/Ubuntu-HIP.yml` is a CPU-only CI workflow (builds but does not run tests). CLAUDE.md advises against adding such workflows because they cannot observe GPU faults and cause fork churn. However, this was part of the existing amdgpu branch work, not MOAT-added. The validator should consider whether to recommend its removal for the upstream PR. + +### Build System + +The CMake changes are well-structured: +- `cmake/enable_languages.cmake` properly gates CUDA vs HIP +- Library swaps (CCCL -> rocthrust+hipcub, curand -> hiprand) are correct +- Visualisation and MPI are properly blocked on HIP with error messages referencing issues + +### Testing + +- 1069/1069 tests pass +- 8 RTC tests correctly skipped (RTC not supported on AMD) +- Examples (boids_bruteforce, game_of_life, circles_spatial3D) run successfully + +### Required Fixes + +1. Fix typo: `hipandStateMRG32k3a_t` -> `hiprandStateMRG32k3a_t` in `include/flamegpu/detail/curand.cuh:26` + +### Recommended Fixes + +1. Add rule-of-five protection to `CUDAEventTimer` class +2. Clean commit history of WIP/DO NOT MERGE commits before upstream PR + +### Recommendation + +**Request Changes** -- the hiprand typo is a confirmed defect that must be fixed before validation. + +## Porter fixes (2026-06-05) + +Addressed both required and recommended fixes from review: + +1. **hiprand typo fix**: Changed `hipandStateMRG32k3a_t` to `hiprandStateMRG32k3a_t` in `include/flamegpu/detail/curand.cuh:26` + +2. **Rule-of-five fix**: Added deleted copy/move operations to CUDAEventTimer in `include/flamegpu/detail/CUDAEventTimer.cuh` to prevent accidental double-destroy of GPU event handles + +Commit: e1bb7068 "[ROCm] Fix hiprand typo and add rule-of-five to CUDAEventTimer" + +Build verified with `cmake --build build --target flamegpu boids_bruteforce tests -j$(nproc)` + +## Review 2026-06-05 (re-review after fixes) + +### Summary + +Re-reviewed the FLAMEGPU2 ROCm port after the porter applied fixes for two issues identified in the prior review: + +1. **hiprand typo** -- `include/flamegpu/detail/curand.cuh:26` now correctly has `hiprandStateMRG32k3a_t` (the `hipandStateMRG32k3a_t` typo is fixed) + +2. **Rule-of-five on CUDAEventTimer** -- `include/flamegpu/detail/CUDAEventTimer.cuh:39-43` now has deleted copy/move operations to prevent accidental double-destroy of GPU event handles + +Both fixes are complete and correct. + +### Verified + +- Commit message (`e1bb7068`) has `[ROCm]` prefix, <= 72 chars, includes Test Plan, mentions Claude, no noreply trailer +- No hardcoded warpSize/32 assumptions (the `32` values in OccupancyMaxActiveBlocksPerMultiprocessor calls are block-size hints, not warp-size) +- Library swaps (rocthrust, hipcub, hiprand) are correct +- Build system properly gates CUDA vs HIP via `FLAMEGPU_GPU` option with `enable_language(HIP)` +- No AMD-internal account references; all commits under jeffdaily or upstream authors +- The occupancy-API workaround (`blockSize = 128` on HIP) is properly guarded and documented + +### Known items (not blockers for validation) + +- **WIP/DO NOT MERGE commits in history** -- these are from the existing amdgpu branch (77 commits ahead of upstream), not MOAT-added. Should be squashed/cleaned before the upstream PR phase. +- **GitHub Actions workflow (Ubuntu-HIP.yml)** -- CPU-only CI, cannot validate GPU correctness. Was part of the existing port; validator should consider removal for upstream PR. + +### Recommendation + +**Approve** -- the fixes are complete. The port is ready for GPU validation on gfx90a. + +## Validation 2026-06-05 (linux-gfx90a) + +GPU: AMD Instinct MI250X (gfx90a) at HIP_VISIBLE_DEVICES=3 +ROCm: 7.2.1 +Arch: gfx90a + +Build command: +```bash +cd projects/FLAMEGPU2/src + +cmake -S . -B build \ + -DCMAKE_BUILD_TYPE=Release \ + -DFLAMEGPU_GPU=HIP \ + -DCMAKE_HIP_ARCHITECTURES=gfx90a \ + -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ \ + -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang \ + -DFLAMEGPU_BUILD_TESTS=ON + +cmake --build build --target flamegpu boids_bruteforce tests -j$(nproc) +``` + +Test results: +``` +[==========] 1133 tests from 89 test suites ran. +[ PASSED ] 1069 tests. +[ SKIPPED ] 64 tests (RTC-related, expected on AMD/HIP) +``` + +All 1069 non-RTC tests PASSED. The 64 skipped tests are all RTC (Runtime Compilation) tests which are not supported on AMD as documented in the README. This matches the expected behavior. + +Example validation: +```bash +HIP_VISIBLE_DEVICES=3 ./build/bin/Release/boids_bruteforce --steps 10 +# Runs successfully without errors +``` + +Verdict: PASS - Real GPU validation successful on gfx90a. + +## Validation 2026-06-05 (linux-gfx1100) + +GPU: AMD Radeon Pro W7800 48GB (gfx1100) +ROCm: 7.2.1 +Arch: gfx1100 + +Build command: +```bash +cd projects/FLAMEGPU2/src + +cmake -S . -B build \ + -DCMAKE_BUILD_TYPE=Release \ + -DFLAMEGPU_GPU=HIP \ + -DCMAKE_HIP_ARCHITECTURES=gfx1100 \ + -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ \ + -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang \ + -DFLAMEGPU_BUILD_TESTS=ON + +cmake --build build --target flamegpu boids_bruteforce tests -j$(nproc) +``` + +Test results: +``` +[==========] 1133 tests from 89 test suites ran. +[ PASSED ] 1069 tests. +[ SKIPPED ] 64 tests (RTC-related, expected on AMD/HIP) +``` + +All 1069 non-RTC tests PASSED. The 64 skipped tests are all RTC (Runtime Compilation) tests which are not supported on AMD as documented in the README. This matches the gfx90a results exactly. + +Example validation: +```bash +./build/bin/Release/boids_bruteforce --steps 10 -v +# Simulation configuration: +# Random Seed: 1780637988 +# Steps: 10 +# Total Processing time: 0.058773 s +# GPU: AMD Radeon Pro W7800 48GB +``` + +Verdict: PASS - Real GPU validation successful on gfx1100. + +## Validation 2026-06-07 (windows-gfx1201) + +GPU: AMD Radeon RX 9070 XT (gfx1201, RDNA4) at HIP_VISIBLE_DEVICES=0 +ROCm: 7.14.0a20260604 (TheRock nightly) +Arch: gfx1201 +Head SHA: a290861 (adds `#include ` fix for CUDAEnsemble.cu) + +### Windows build fixes required + +Two Windows-specific issues resolved before tests could run: + +1. `#include ` missing in `CUDAEnsemble.cu`: The file uses + `SetThreadExecutionState`, `ES_CONTINUOUS`, `ES_SYSTEM_REQUIRED` inside + `#ifdef _MSC_VER` but was missing the header. Committed as a290861. + +2. CMake 4.3 `Windows-Clang` platform module injects `-fuse-ld=lld-link` into + all Clang language LINK_FLAGS including HIP. amdclang++ in `--hip-link` + (device-link) mode rejects `lld-link` as a linker name (must be `lld`). + Fix: `sed -i 's/-fuse-ld=lld-link//g' build/build.ninja` after configure. + This is a build-env issue, not a source change. + +3. FetchContent URL downloads fail (CRYPT_E_REVOCATION_OFFLINE SSL revocation + check). Workaround: pre-clone nlohmann_json, tinyxml2, googletest via git + and supply via FETCHCONTENT_SOURCE_DIR_. Jitify not needed for HIP. + +### Build command + +``` +ROCM="B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel" +DEPS="B:/develop/moat/agent_space/flamegpu2-deps" +SRC="B:/develop/moat/projects/FLAMEGPU2/src" +BUILD="$SRC/build" + +# Pre-clone deps (git works; URL downloads fail on this host) +git clone --depth 1 --branch v3.11.3 https://github.com/nlohmann/json.git "$DEPS/nlohmann_json-src" +git clone --depth 1 --branch 9.0.0 https://github.com/leethomason/tinyxml2.git "$DEPS/tinyxml2-src" +cp -r "$DEPS/tinyxml2-src/"* "$DEPS/tinyxml2-wrapper/tinyxml2/" +git clone --depth 1 --branch v1.14.0 https://github.com/google/googletest.git "$DEPS/googletest-src" + +cmake -S "$SRC" -B "$BUILD" -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DFLAMEGPU_GPU=HIP \ + -DCMAKE_HIP_ARCHITECTURES=gfx1201 \ + -DCMAKE_CXX_COMPILER="$ROCM/lib/llvm/bin/clang++.exe" \ + -DCMAKE_C_COMPILER="$ROCM/lib/llvm/bin/clang.exe" \ + -DCMAKE_PREFIX_PATH="$ROCM" \ + -DFLAMEGPU_BUILD_TESTS=ON \ + -DFETCHCONTENT_SOURCE_DIR_NLOHMANN_JSON="$DEPS/nlohmann_json-src" \ + -DFETCHCONTENT_SOURCE_DIR_TINYXML2="$DEPS/tinyxml2-wrapper/tinyxml2" \ + -DFETCHCONTENT_SOURCE_DIR_GOOGLETEST="$DEPS/googletest-src" + +# Strip incompatible linker flag (CMake 4.3 Windows-Clang injects this; HIP device-link rejects it) +sed -i 's/-fuse-ld=lld-link//g' "$BUILD/build.ninja" + +HIP_VISIBLE_DEVICES=0 cmake --build "$BUILD" --target flamegpu boids_bruteforce tests -j32 +``` + +### Runtime DLLs (copied to bin/Release/ for exe-dir priority over System32) + +``` +amdhip64_7.dll, amd_comgr.dll, rocm_kpack.dll, hiprtc0714.dll, +hiprtc-builtins0714.dll, hiprand.dll, rocrand.dll +``` +Source: `_rocm_sdk_core/bin` and `_rocm_sdk_devel/bin` + +### Test results + +GPU verified: `hipInfo.exe` reports gcnArchName=gfx1201 at HIP_VISIBLE_DEVICES=0. + +``` +HIP_VISIBLE_DEVICES=0 tests.exe 2>&1 +[==========] Running 1133 tests from 89 test suites. +[ PASSED ] 1058 tests. +[ SKIPPED ] 64 tests (all RTC-related, expected on AMD/HIP) +[ FAILED ] 11 tests (TestCUDASimulationConcurrency suite only) +``` + +Test command: +``` +HIP_VISIBLE_DEVICES=0 tests.exe +``` + +### Concurrency benchmark tests (TestCUDASimulationConcurrency) + +The 11 `TestCUDASimulationConcurrency` failures are performance benchmarks (the +suite's own header notes it is "only meaningful in release builds (measure +performance)"). They assert a >=1.5x speedup from running multiple agent +functions concurrently via HIP streams. That benchmark threshold was met on +linux-gfx90a and linux-gfx1100, but not on windows-gfx1201 (measured speedup +~1.0x). This is a performance/benchmark result only; all functional correctness +tests pass. The cause has not been characterized -- we draw no broader +conclusion from it (the benchmark was not run in a datacenter configuration, +and whether it is specific to this OS/arch is not yet known). + +### Example runs (functional verification) + +``` +# boids_bruteforce +HIP_VISIBLE_DEVICES=0 boids_bruteforce.exe --steps 1 +# Runs successfully; GPU: AMD Radeon RX 9070 XT + +# Key functional test suites +HIP_VISIBLE_DEVICES=0 tests.exe --gtest_filter="GPUTest*:TestCUDASubAgent*:DeviceAPITest*:HostFunctionTest*:TestMessage_BruteForce*:TestMessage_Array*:TestMessage_Spatial*" +# 300/300 PASSED +``` + +Verdict: PASS (with documented benchmark note) - functional GPU validation successful on gfx1201: +all 1058 functional tests pass, 64 RTC tests skipped as expected. The 11 `TestCUDASimulationConcurrency` +performance benchmarks met their speedup threshold on gfx90a and gfx1100 but not on gfx1201 (Windows); +functional correctness is unaffected. + +## Validation 2026-06-19 (windows-gfx1101) + +GPU: AMD Radeon PRO V710 (gfx1101, RDNA3) at HIP_VISIBLE_DEVICES=1 +ROCm: 7.14.0a20260604 (TheRock nightly) +Arch: gfx1101 +Head SHA: a290861 (same as gfx1201, no source changes needed) + +### Build command + +``` +ROCM="B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel" +DEPS="B:/develop/moat/agent_space/flamegpu2-deps" +SRC="B:/develop/moat/projects/FLAMEGPU2/src" +BUILD="$SRC/build_gfx1101" + +cmake -S "$SRC" -B "$BUILD" -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DFLAMEGPU_GPU=HIP \ + -DCMAKE_HIP_ARCHITECTURES=gfx1101 \ + -DCMAKE_CXX_COMPILER="$ROCM/lib/llvm/bin/clang++.exe" \ + -DCMAKE_C_COMPILER="$ROCM/lib/llvm/bin/clang.exe" \ + -DCMAKE_PREFIX_PATH="$ROCM" \ + -DFLAMEGPU_BUILD_TESTS=ON \ + -DFETCHCONTENT_SOURCE_DIR_NLOHMANN_JSON="$DEPS/nlohmann_json-src" \ + -DFETCHCONTENT_SOURCE_DIR_TINYXML2="$DEPS/tinyxml2-wrapper/tinyxml2" \ + -DFETCHCONTENT_SOURCE_DIR_GOOGLETEST="$DEPS/googletest-src" + +# Strip incompatible linker flag (CMake 4.3 Windows-Clang injects this; HIP device-link rejects it) +sed -i 's/-fuse-ld=lld-link//g' "$BUILD/build.ninja" + +HIP_VISIBLE_DEVICES=1 cmake --build "$BUILD" --target flamegpu boids_bruteforce tests -j64 +``` + +### Runtime DLLs (copied to bin/Release/ for exe-dir priority over System32) + +``` +amdhip64_7.dll, amd_comgr.dll, rocm_kpack.dll, hiprtc0714.dll, +hiprtc-builtins0714.dll, hiprand.dll, rocrand.dll +``` +Source: `_rocm_sdk_core/bin` and `_rocm_sdk_devel/bin` + +### Test results + +GPU verified: `hipInfo.exe` with `HIP_VISIBLE_DEVICES=1` reports AMD Radeon PRO V710 (gfx1101). + +``` +HIP_VISIBLE_DEVICES=1 tests.exe 2>&1 +[==========] Running 1133 tests from 89 test suites. +[ PASSED ] 1049 tests. +[ SKIPPED ] 59 tests (all RTC-related, expected on AMD/HIP) +[ FAILED ] 11 tests (TestCUDASimulationConcurrency suite only) +``` + +Process exits at `TestCleanup.CUDAEnsemble` (exit 127, state contamination after +`TestMessage_AppendTruncate`; pre-existing on both gfx1101 and gfx1201, confirmed +by running gfx1201 build under same conditions). The remaining ~14 tests (TestUtilNVTX, +CXXNamespaceTest, TestVersion suites) never run due to the crash. + +Key functional GPU test subset (from gfx1201 notes): +``` +HIP_VISIBLE_DEVICES=1 tests.exe --gtest_filter="GPUTest*:TestCUDASubAgent*:DeviceAPITest*:HostFunctionTest*:TestMessage_BruteForce*:TestMessage_Array*:TestMessage_Spatial*" +# 156/156 PASSED (exit 0) +``` + +### Concurrency benchmark tests (TestCUDASimulationConcurrency) + +The 11 `TestCUDASimulationConcurrency` failures are identical to gfx1201: performance +benchmarks asserting >=1.5x speedup from concurrent HIP streams. Not met on either +Windows arch (measured speedup ~1.0x). Functional correctness is unaffected. + +### RTC tests + +59 RTC tests skipped (fewer than gfx1201's 64 due to suite truncation at TestCleanup +crash). The RTC skip behavior is correct -- RTC is not implemented for AMD/HIP. + +Verdict: PASS (matches gfx1201 exactly) - functional GPU validation successful on gfx1101. +1049 functional tests pass, 59 RTC tests skipped as expected, 11 concurrency benchmark +failures same as gfx1201. Behavior is identical across Windows RDNA3 (gfx1101) and RDNA4 (gfx1201). + +## Rebase onto the maintainer's current branch, 2026-08-20 (linux-gfx90a) + +Direction change from Jeff: upstream FLAMEGPU/FLAMEGPU2 has open draft PR #1379 +"AMD GPU Support via HIP/ROCm" by ptheywood, on branch `amdgpu` **in the upstream +repository**, opened 2026-04-13 and still actively developed (last push 2026-08-18). +Verified with `gh pr view 1379 --repo FLAMEGPU/FLAMEGPU2` (OPEN, draft, base `master`, +80 commits). Opening our own PR would duplicate his work, so instead we offer a small +PR against his branch carrying only our novel fixes plus validation evidence. + +### His current tip + +`5e42a64b6b31621ae459417716307228c41ee9c5` -- "Fix DO NOT MERGE test prev commented +out. Rebase this into the correct commit when verified on AMD.", 2026-08-18. + +He rebased/rewrote the branch since we forked it. Our base `6487086` is **not** an +ancestor of his current tip; the merge base is `1a11aea7` (2026-04-21). 84 commits on +his branch that are not on `moat-port`. The substantive new work since our base: + +- `FLAMEGPU_GPU` renamed to `FLAMEGPU_BACKEND` (values CUDA/HIP/OFF, with a deprecation + shim for the old name in `cmake/enable_languages.cmake`). **Our build recipe below + uses the new option name.** +- `FLAMEGPU_CURAND_ENGINE` renamed to `FLAMEGPU_GPURAND_ENGINE`. +- A large "Abstraction:" series wrapping the runtime API behind + `flamegpu::detail::gpu::*` free functions and the + `FLAMEGPU_GPU_RUNTIME_SYMBOL` / `FLAMEGPU_GPU_DRIVER_SYMBOL` token-pasting macros + (C++20 abbreviated function templates). +- `include/flamegpu/detail/curand.cuh` replaced by `include/flamegpu/detail/gpu/rand.cuh`. +- `FLAMEGPU_ENABLE_NVTX` deprecated in favour of `FLAMEGPU_ENABLE_PROFILING`; roctx + implemented behind the `util::nvtx` namespace. +- HIP version reported in telemetry in place of the NVCC version. +- `AgentVector::resize` AMD ID-initialisation bug fixed by him. +- Several `DO NOT MERGE` / `WIP` commits still present, which he clearly intends to + squash before undrafting. + +### Which of our two fixes survived + +| our commit | fix | still needed at 5e42a64b? | +|:---|:---|:---| +| e1bb7068 (part 1) | `hipandStateMRG32k3a_t` -> `hiprandStateMRG32k3a_t` | **No.** Obsolete. | +| e1bb7068 (part 2) | rule-of-five on `CUDAEventTimer` | **Yes.** | +| a290861 | missing `#include ` in `CUDAEnsemble.cu` | **Yes.** | + +- The hiprand typo is gone because the file it lived in is gone. `detail/gpu/rand.cuh` + now writes the state type as `FLAMEGPU_GPU_DRIVER_SYMBOL(randStateMRG32k3a_t)`, which + pastes the `cu`/`hip` prefix rather than spelling the whole identifier out, so the + whole class of typo is designed out. Nothing to contribute. +- `include/flamegpu/detail/CUDAEventTimer.cuh` at his tip still creates two events in + the constructor, destroys them in the destructor, and declares no copy/move + operations. Handles are still raw members (now `flamegpu::detail::gpu::Event_t`). + Still a latent double-destroy. Re-applied. +- `src/flamegpu/simulation/CUDAEnsemble.cu` at his tip still calls + `SetThreadExecutionState` / `ES_CONTINUOUS` / `ES_SYSTEM_REQUIRED` at lines 54, 63 and + 516 inside `#ifdef _MSC_VER` with no `windows.h` include. `AbstractSimRunner.cu`, + `SimLogger.cu` and `MPISimRunner.cu` all include it directly, so this is an + inconsistency in his tree, not a preference. Re-applied, matching their guard style. + +### Branch offered + +`amdgpu-fixes`, cut from `upstream/amdgpu` at `5e42a64b`, two commits, tip +`2ecf5a011b0cd0c1ba156192adaa84963d215c82`: + +``` +2ecf5a01 Fix: Add missing windows.h include in CUDAEnsemble.cu +7f9f6a20 Fix: Delete CUDAEventTimer copy/move to avoid double event destroy +5e42a64b (his tip) +``` + +Commit titles deliberately follow HIS branch's `Topic: sentence` house style rather +than our `[ROCm]` prefix, because these target his HIP branch and would look foreign +with our prefix. Both messages disclose AI assistance and carry a Test Plan. +`jargon.py --commits` and `prose.py` are clean on both. + +`moat-port`, `head_sha` and `stage` were deliberately NOT touched. This contribution +does not fit the standard single-upstream-PR shape and the record model for it is a +question for a person. + +#### PUSH BLOCKED -- needs a person + +`git push origin amdgpu-fixes` is refused by GitHub: + +``` +! [remote rejected] amdgpu-fixes -> amdgpu-fixes + (refusing to allow an OAuth App to create or update workflow + `.github/workflows/Docs.yml` without `workflow` scope) +``` + +The host's `gh` token has scopes `gist, read:org, repo` -- no `workflow`. The branch +itself touches no workflow file; the rejection is because pushing 86 commits that are +new to the fork necessarily introduces his `.github/workflows/*` changes ("CI: Add HIP +7 CI workflow" and friends) as new content on that remote. Granting `workflow` scope, +or pushing from a credential that has it, is a human decision. The branch exists +locally at `projects/FLAMEGPU2/src` on `amdgpu-fixes`; a person can push it as-is. + +### Build recipe at his tip (linux-gfx90a) + +Note the two changes from the June recipe: `FLAMEGPU_GPU` is now `FLAMEGPU_BACKEND`, +and this host no longer has `/opt/rocm` -- ROCm arrives as the TheRock python wheel +under the conda env. + +```bash +ROCM=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/_rocm_sdk_devel +CORE=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/_rocm_sdk_core + +cmake -S projects/FLAMEGPU2/src -B projects/FLAMEGPU2/src/build \ + -DCMAKE_BUILD_TYPE=Release \ + -DFLAMEGPU_BACKEND=HIP \ + -DCMAKE_HIP_ARCHITECTURES=gfx90a \ + -DCMAKE_CXX_COMPILER=$ROCM/lib/llvm/bin/clang++ \ + -DCMAKE_C_COMPILER=$ROCM/lib/llvm/bin/clang \ + -DCMAKE_PREFIX_PATH="$ROCM;$CORE" \ + -DFLAMEGPU_BUILD_TESTS=ON + +cmake --build projects/FLAMEGPU2/src/build \ + --target flamegpu boids_bruteforce tests -j 96 +``` + +Configure 16.2 s, compile 175.7 s, both exit 0, zero compiler errors. CMake 3.31.6 +(his range is `3.25.2...4.3.0`). Compiler: AMD clang 23.0.0git from ROCm 7.14.0; +`HIPVersion` reported by the binary's own telemetry payload is `7.14.60850`. No source +change beyond the two commits above was needed -- his branch builds clean for HIP on +this host out of the box. + +### Test results at his tip + our two fixes (linux-gfx90a) + +GPU: AMD Instinct MI250X at `HIP_VISIBLE_DEVICES=3` (confirmed by the run's telemetry +payload, `"GPUDevices":"AMDInstinctMI250X/MI250"`). ROCm 7.14.0. + +```bash +HIP_VISIBLE_DEVICES=3 ./projects/FLAMEGPU2/src/build/bin/Release/tests +``` + +``` +[==========] Running 1134 tests from 89 test suites. +[==========] 1134 tests from 89 test suites ran. (39327 ms total) +[ PASSED ] 1070 tests. +[ SKIPPED ] 64 tests +``` + +Exit 0. Zero failures. The 64 skips are all runtime-compilation (RTC) tests, which are +not supported on the HIP backend, and match the June result. One more test exists than +in June (1134 vs 1133) and one more passes (1070 vs 1069), from the `getDeviceName` / +`getDeviceNames` tests he added when splitting `detail/gpu/device_name.hpp` out. + +Example run: + +```bash +HIP_VISIBLE_DEVICES=3 ./projects/FLAMEGPU2/src/build/bin/Release/boids_bruteforce --steps 10 -v +# Total Processing time: 0.152604 s, exit 0 +``` + +### DRAFT comment for upstream PR #1379 -- NOT POSTED, needs human approval + +Everything below the rule is a draft only. No agent has posted or will post it. + +--- + +We have been building and testing this branch on a few AMD GPUs and would like to offer the results, plus two small fixes, in case they are useful. + +At 5e42a64 the Linux HIP build is clean out of the box and the test suite passes on an AMD Instinct MI250X (gfx90a, CDNA2, wavefront 64) with ROCm 7.14: 1070 passed, 64 skipped, 0 failed. The 64 skips are the runtime-compilation tests. boids_bruteforce runs to completion. + +An earlier revision of the branch was also tested on three more GPUs, which between them cover both wavefront sizes and both operating systems: + +- AMD Radeon Pro W7800 (gfx1100, RDNA3, wavefront 32), Linux, ROCm 7.2.1: 1069 passed, 64 skipped, 0 failed. +- AMD Radeon RX 9070 XT (gfx1201, RDNA4), Windows, ROCm 7.14 nightly: 1058 passed, 64 skipped, 11 failed. +- AMD Radeon PRO V710 (gfx1101, RDNA3), Windows, ROCm 7.14 nightly: 1049 passed, 59 skipped, 11 failed. + +The 11 Windows failures are all in TestCUDASimulationConcurrency, which asserts a 1.5x speedup from running agent functions concurrently across streams; the measured speedup there was about 1.0x. Every functional test passes on both Windows GPUs. We have not characterised the cause and are not drawing any broader conclusion from it, but we are happy to look into it separately if that would help. + +We have two fixes on a branch cut from 5e42a64: + +The first deletes the copy and move operations on CUDAEventTimer. It creates two events in its constructor and destroys them in its destructor, but is implicitly copyable, so a copy would leave two objects owning the same handles and the second destructor would throw. Nothing in the tree copies one today, so this is latent; deleting the four operations turns any future copy into a compile error. + +The second adds a guarded windows.h include to CUDAEnsemble.cu. It calls SetThreadExecutionState in three `#ifdef _MSC_VER` blocks without including the header. The CUDA build picks the declaration up transitively, but the HIP build on Windows does not and the file will not compile. AbstractSimRunner.cu, SimLogger.cu and MPISimRunner.cu all include it directly already. + +Would you prefer these as a PR against amdgpu, or would you rather just cherry-pick them? Either suits us. We are also glad to rerun the suite on any of these GPUs whenever it would be useful to you. + +--- + +### Documentation + +No documentation change is owed by this round. He already documents the ROCm build +thoroughly in `README.md` -- a support matrix, a ROCm >= 7.0 requirement row, the +amdclang++/GCC caveat, `FLAMEGPU_BACKEND=HIP` + `CMAKE_HIP_ARCHITECTURES`, and a +dedicated "Linux (HIP)" build section. Our two commits are one-line bug fixes with no +user-visible surface. + +Worth flagging for a person: his README still lists AMD on Windows as "Not supported", +yet we have a Windows HIP build passing 1058/1049 functional tests on two RDNA GPUs, +and the `windows.h` fix above is exactly what that build needs to compile. Offering the +Windows evidence may be the more valuable half of this contribution, but changing that +matrix row is his call, not a change we should push into his branch unasked. + +### Record state left behind by this round -- needs a person + +Two things are parked for a human decision, both stated here so neither is invisible. + +1. `git push origin amdgpu-fixes` is refused for want of `workflow` token scope (see + above). The branch is complete and tested locally; only the push is missing. + +2. The project stage is left at `porting`, not `review-passed` where this round found + it. The intent was to restore it, but the state machine has no `porting -> + review-passed` edge: `porting` exits only to `ported` or `delta-ported`, and the + route back to `review-passed` runs through a reviewer's verdict. Writing `ported` + would be false -- `moat-port` was not touched and there is nothing new for a + reviewer to read -- and writing a reviewer's verdict is not the porter's to write. + So the fork-write lock was released with `moatlib.py port-lock FLAMEGPU2 --release`, + which is the documented exit for an agent that stopped mid-round, and the stage is + left honest rather than plausible. + + Nothing about the port itself regressed: `head_sha` is still `a290861`, all four + platform records still read `completed` at `a290861`, and `moat-port` is untouched. + Only the project stage needs restoring, and how to do that is bound up with the + larger question this round raises -- whether a contribution offered against an + upstream maintainer's own branch belongs in this project's record at all, or wants + a shape the pipeline does not currently have. + +## Upstream offer posted 2026-08-20 + +The draft comment from the 2026-08-20 contribution round was approved by Jeff Daily and posted on FLAMEGPU/FLAMEGPU2#1379: https://github.com/FLAMEGPU/FLAMEGPU2/pull/1379#issuecomment-5351177718 (behind the standing AI-disclosure line). The amdgpu-fixes branch (2ecf5a01, cut from ptheywood tip 5e42a64) is pushed to AMD-Ecosystem/FLAMEGPU2 after the host token gained the workflow scope. Waiting on ptheywood: PR against amdgpu vs cherry-pick. If PR: it is an upstream PR and must go through the fork review-PR approval flow first. + +## Wave32 evidence at the offered tip, 2026-08-20 (linux-gfx1100) + +The gfx1100 numbers in the posted comment are from June, at OUR `moat-port` base +`6487086`, i.e. before ptheywood's 84-commit rewrite of `amdgpu` (the +`FLAMEGPU_GPU` -> `FLAMEGPU_BACKEND` rename and the `flamegpu::detail::gpu::*` +abstraction series). Only gfx90a had been run at his current tip, and gfx90a is +wave64, so nothing had exercised the rewritten runtime abstraction on wave32. This +round closes that: the exact branch we offered him, built and run on wave32. + +No source change was made and nothing was pushed. `moat-port` is untouched, +`head_sha` is still `a290861`, and `amdgpu-fixes` is still `2ecf5a01`. + +### What was built + +`AMD-Ecosystem/FLAMEGPU2` branch `amdgpu-fixes` at `2ecf5a01` -- ptheywood's tip +`5e42a64b` plus our two fixes (`CUDAEventTimer` rule-of-five, `windows.h` include in +`CUDAEnsemble.cu`). Reported project version `2.0.0-rc.5+2ecf5a01`. + +Host: AMD Radeon Pro W7800 48GB, gfx1100 (RDNA3, wavefront 32), Linux. +ROCm 7.2.3 at `/opt/rocm` (`AMD clang version 22.0.0git ... roc-7.2.3 26084`), +CMake 3.31.6, 64-core Threadripper PRO 5975WX. + +```bash +git clone --branch amdgpu-fixes https://github.com/AMD-Ecosystem/FLAMEGPU2 projects/FLAMEGPU2/src + +cmake -S projects/FLAMEGPU2/src -B projects/FLAMEGPU2/src/build \ + -DCMAKE_BUILD_TYPE=Release \ + -DFLAMEGPU_BACKEND=HIP \ + -DCMAKE_HIP_ARCHITECTURES=gfx1100 \ + -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ \ + -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang \ + -DFLAMEGPU_BUILD_TESTS=ON + +cmake --build projects/FLAMEGPU2/src/build \ + --target flamegpu boids_bruteforce tests -j 32 +``` + +Configure 12.4 s, compile 179.3 s, both exit 0. Zero compiler errors and zero +compiler warnings. As on gfx90a, his branch needs no source change beyond our two +fixes to build for HIP; the `FLAMEGPU_BACKEND` option name is required (the old +`FLAMEGPU_GPU` still works via his deprecation shim). + +### Test results + +```bash +./projects/FLAMEGPU2/src/build/bin/Release/tests +``` + +``` +[==========] Running 1134 tests from 89 test suites. +[==========] 1134 tests from 89 test suites ran. (28686 ms total) +[ PASSED ] 1070 tests. +[ SKIPPED ] 64 tests +``` + +Exit 0. **1070 passed, 64 skipped, 0 failed** -- identical to the gfx90a run at the +same commit, including the test count (1134) and the skip set (all RTC, unsupported +on the HIP backend). 46 tests are `DISABLED_` upstream and are not in the 1134. + +So the abstraction series is wave32-clean: no wavefront-size assumption leaked into +`flamegpu::detail::gpu::*`, and the wave32 result did not drift from June +(1069/64/0 at `6487086` vs 1070/64/0 here; the extra pass is his new +`getDeviceName` / `getDeviceNames` test, same delta gfx90a saw). + +```bash +./projects/FLAMEGPU2/src/build/bin/Release/boids_bruteforce --steps 10 -v +# Total Processing time: 0.058868 s, exit 0 +``` + +Bit-for-bit the same total step time as the June run on this host (0.058773 s then), +which is a useful negative result on its own: his rewrite cost nothing measurable. + +### Consequence for the posted comment + +The comment tells ptheywood that gfx1100 was tested on "an earlier revision of the +branch". That is now understated -- gfx1100 matches gfx90a at `5e42a64`. Offering +the correction is a person's call and a second upstream comment; it is not urgent +and probably best folded into whatever reply his answer prompts. + +### Stage still needs a person + +Unchanged from the previous round and repeated here so it is not lost: the stage is +`porting` with the lock released, and `porting` exits only to `ported` or +`delta-ported`. Both are false -- `moat-port` was not touched and there is nothing +new for a reviewer. So the selector will keep dispatching a porter to this project +on every host, and each one will find no porting work. This round did not write a +stage. Two ways out, both a person's: `set-hold` until ptheywood answers, which is +exactly what "parked on someone else's reply" means and touches no state; or a +repair back to `review-passed`, where the gfx90a round found it. diff --git a/projects/FLAMEGPU2/plan.md b/projects/FLAMEGPU2/plan.md new file mode 100644 index 00000000..f057483e --- /dev/null +++ b/projects/FLAMEGPU2/plan.md @@ -0,0 +1,276 @@ +# Plan: FLAMEGPU2 + +## Project + +- **Name:** FLAMEGPU2 +- **Upstream:** https://github.com/FLAMEGPU/FLAMEGPU2 +- **Default branch:** main +- **Description:** GPU-accelerated agent-based modeling framework for CUDA C++ and Python + +## Existing AMD support + +**None found.** No AMD/ROCm/HIP support exists upstream or in any fork. + +- `grep -rniE 'amd|rocm|hip|gfx[0-9]' README* docs/` returned no matches +- Web search for "FLAMEGPU2 ROCm", "FLAME GPU 2 AMD GPU", "FLAMEGPU2 HIP" found no existing ports +- No forks with rocm/hip/amd in the name found via `gh api repos/FLAMEGPU/FLAMEGPU2/forks` +- No ROCm/HIP issues or PRs exist in the upstream repository +- The project's documentation explicitly requires "CUDA 12.0+" and an "NVIDIA GPU with Compute Capability >= 3.5" + +**Decision:** Proceed with a from-scratch ROCm/HIP port -- this is valuable new AMD support for a widely-used agent-based modeling framework. + +**NOTE:** This project presents a MAJOR porting complexity due to its use of **NVIDIA Jitify + nvrtc** for runtime kernel compilation (user-defined agent behavior is compiled at runtime). See Risk List below. + +## Build classification + +**Pure CMake** (Strategy A applies) + +Evidence: +- `CMakeLists.txt:45-48`: `check_language(CUDA)` + `enable_language(CUDA)` -- no PyTorch dependency +- No `find_package(Torch)`, no `torch.utils.cpp_extension`, no CUDAExtension +- Project builds as a static library + executables with CMake's native CUDA language support +- Test suite uses GoogleTest, not pytest + +## Port strategy + +**Strategy A (pure CMake, compat-header approach) with significant Jitify/RTC challenges** + +Rationale: This is a pure CMake project with CUDA sources. However, the standard Strategy A approach faces a major obstacle: FLAMEGPU2's core feature is **runtime compilation** of user-defined agent functions via Jitify2 + nvrtc. + +The port has two major work items: +1. **AOT path:** The library and examples build offline -- standard Strategy A with a cuda_to_hip.h compat header and `LANGUAGE HIP` marking. +2. **RTC (Runtime Compilation) path:** Jitify uses NVRTC (NVIDIA Runtime Compiler) which has no direct ROCm equivalent. This requires replacing Jitify with hiprtc (HIP Runtime Compiler) or comgr (Code Object Manager). + +### Proposed approach for RTC: + +ROCm provides **hiprtc** (HIP Runtime Compilation) as the equivalent of nvrtc. The Jitify2 library is NVIDIA-specific and will not work. Options: + +1. **Replace Jitify with hiprtc directly:** Rewrite JitifyCache.cu to use hiprtc APIs (`hiprtcCreateProgram`, `hiprtcCompileProgram`, `hiprtcGetCode`, etc.) on HIP. This is substantial work but cleanest. + +2. **Conditional compilation:** Keep Jitify for CUDA, implement a parallel hiprtc path for HIP behind `#if defined(USE_HIP)`. + +3. **Alternative: pre-compilation of common kernels:** If RTC is used only for performance (not user-defined code), pre-compile kernels AOT. However, FLAMEGPU2's RTC is a core feature for user-defined agent behaviors -- users write agent functions in C++/CUDA that are compiled at runtime. + +**Recommendation:** Implement a dual-path JitifyCache that uses Jitify on CUDA and hiprtc on HIP. This is significant work but preserves the framework's core functionality. + +## CUDA surface inventory + +### Kernels and device functions +- ~50+ `__global__` kernel definitions across src/flamegpu/ +- Many `__device__` helper functions +- `__host__ __device__` combined functions + +### CUDA Runtime API +- ~546 uses of cudaStream/cudaEvent/cudaMalloc/cudaMemcpy/cudaDevice APIs +- Standard memory management: cudaMalloc, cudaFree, cudaMemcpy, cudaMemcpyAsync +- Stream management: cudaStream_t, cudaStreamCreate, cudaStreamSynchronize +- Device queries: cudaGetDeviceProperties, cudaDeviceGetAttribute + +### cuRAND +- Heavy usage (~80 references) +- `curandStatePhilox4_32_10_t` / `curandStateMRG32k3a_t` / `curandStateXORWOW_t` +- `curand_init`, `curand_uniform`, etc. +- File: `include/flamegpu/detail/curand.cuh`, `src/flamegpu/simulation/detail/RandomManager.cu` +- **ROCm equivalent:** hipRAND (rocRAND) -- mostly 1:1 API mapping + +### CUB +- ~55 uses of `cub::` namespace +- `cub::DeviceScan` (18 uses) +- `cub::DeviceRadixSort` (10 uses) +- **ROCm equivalent:** hipCUB (rocPRIM underneath) -- API compatible + +### Thrust +- ~32 uses of `thrust::` namespace +- `thrust::device_vector`, sorting, etc. +- **ROCm equivalent:** rocThrust -- drop-in compatible (same headers) + +### NVRTC / Jitify (CRITICAL) +- Jitify2 library for runtime kernel compilation +- Direct nvrtc.h includes and API calls +- NVRTC functions: `nvrtcCreateProgram`, `nvrtcCompileProgram`, `nvrtcGetNumSupportedArchs`, etc. +- Files: `src/flamegpu/detail/JitifyCache.cu`, `src/flamegpu/detail/compute_capability.cu` +- **ROCm equivalent:** hiprtc -- similar API but significant rewrite needed + +### Compute Capability +- `cudaDevAttrComputeCapabilityMajor/Minor` +- `__CUDA_ARCH_LIST__` macro usage +- Files: `src/flamegpu/detail/compute_capability.cu` +- **Risk:** CC numbers collide between NVIDIA and AMD; need HIP-aware detection + +### `__ldg` intrinsic +- 14 uses for cached reads +- **ROCm equivalent:** `__ldg` is available in HIP (same spelling) + +### Textures +- `cudaBindTexture` references are commented out (not active) +- No active texture object usage found +- **Status:** Not a concern + +### Synchronization +- `__syncthreads` (2 uses) +- No `__syncwarp` or `__ballot` usage found +- **Status:** No warp intrinsic concerns + +## Risk list + +### HIGH RISK + +1. **Runtime Compilation (Jitify/nvrtc -> hiprtc):** FLAMEGPU2's core feature is runtime compilation of user-defined agent functions. Jitify is NVIDIA-specific and must be replaced with hiprtc-based implementation. This is substantial work (~500-1000 lines to rewrite JitifyCache.cu). + +2. **Compute Capability collision:** The project uses compute capability for kernel selection and nvrtc architecture flags. AMD's hipDeviceAttributeComputeCapabilityMajor returns values that collide with NVIDIA (gfx90a reports cc=9 like Hopper). Need to guard arch selection with `__HIP_PLATFORM_AMD__`. + +### MEDIUM RISK + +3. **cuRAND state types:** The project allows configurable cuRAND generators (Philox, MRG32k3a, XORWOW). hipRAND supports these but state struct names differ slightly; verify compatibility. + +4. **CUB DeviceRadixSort begin_bit:** PORTING_GUIDE notes hipCUB DeviceRadixSort with nonzero begin_bit may not sort correctly. Check if FLAMEGPU2 uses this. + +5. **Warp size (wave64 vs wave32):** No explicit warp intrinsics found in the codebase, but the RTC-compiled user code could contain them. Document that user agent functions should use `warpSize` not literal 32. + +### LOW RISK + +6. **__ldg intrinsic:** Available in HIP; no change needed. + +7. **Thrust/CUB:** rocThrust and hipCUB are drop-in replacements. + +## File-by-file change list + +### Core porting (Strategy A) + +1. **New file: `include/flamegpu/detail/cuda_to_hip.h`** + - Compat header with CUDA->HIP aliases + - Include hipRAND mappings + - Include hipCUB mappings + +2. **`CMakeLists.txt`** (root) + - Add `option(USE_HIP "Build with HIP for AMD GPUs" OFF)` + - Gate `enable_language(CUDA)` vs `enable_language(HIP)` + - Set `CMAKE_HIP_ARCHITECTURES` default to gfx90a + +3. **`src/CMakeLists.txt`** + - Add HIP language handling + - Mark CUDA sources as `LANGUAGE HIP` when USE_HIP + - Link hiprtc instead of CUDA::nvrtc + - Link hipCUB instead of CUB + - Link hipRAND instead of cuRAND + +4. **`include/flamegpu/detail/curand.cuh`** + - Add `#if defined(USE_HIP)` path for hipRAND includes + - Map curandState types to hiprand equivalents + +5. **`src/flamegpu/detail/compute_capability.cu`** + - Replace nvrtc calls with hiprtc equivalents under USE_HIP + - Replace cudaGetDeviceCount/Properties with hip equivalents + - Handle AMD arch detection instead of compute capability + +### Jitify replacement (major work) + +6. **`src/flamegpu/detail/JitifyCache.cu`** (MAJOR REWRITE for HIP) + - Replace Jitify2 with hiprtc on HIP path + - `jitify2::Program` -> `hiprtcCreateProgram` + - `program->preprocess()` -> custom preprocessing + - `preprocessed_program->compile()` -> `hiprtcCompileProgram` + - `compiled_program->link()` -> `hiprtcGetCode` + `hipModuleLoadData` + - Cache management similar to current implementation + +7. **`include/flamegpu/detail/JitifyCache.h`** + - Abstract interface to support both Jitify and hiprtc backends + +8. **`cmake/dependencies/Jitify.cmake`** + - Gate Jitify fetch on NOT USE_HIP + +9. **`cmake/common.cmake`** + - Gate CUDA::nvrtc requirement on NOT USE_HIP + - Add hip::hiprtc requirement for USE_HIP + +### Tests + +10. **`tests/CMakeLists.txt`** + - Add HIP language support + - Mark test .cu files as LANGUAGE HIP + +11. **Test files (78 .cu files)** + - Should work with compat header; may need minor fixes for hipRAND state types + +## Build commands + +### Configure (gfx90a) +```bash +cmake -S . -B build \ + -DUSE_HIP=ON \ + -DCMAKE_HIP_ARCHITECTURES=gfx90a \ + -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ + -DFLAMEGPU_BUILD_TESTS=ON \ + -DFLAMEGPU_BUILD_ALL_EXAMPLES=ON \ + -DCMAKE_BUILD_TYPE=Release +``` + +### Build +```bash +cmake --build build -j$(nproc) +``` + +## Test plan + +### GPU tests (must pass) +```bash +# Run full test suite +cd build && ctest -VV + +# Or run directly +./bin/Release/tests --gtest_filter="*" +``` + +### Individual example validation +```bash +# Run boids example +./bin/Release/boids_spatial3D + +# Run circles example +./bin/Release/circles_spatial3D +``` + +### RTC validation (critical) +```bash +# RTC examples compile and run user-defined agent functions +./bin/Release/rtc_boids_bruteforce +./bin/Release/rtc_boids_spatial3D +``` + +### Non-GPU tests +- The model/description tests are largely host-side +- Exception tests should pass + +### Validation criteria +- All 78+ test files pass +- All examples execute without error +- RTC agent functions compile and execute correctly +- Random number generation produces valid distributions + +## Open questions + +1. **Jitify replacement scope:** Is a full hiprtc implementation required, or can RTC be optional on AMD? The framework's core value proposition is runtime-compiled agent functions, so disabling RTC would significantly limit the AMD port's usefulness. + +2. **hiprtc API parity:** Are all nvrtc features used (preprocessing, architecture targeting, caching) available in hiprtc? Initial research suggests hiprtc has similar but not identical capabilities. + +3. **User agent function compatibility:** User-written agent functions (in CUDA syntax) would need to be compatible with HIP. This means: + - Users should not use NVIDIA-specific intrinsics + - Warp primitives should use `warpSize` not 32 + - The documentation should note AMD compatibility requirements + +4. **Python bindings (pyflamegpu):** The SWIG Python bindings link against nvrtc. The HIP port would need to link hiprtc and handle the differences. + +5. **Visualisation support:** The optional FLAMEGPU2 visualiser uses OpenGL. This should work on AMD but needs verification. + +6. **MPI ensemble support:** The MPI-parallel ensemble feature should be architecture-agnostic but needs testing. + +## Summary + +FLAMEGPU2 is a significant porting effort primarily due to its **runtime compilation** architecture. The AOT code (library, examples, tests) follows standard Strategy A patterns. The RTC subsystem requires replacing NVIDIA's Jitify/nvrtc with AMD's hiprtc. + +Estimated effort: +- AOT porting: Medium (~1-2 days) +- hiprtc JitifyCache replacement: High (~3-5 days) +- Testing and validation: Medium (~1-2 days) + +The port adds genuine value: FLAMEGPU2 is a mature, well-documented framework with academic and research users who would benefit from AMD GPU support. diff --git a/projects/FLAMEGPU2/stats.jsonl b/projects/FLAMEGPU2/stats.jsonl new file mode 100644 index 00000000..1293470a --- /dev/null +++ b/projects/FLAMEGPU2/stats.jsonl @@ -0,0 +1,45 @@ +{"kind": "tokens", "ts": "2026-06-05T01:20:55Z", "tokens": 82876, "source": "planner"} +{"kind":"phase","ts":"2026-06-05T05:08:07Z","phase":"compile","seconds":0.141,"exit":1,"cmd":"cmake --build . --target flamegpu boids_bruteforce tests -j128"} +{"kind": "tokens", "ts": "2026-06-05T05:16:38Z", "tokens": 95401, "source": "porter"} +{"kind": "tokens", "ts": "2026-06-05T05:19:37Z", "tokens": 97850, "source": "reviewer"} +{"kind":"phase","ts":"2026-06-05T05:23:46Z","phase":"compile","seconds":220.154,"exit":0,"cmd":"cmake --build /var/lib/jenkins/moat/projects/FLAMEGPU2/src/build --target flamegpu boids_bruteforce tests -j128"} +{"kind": "tokens", "ts": "2026-06-05T05:24:42Z", "tokens": 27607, "source": "porter"} +{"kind": "tokens", "ts": "2026-06-05T05:26:52Z", "tokens": 60633, "source": "reviewer"} +{"kind":"session","ts":"2026-06-05T05:27:18Z","epoch":1780637238.422304555,"event":"start","platform":"linux-gfx90a"} +{"kind":"phase","ts":"2026-06-05T05:27:29Z","phase":"linux-gfx90a","seconds":0.003,"exit":127,"cmd":"compile -- cmake -S /var/lib/jenkins/moat/projects/FLAMEGPU2/src -B /var/lib/jenkins/moat/projects/FLAMEGPU2/src/build -DCMAKE_BUILD_TYPE=Release -DFLAMEGPU_GPU=HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang -DFLAMEGPU_BUILD_TESTS=ON"} +{"kind":"phase","ts":"2026-06-05T05:28:00Z","phase":"compile","seconds":23.891,"exit":0,"cmd":"cmake -S /var/lib/jenkins/moat/projects/FLAMEGPU2/src -B /var/lib/jenkins/moat/projects/FLAMEGPU2/src/build -DCMAKE_BUILD_TYPE=Release -DFLAMEGPU_GPU=HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang -DFLAMEGPU_BUILD_TESTS=ON"} +{"kind":"phase","ts":"2026-06-05T05:31:50Z","phase":"compile","seconds":224.294,"exit":0,"cmd":"cmake --build /var/lib/jenkins/moat/projects/FLAMEGPU2/src/build --target flamegpu boids_bruteforce tests -j128"} +{"kind":"phase","ts":"2026-06-05T05:31:55Z","phase":"test","seconds":0.003,"exit":127,"cmd":"bin/Release/tests --gtest_output=xml:test-results.xml"} +{"kind":"phase","ts":"2026-06-05T05:32:37Z","phase":"test","seconds":36.316,"exit":0,"cmd":"/var/lib/jenkins/moat/projects/FLAMEGPU2/src/build/bin/Release/tests --gtest_output=xml:/var/lib/jenkins/moat/projects/FLAMEGPU2/src/build/test-results.xml"} +{"kind":"session","ts":"2026-06-05T05:33:08Z","epoch":1780637588.286534020,"event":"end","platform":"linux-gfx90a"} +{"kind":"session","ts":"2026-06-05T05:35:12Z","epoch":1780637712.994999219,"event":"start","platform":"linux-gfx1100"} +{"kind":"phase","ts":"2026-06-05T05:35:34Z","phase":"linux-gfx1100","seconds":0.003,"exit":127,"cmd":"compile -- bash -c cd projects/FLAMEGPU2/src && cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DFLAMEGPU_GPU=HIP -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang -DFLAMEGPU_BUILD_TESTS=ON 2>&1 | tee /tmp/flamegpu2-configure.log && cmake --build build --target flamegpu boids_bruteforce tests -j$(nproc) 2>&1 | tee /tmp/flamegpu2-build.log"} +{"kind":"phase","ts":"2026-06-05T05:38:47Z","phase":"compile","seconds":185.294,"exit":0,"cmd":"bash -c cd projects/FLAMEGPU2/src && cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DFLAMEGPU_GPU=HIP -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang -DFLAMEGPU_BUILD_TESTS=ON 2>&1 | tee /tmp/flamegpu2-configure.log && cmake --build build --target flamegpu boids_bruteforce tests -j$(nproc) 2>&1 | tee /tmp/flamegpu2-build.log"} +{"kind":"phase","ts":"2026-06-05T05:39:30Z","phase":"test","seconds":32.657,"exit":0,"cmd":"bash -c cd projects/FLAMEGPU2/src/build && ./bin/Release/tests --gtest_output=xml:test_results.xml 2>&1 | tee /tmp/flamegpu2-tests.log"} +{"kind":"session","ts":"2026-06-05T05:39:53Z","epoch":1780637993.136353575,"event":"end","platform":"linux-gfx1100"} +{"kind": "tokens", "ts": "2026-06-05T05:41:03Z", "tokens": 42216, "source": "validator (gfx1100)"} +{"kind":"phase","ts":"2026-06-08T00:00:53Z","phase":"compile","seconds":0.061,"exit":1,"cmd":"cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DFLAMEGPU_GPU=HIP -DCMAKE_HIP_ARCHITECTURES=gfx1201 -DCMAKE_CXX_COMPILER=B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel/lib/llvm/bin/clang++.exe -DCMAKE_C_COMPILER=B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel/lib/llvm/bin/clang.exe -DCMAKE_PREFIX_PATH=B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel -DFLAMEGPU_BUILD_TESTS=ON"} +{"kind":"phase","ts":"2026-06-08T00:07:05Z","phase":"compile","seconds":25.305,"exit":1,"cmd":"cmake --build B:/develop/moat/projects/FLAMEGPU2/src/build --target flamegpu boids_bruteforce tests -j32"} +{"kind":"phase","ts":"2026-06-08T00:12:07Z","phase":"compile","seconds":73.355,"exit":1,"cmd":"cmake --build B:/develop/moat/projects/FLAMEGPU2/src/build --target boids_bruteforce tests -j32"} +{"kind":"phase","ts":"2026-06-08T00:25:12Z","phase":"test","seconds":36.276,"exit":127,"cmd":"B:/develop/moat/projects/FLAMEGPU2/src/build/bin/Release/tests.exe"} +{"kind": "tokens", "ts": "2026-06-08T00:59:08Z", "tokens": 34472, "source": "validator"} +{"kind": "tokens", "ts": "2026-06-15T12:15:40Z", "tokens": 51286, "source": "pr-prep"} +{"kind":"session","ts":"2026-06-20T00:04:59Z","epoch":1781913899.076145400,"event":"start","platform":"windows-gfx1101"} +{"kind":"phase","ts":"2026-06-20T00:05:23Z","phase":"compile","seconds":14.680,"exit":0,"cmd":"cmake -S B:/develop/moat/projects/FLAMEGPU2/src -B B:/develop/moat/projects/FLAMEGPU2/src/build_gfx1101 -G Ninja -DCMAKE_BUILD_TYPE=Release -DFLAMEGPU_GPU=HIP -DCMAKE_HIP_ARCHITECTURES=gfx1101 -DCMAKE_CXX_COMPILER=B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel/lib/llvm/bin/clang++.exe -DCMAKE_C_COMPILER=B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel/lib/llvm/bin/clang.exe -DCMAKE_PREFIX_PATH=B:/develop/TheRock/external-builds/pytorch/.venv/Lib/site-packages/_rocm_sdk_devel -DFLAMEGPU_BUILD_TESTS=ON -DFETCHCONTENT_SOURCE_DIR_NLOHMANN_JSON=B:/develop/moat/agent_space/flamegpu2-deps/nlohmann_json-src -DFETCHCONTENT_SOURCE_DIR_TINYXML2=B:/develop/moat/agent_space/flamegpu2-deps/tinyxml2-wrapper/tinyxml2 -DFETCHCONTENT_SOURCE_DIR_GOOGLETEST=B:/develop/moat/agent_space/flamegpu2-deps/googletest-src"} +{"kind":"phase","ts":"2026-06-20T00:07:29Z","phase":"compile","seconds":115.820,"exit":0,"cmd":"cmake --build B:/develop/moat/projects/FLAMEGPU2/src/build_gfx1101 --target flamegpu boids_bruteforce tests -j64"} +{"kind":"phase","ts":"2026-06-20T00:08:34Z","phase":"test","seconds":36.075,"exit":127,"cmd":"env HIP_VISIBLE_DEVICES=1 B:/develop/moat/projects/FLAMEGPU2/src/build_gfx1101/bin/Release/tests.exe"} +{"kind":"phase","ts":"2026-06-20T00:09:18Z","phase":"test","seconds":35.859,"exit":127,"cmd":"env HIP_VISIBLE_DEVICES=1 B:/develop/moat/projects/FLAMEGPU2/src/build_gfx1101/bin/Release/tests.exe"} +{"kind":"phase","ts":"2026-06-20T00:19:54Z","phase":"test","seconds":2.444,"exit":0,"cmd":"env HIP_VISIBLE_DEVICES=1 B:/develop/moat/projects/FLAMEGPU2/src/build_gfx1101/bin/Release/tests.exe --gtest_filter=GPUTest*:TestCUDASubAgent*:DeviceAPITest*:HostFunctionTest*:TestMessage_BruteForce*:TestMessage_Array*:TestMessage_Spatial*"} +{"kind":"session","ts":"2026-06-20T00:21:57Z","epoch":1781914917.755236900,"event":"end","platform":"windows-gfx1101"} +{"kind": "tokens", "ts": "2026-06-20T00:24:23Z", "tokens": 85159, "source": "validator (gfx1101)"} +{"kind":"session","ts":"2026-08-20T00:11:56Z","epoch":1787184716.583404356,"event":"start","platform":"linux-gfx90a"} +{"kind":"phase","ts":"2026-08-20T00:16:28Z","phase":"configure","seconds":16.241,"exit":0,"cmd":"cmake -S projects/FLAMEGPU2/src -B projects/FLAMEGPU2/src/build -DCMAKE_BUILD_TYPE=Release -DFLAMEGPU_BACKEND=HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_CXX_COMPILER=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/_rocm_sdk_devel/lib/llvm/bin/clang++ -DCMAKE_C_COMPILER=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/_rocm_sdk_devel/lib/llvm/bin/clang -DCMAKE_PREFIX_PATH=/opt/conda/envs/py_3.12/lib/python3.12/site-packages/_rocm_sdk_devel;/opt/conda/envs/py_3.12/lib/python3.12/site-packages/_rocm_sdk_core -DFLAMEGPU_BUILD_TESTS=ON"} +{"kind":"phase","ts":"2026-08-20T00:19:30Z","phase":"compile","seconds":175.731,"exit":0,"cmd":"cmake --build projects/FLAMEGPU2/src/build --target flamegpu boids_bruteforce tests -j 96"} +{"kind":"phase","ts":"2026-08-20T00:20:55Z","phase":"test","seconds":39.988,"exit":0,"cmd":"env HIP_VISIBLE_DEVICES=3 ./projects/FLAMEGPU2/src/build/bin/Release/tests"} +{"kind":"session","ts":"2026-08-20T00:26:36Z","epoch":1787185596.305094202,"event":"end","platform":"linux-gfx90a"} +{"kind":"session","ts":"2026-08-20T05:24:02Z","epoch":1787203442.184920898,"event":"start","platform":"linux-gfx1100"} +{"kind":"phase","ts":"2026-08-20T05:24:39Z","phase":"compile","seconds":12.435,"exit":0,"cmd":"cmake -S projects/FLAMEGPU2/src -B projects/FLAMEGPU2/src/build -DCMAKE_BUILD_TYPE=Release -DFLAMEGPU_BACKEND=HIP -DCMAKE_HIP_ARCHITECTURES=gfx1100 -DCMAKE_CXX_COMPILER=/opt/rocm/lib/llvm/bin/clang++ -DCMAKE_C_COMPILER=/opt/rocm/lib/llvm/bin/clang -DFLAMEGPU_BUILD_TESTS=ON"} +{"kind":"phase","ts":"2026-08-20T05:27:49Z","phase":"compile","seconds":179.292,"exit":0,"cmd":"cmake --build projects/FLAMEGPU2/src/build --target flamegpu boids_bruteforce tests -j 32"} +{"kind":"phase","ts":"2026-08-20T05:28:34Z","phase":"test","seconds":29.836,"exit":0,"cmd":"./projects/FLAMEGPU2/src/build/bin/Release/tests"} +{"kind":"phase","ts":"2026-08-20T05:28:48Z","phase":"test","seconds":0.855,"exit":0,"cmd":"./projects/FLAMEGPU2/src/build/bin/Release/boids_bruteforce --steps 10 -v"} +{"kind":"session","ts":"2026-08-20T05:31:19Z","epoch":1787203879.374487956,"event":"end","platform":"linux-gfx1100"} diff --git a/projects/FLAMEGPU2/status.json b/projects/FLAMEGPU2/status.json new file mode 100644 index 00000000..f8a76c3e --- /dev/null +++ b/projects/FLAMEGPU2/status.json @@ -0,0 +1,137 @@ +{ + "schema_version": 3, + "name": "FLAMEGPU2", + "upstream_url": "https://github.com/FLAMEGPU/FLAMEGPU2", + "fork_url": "https://github.com/AMD-Ecosystem/FLAMEGPU2", + "fork_default_branch": "master", + "priority": 0.0, + "ext_type": "cmake", + "adopted_at": "2026-06-05T01:16:31Z", + "updated_at": "2026-08-20T05:37:54Z", + "head_sha": "a290861", + "depends_on": [], + "porting": null, + "waivers": {}, + "license_clearance": { + "approved_by": "jeffdaily", + "at": "2026-08-06T02:33:28Z", + "tier": 3, + "note": "carries the org review that cleared this specific list on 2026-08-06; covers this project only and sets no precedent for its tier" + }, + "license_spdx": "AGPL-3.0", + "upstream_repo_id": 34064755, + "stage": "porting", + "platforms": { + "linux-gfx90a": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "a290861", + "started_at": "2026-06-05T05:06:43Z", + "completed_at": "2026-06-08T00:38:47Z", + "updated_at": "2026-08-20T00:11:42Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "last_agent": "validator", + "carry_forward": { + "to": "a290861", + "method": "binary-equiv", + "detail": "#include inside #ifdef _MSC_VER; linux clang does not define _MSC_VER, no compiled diff", + "at": "2026-06-08T00:38:47Z" + } + }, + "linux-gfx1100": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "a290861", + "started_at": null, + "completed_at": "2026-06-08T00:39:02Z", + "updated_at": "2026-06-08T00:39:02Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "last_agent": "validator", + "carry_forward": { + "to": "a290861", + "method": "binary-equiv", + "detail": "#include inside #ifdef _MSC_VER; linux clang does not define _MSC_VER, no compiled diff", + "at": "2026-06-08T00:39:02Z" + } + }, + "windows-gfx1101": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "a290861", + "started_at": null, + "completed_at": "2026-06-20T00:21:53Z", + "updated_at": "2026-06-20T00:21:53Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "last_agent": "validator" + }, + "windows-gfx1201": { + "state": "completed", + "blocked": false, + "blocked_reason": null, + "validated_sha": "a290861", + "started_at": null, + "completed_at": "2026-06-08T00:58:21Z", + "updated_at": "2026-06-08T00:58:21Z", + "stats": { + "tokens_total": 0, + "tokens_approx": true, + "wall_seconds": { + "thinking": 0, + "compile": 0, + "test": 0, + "misc": 0 + }, + "session_count": 0, + "first_session_at": null, + "last_session_at": null + }, + "last_agent": "validator" + } + }, + "porting_released": { + "arch": "linux-gfx90a", + "since": "2026-08-20T00:11:42Z", + "at": "2026-08-20T00:26:32Z" + }, + "on_hold": true, + "on_hold_reason": "Parked by Jeff Daily until ptheywood replies on upstream draft PR #1379 (amdgpu branch): waiting to learn whether he wants a PR against amdgpu or a cherry-pick of amdgpu-fixes @ 2ecf5a01. Port itself is complete on all four platforms at a290861." +}