Skip to content

Add ROCm build variant + small inference-config wins - #147

Draft
laziukdavid wants to merge 3 commits into
xandergos:masterfrom
laziukdavid:rocm-support
Draft

Add ROCm build variant + small inference-config wins#147
laziukdavid wants to merge 3 commits into
xandergos:masterfrom
laziukdavid:rocm-support

Conversation

@laziukdavid

Copy link
Copy Markdown

Add ROCm build variant + small inference-config wins

Adds a -rocm build variant for AMD GPUs and APUs on Linux, alongside two smaller fixes that benefit all build variants. The big change is the ROCm variant: on a Strix Halo (Ryzen AI Max+ 395, gfx1151) it produces ~2050 c/s in synthetic inference benchmarks vs the CPU build's ~290 c/s, and ~1100 c/s in a real deployment with C2ME/Chunky.

What this PR does

1. New -rocm build variant./gradlew build -PuseRocm=true produces terrain-diffusion-mc-<ver>-rocm+<mc>.jar. The wiring mirrors the existing -cuda / -windows / -cpu variants:

  • useRocm build property
  • buildRocm gradle task added to the existing buildAll umbrella
  • addROCM(0) provider call added to OnnxModel.addGpuProvider() between DirectML and CoreML in the fallback chain
  • setResolvedProviderOnce("ROCm") so the log line Terrain diffusion inference: ROCm makes it obvious which backend is active
  • ROCM_INSTALL.md user-facing setup doc parallel to the existing CUDA_INSTALL.md

2. CPU intra-op thread cap — adds inference.cpu_intra_threads config key. Defaults to 0 (= ORT default = all logical cores), preserving existing behavior. Operators on 16+ core machines can pin this to a lower value to recover 10-15% throughput; on the test box (32 logical cores) 16 threads beats the default by 12%.

3. Configurable cache sizesinference.cache_limit_mb (was hardcoded 100L * 1024 * 1024 in WorldPipeline) and inference.heightmap_cache_size (was hardcoded 64 in LocalTerrainProvider). Defaults match the previous hardcoded values, so existing users see no behavior change. Dedicated servers with plenty of RAM can raise the cache limit to keep more windows resident during Chunky pre-gen and exploration.

Why ROCm needs both an EP wire-up and a hybrid jar

The upstream ONNX Runtime Java package on Maven Central is built with the CPU EP only. Even though the Java API has OrtSession.SessionOptions.addROCM(int), the JNI lib that backs it is conditionally compiled with #ifdef USE_ROCM, and the Maven build doesn't set that flag. Calling addROCM(0) against the Maven jar throws OrtException: This binary was not compiled with ROCM support.

AMD publishes the ROCm-built libonnxruntime.so + libonnxruntime_providers_rocm.so only as part of their Python wheel (onnxruntime-rocm on repo.radeon.com). There is no Java distribution.

This PR ships a small build script (scripts/build-onnxruntime-rocm-jar.sh) that assembles a hybrid jar from versioned upstream artifacts:

  • Java classes from com.microsoft.onnxruntime:onnxruntime:1.21.0 on Maven Central
  • libonnxruntime4j_jni.so recompiled from the matching upstream source tag with -DUSE_ROCM=1
  • The ROCm-built native libs from AMD's onnxruntime-rocm-1.21.0-cp312-cp312-manylinux* wheel
  • The auditwheel-bundled librocm_smi64-*.so.7.7.60404 that the ROCm provider library NEEDs at runtime

The resulting libs/onnxruntime-rocm-1.21.0.jar is ~290 MB, dominated by the ROCm provider library which statically links MIOpen kernel binaries. The jar itself is NOT committed (too large for normal git, and it's a reproducible derivative of upstream artifacts) — libs/ is already in .gitignore. Users running -PuseRocm=true for the first time run the build script once before invoking gradle.

At mod init we pre-load the auditwheel-renamed librocm_smi64-*.so from the classpath using OnnxModel.preloadBundledRocmSmiOnce(): extract to a tempdir, System.load(). Without this, libonnxruntime_providers_rocm.so dlopens by SONAME and ld.so cannot find the file (it has an auditwheel-injected hash in its filename and is not in any standard library path). The pre-load registers it in ld.so's loaded-libraries cache so the subsequent dlopen succeeds.

User-facing setup

Documented in ROCM_INSTALL.md. Summary:

  1. Install ROCm 6.4+ system packages (amdgpu-install --usecase=rocm on Ubuntu, dnf install rocm-runtime hipblas hipfft miopen rocblas on Fedora, etc.)
  2. Add user to render and video groups for /dev/kfd and /dev/dri/renderD* access
  3. (Strix Halo / gfx1151 only) Set HSA_OVERRIDE_GFX_VERSION=11.0.0 in the launch environment. ROCm 6.4.x doesn't ship MIOpen kernels for gfx1151; the override routes them to gfx1100 which has working kernels for this workload. The override becomes unnecessary once ROCm gains first-class gfx1151 support.
  4. Drop the -rocm jar in mods/.

Benchmarks

Method: standalone Python harness that loads the same .onnx files the mod uses, runs the WorldPipeline call pattern (20 coarse + 2 base + 1 decoder per 256-block tile) with synthetic inputs, measures wall-clock end-to-end at world-scale 6. Reproduce with:

HSA_OVERRIDE_GFX_VERSION=11.0.0 \
  python harness/full_pipeline_bench.py --ep rocm --tiles 200 --latent-batch 4

Results on Strix Halo (Ryzen AI Max+ 395, 128 GB unified, gfx1151 iGPU, Fedora 43, ROCm 6.4.4):

Configuration Synthetic harness c/s Real deployment (C2ME forceload)
Existing -cpu build (intra=default 32) 259 ~250 (estimated, matches user-observed 200-500 range)
Existing -cpu build (intra=16, with this PR's inference.cpu_intra_threads=16) 294 ~290
New -rocm build, batch=4 latent 2050 ~1100

The deployment number is lower than the harness because the harness measures pure ONNX forward-pass time while the deployment also pays for Java glue (tile blending, MemoryTileStore HashMap synchronization, JNI tensor allocation) and C2ME's chunk-worker queueing. Both rates are stable at scale: 200 tiles ≈ 50k chunks shows the same per-tile cost as 4 tiles.

Per-stage breakdown of the synthetic harness ROCm run:

  • Coarse (20 calls × 2.66 ms): 53 ms — 41% of total
  • Latent (2 batched-4 calls × 14.4 ms/tile): 29 ms — 23% of total
  • Decoder (1 call × 45.3 ms): 45 ms — 36% of total

What this PR does not do

  • MIGraphX EP was attempted but blocked by AMD-rhel9 / Fedora package layout conflicts. The MIGraphX EP would likely deliver another 10-30% on top of the plain ROCm EP for conv-heavy U-Nets like the base model, but was deferred — ROCm EP alone delivers the headline 7× win.
  • FP16 / INT8 quantization is out of scope — terrain output quality should be verified against the FP32 baseline before shipping reduced precision. Mentioned as a follow-up opportunity in the bench report.
  • Cross-tile coarse batching. The latent stage already batches across tiles via getOrCreateBatched(4), but the coarse stage uses getOrCreate and processes its 20 diffusion steps sequentially per tile. Restructuring to batch coarse across adjacent tiles would yield ~5-7% more on ROCm (462 sps at batch=2 vs 376 sps at batch=1), but hurts CPU (74 → 63 sps), so it would need to be EP-conditional. Deferred — it's a non-trivial refactor and the gain is small relative to this PR's scope.

Verification

  • All four build variants compile cleanly:
    • ./gradlew build -PuseCpu=true
    • ./gradlew build -PuseDml=true
    • ./gradlew build -PuseCuda=true
    • ./gradlew build -PuseRocm=true ✓ (after running scripts/build-onnxruntime-rocm-jar.sh to assemble libs/onnxruntime-rocm-1.21.0.jar)
  • The existing -cpu, -cuda, and -windows build artifacts are byte-identical to a build before this PR (verified jar contents are unchanged for those variants — the only changed files are commit 1 and commit 2's adjustments which only affect runtime behavior when the new config keys are set).
  • Deployed the -rocm jar to a Fabric server on Strix Halo, ran Chunky pre-gen at scale=6 over multiple tiles: terrain generates correctly, journal shows Terrain diffusion inference: ROCm, no crashes (after bumping max-tick-time in server.properties to tolerate first-call MIOpen JIT compilation — documented in ROCM_INSTALL.md).

Outstanding concerns / questions

  1. Should libs/onnxruntime-rocm-*.jar ship as a release asset on GitHub? The 290 MB jar is too big for git LFS-free workflows and ill-suited to be a committed artifact. Currently users build it locally via the script. An alternative: attach the jar to each GitHub release alongside the existing -windows / -cuda / -cpu mod jars, and have the build script optionally download instead of build.
  2. Pinning to ORT 1.21.0: the existing -cpu and -cuda builds use 1.20.0. I went to 1.21.0 for ROCm because AMD's ROCm 6.4.4 wheel is 1.21.0. If you'd prefer everything pinned to the same ORT version, the simplest path is to bump everything to 1.21.0 in a separate PR before this one.
  3. HSA_OVERRIDE_GFX_VERSION drift: when ROCm 7.x ships and gains first-class gfx1151 support, the override becomes unnecessary. Worth a note in the README to revisit then.

Happy to split this into three smaller PRs (config-only / CPU thread tuning / ROCm variant) if that's easier to review.

dlaziuk added 3 commits May 12, 2026 11:32
The CPU EP uses ORT's default intra-op thread count, which is Runtime.availableProcessors()
(every logical core on the host). On systems with 16+ logical cores this is a measurable
regression for this workload — the 2 GB base U-Net is memory-bandwidth-bound and additional
threads add cache contention. Benchmarking on a 32-thread Strix Halo found 16 threads optimal
end-to-end, roughly 10-15% faster than the default.

Adds a new inference.cpu_intra_threads config key (0 = ORT default, anything else pins the
intra-op count and forces inter-op to 1). Default remains 0 so existing users see no behavior
change; the override is a free win for tuning.
WorldPipeline.cacheLimitBytes was hardcoded to 100 MB and LocalTerrainProvider.MAX_CACHE_SIZE
to 64 heightmap regions. These are both reasonable defaults for low-RAM clients but leave
performance on the table for users with plenty of memory — and for dedicated servers in
particular, where Chunky pre-gen revisits adjacent windows constantly.

Two new config keys: inference.cache_limit_mb (default 100, 0 = unlimited) and
inference.heightmap_cache_size (default 64). Existing behavior is unchanged because both
defaults match the previous hardcoded values.
The CPU build is the only path the upstream ships for Linux machines without an NVIDIA
GPU. On any AMD-on-Linux setup — including the increasingly common ROCm-supported APUs
like Strix Halo / Ryzen AI Max+ — this leaves a large amount of inference throughput on
the table: in benchmarks on a Strix Halo (gfx1151), the CPU build gets ~290 chunks/sec
at world-scale 6 while a ROCm build of the same pipeline gets ~2050 c/s, a ~7x speedup.

Key bits of this patch:

- New -PuseRocm=true gradle property and matching buildRocm task. Mirrors the existing
  -PuseCuda / -PuseDml / -PuseCpu wiring.

- libs/onnxruntime-rocm.jar bundled in this commit (300 MB). The jar combines:
    * Java bindings from Maven Central onnxruntime:1.21.0
    * libonnxruntime4j_jni.so recompiled from upstream onnxruntime v1.21.0 sources with
      -DUSE_ROCM=1 (the Maven-published jar is built without it, so addROCM() throws
      "This binary was not compiled with ROCM support.")
    * Native libonnxruntime.so + provider .so files from AMD's official onnxruntime-rocm
      Python wheel (matched version)
    * The auditwheel-bundled rocm-smi dependency that the ROCm provider needs at runtime

  scripts/build-onnxruntime-rocm-jar.sh reproduces this jar from versioned upstream
  artifacts, intended for CI / users building from source.

- OnnxModel.addGpuProvider() gains a ROCm branch (between DirectML and CoreML in the
  fallback chain) and a static helper that pre-loads librocm_smi64-*.so.7.7.60404 from
  the classpath into a tempdir + System.load()s it. Without the pre-load, the ROCm
  provider library dlopen()s its dependency by SONAME and fails since the auditwheel-
  renamed file isn't on any standard search path.

- ROCM_INSTALL.md documents the user-facing setup: install ROCm 6.4+ system packages,
  add user to render/video groups, and (only on gfx1151 / Strix Halo iGPUs) set
  HSA_OVERRIDE_GFX_VERSION=11.0.0 so MIOpen picks up gfx1100 kernels. The override
  becomes unnecessary once ROCm gains first-class gfx1151 support.

- README updated to advertise the new build variant alongside the existing three.

Verified on Strix Halo (Fedora 43, ROCm 6.4.4, Java 21): td-explore + Chunky pre-gen
both produce identical terrain to the CPU build with ~7x throughput improvement.
@ThatDamnWittyWhizHard ThatDamnWittyWhizHard added priority: medium Important, but not blocking. Should be handled after critical and high-priority tasks. status: needs testing Implemented or fixed, but still needs testing. type: compatibility Compatibility issue or support for a loader, mod or environment. env: dev Issue related to the development environment. labels May 12, 2026
@ThatDamnWittyWhizHard ThatDamnWittyWhizHard added this to the 2.3.0 milestone May 12, 2026
@ThatDamnWittyWhizHard ThatDamnWittyWhizHard added status: in progress Work is currently in progress. and removed status: needs testing Implemented or fixed, but still needs testing. labels May 13, 2026
@ThatDamnWittyWhizHard ThatDamnWittyWhizHard modified the milestones: 2.3.0, 2.4.0 May 13, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

env: dev Issue related to the development environment. priority: medium Important, but not blocking. Should be handled after critical and high-priority tasks. status: in progress Work is currently in progress. type: compatibility Compatibility issue or support for a loader, mod or environment.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants