moss-infer-hip: the native AMD (ROCm/HIP) program

.github/workflows/build-hip.yml

@@ -0,0 +1,97 @@
+# Reusable build of moss-infer-hip, the standalone ROCm program (hip/), called by ci.yml and release.yml. Linux x86-64
+# only, in AMD's ROCm 7.2.4 development container pinned by digest (ROCm's clang, the HIP runtime and headers), with
+# rocWMMA's headers and LLVM's dev tools from the container's ROCm 7.2.4 apt repository: one binary for gfx1100, gfx1200
+# and gfx1201, warnings as errors, the device check of every kernel source on every architecture (part of the build),
+# the checks and tools compiled, then the CPack archive and its program's RUNPATH, --version and --list-devices. Hosted
+# runners have no AMD GPU, so nothing runs a kernel here. No CUDA is involved.
+name: build-hip
+
+on:
+  workflow_call:
+    inputs:
+      retention-days:
+        description: Days to keep the uploaded package
+        type: number
+        default: 7
+
+permissions:
+  contents: read
+
+jobs:
+  hip:
+    name: linux-hip
+    runs-on: ubuntu-22.04
+    timeout-minutes: 120
+    # The ROCm the release builds with: its clang emits the WMMA counts check_device.sh expects. Pulled by the runner
+    # (about 1 GB compressed); no cache is needed beyond it.
+    container:
+      image: rocm/dev-ubuntu-22.04:7.2.4@sha256:a50a8547101ac9e7c35848ff609a56335214b789b468c9eb05dcd006a4a38c25
+    defaults:
+      run:
+        shell: bash
+    env:
+      ROCM_PATH: /opt/rocm
+      # The container's ROCm 7.2.4 apt repository: rocWMMA's headers, and the LLVM tools check_device.sh reads the device
+      # code with (llvm-cxxfilt, llvm-nm), which the container's compiler package leaves out.
+      ROCWMMA_VERSION: 2.2.0.70204-93~22.04
+      ROCM_LLVM_DEV_VERSION: 22.0.0.26084.70204-93~22.04
+    steps:
+      - uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1
+
+      # unzip: get-cmake extracts its Ninja download with it, and the image has none.
+      - name: rocWMMA headers, LLVM tools, zlib and unzip
+        run: |
+          apt-get update
+          apt-get install -y --no-install-recommends "rocwmma-dev=$ROCWMMA_VERSION" \
+            "rocm-llvm-dev=$ROCM_LLVM_DEV_VERSION" zlib1g-dev binutils unzip
+          apt-get clean
+          "$ROCM_PATH/lib/llvm/bin/clang++" --version
+          cat "$ROCM_PATH/.info/version"
+
+      - uses: lukka/get-cmake@fffaaafeea488556c2c12dad60690008bc1caacb # v4.4.2
+
+      - name: Configure
+        run: |
+          cmake -S hip -B build-hip -G Ninja -DCMAKE_BUILD_TYPE=Release "-DHIP_TARGET=gfx1100;gfx1200;gfx1201" \
+            -DROCM_PATH="$ROCM_PATH" -DMOSSINFER_STATIC_LIBSTDCXX=ON
+
+      # The build runs check_device.sh on every kernel source for every architecture (mossinfer_hip_device_check).
+      - name: Build
+        run: |
+          cmake --build build-hip --parallel 4 2>&1 | tee build-hip/build.log
+          test "${PIPESTATUS[0]}" = 0
+          checks=$(grep -c "kernels checked" build-hip/build.log)
+          echo "device checks: $checks"
+          test "$checks" -eq 12
+
+      - name: Package
+        run: |
+          cd build-hip
+          cpack
+          ls -l mossinfer-*-hip-*
+
+      # The archive's program: RUNPATH /opt/rocm/lib and nothing of the build tree, ROCm 7's HIP runtime, no libstdc++
+      # of the build's; it starts, and with no GPU in the container it lists no device and exits 0.
+      - name: Package check
+        run: |
+          mkdir pkg
+          tar -xzf build-hip/mossinfer-*-linux-hip-x86_64.tar.gz -C pkg
+          exe=$(echo pkg/mossinfer-*-linux-hip-x86_64/bin/moss-infer-hip)
+          readelf -d "$exe" | tee dynamic.txt
+          grep -Fq "Library runpath: [/opt/rocm/lib]" dynamic.txt
+          grep -Fq "Shared library: [libamdhip64.so.7]" dynamic.txt
+          if grep -Fq "Shared library: [libstdc++" dynamic.txt; then echo "libstdc++ is linked dynamically"; exit 1; fi
+          ls pkg/mossinfer-*-linux-hip-x86_64/share/doc/mossinfer/hip.md
+          version=$("$exe" --version)
+          echo "$version"
+          [[ "$version" == *"(kernels: gfx1100,gfx1200,gfx1201)" ]]
+          devices=$("$exe" --list-devices)
+          echo "$devices"
+          [ "$devices" = "no HIP devices" ]
+
+      - uses: actions/upload-artifact@cf430e030ddbb5b0abf93d22962f4752f3646cd9 # v7.0.2
+        with:
+          name: mossinfer-hip-linux
+          path: build-hip/mossinfer-*-linux-hip-x86_64.tar.gz
+          retention-days: ${{ inputs.retention-days }}
+          if-no-files-found: error

.github/workflows/ci.yml

@@ -32,3 +32,7 @@ jobs:
     uses: ./.github/workflows/build-vulkan.yml
     with:
       retention-days: 7
+  hip:
+    uses: ./.github/workflows/build-hip.yml
+    with:
+      retention-days: 7

.github/workflows/release.yml

@@ -1,8 +1,8 @@
 # Tag vX.Y.Z (a release) or vX.Y.Z-rcN, N >= 1 (a release candidate, published as a pre-release that does not become
 # "latest"): build and check all four standard flavours, the Windows CUDA 13 SM75 flavour and the Linux CUDA 12 SM70
-# flavour (build.yml) and moss-infer-vulkan for Linux and Windows (build-vulkan.yml), then publish
-# their archives (the full package and the -system-cuda variant of each CUDA flavour, one archive per Vulkan platform)
-# and a SHA256SUMS file as a GitHub release.
+# flavour (build.yml), moss-infer-vulkan for Linux and Windows (build-vulkan.yml) and moss-infer-hip for Linux
+# (build-hip.yml), then publish their archives (the full package and the -system-cuda variant of each CUDA flavour, one
+# archive per Vulkan platform, the HIP archive) and a SHA256SUMS file as a GitHub release.
 # X.Y.Z must equal the MI_VERSION_* of
 # include/mossinfer/mossinfer.h and an rc must come before vX.Y.Z (.github/scripts/release_version.py), the tagged
 # commit must be on main with a successful ci run (.github/scripts/release_gate.sh), and the release must not exist
@@ -99,9 +99,15 @@ jobs:
     with:
       retention-days: 30
 
+  hip:
+    needs: version
+    uses: ./.github/workflows/build-hip.yml
+    with:
+      retention-days: 30
+
   # Checks and gathers the release assets with a read-only token; on a dispatch its artifact is the dry run's result.
   package:
-    needs: [version, build, vulkan]
+    needs: [version, build, vulkan, hip]
     runs-on: ubuntu-22.04
     env:
       ARCHIVE_VERSION: ${{ needs.version.outputs.archive_version }}
@@ -119,7 +125,7 @@ jobs:
           for full in mossinfer-*.zip mossinfer-*.tar.gz; do
             [ -f "$full" ] || continue
             case "$full" in "mossinfer-$ARCHIVE_VERSION-"*) ;; *) echo "$full is not version $ARCHIVE_VERSION"; exit 1 ;; esac
-            case "$full" in *-system-cuda.* | *-vulkan-*) continue ;; esac
+            case "$full" in *-system-cuda.* | *-vulkan-* | *-hip-*) continue ;; esac
             case "$full" in *.tar.gz) ext="tar.gz"; base="${full%.tar.gz}" ;; *) ext="zip"; base="${full%.zip}" ;; esac
             test -f "$base-system-cuda.$ext" || { echo "missing $base-system-cuda.$ext"; exit 1; }
             n=$((n + 1))
@@ -135,6 +141,8 @@ jobs:
           for vk in "mossinfer-$ARCHIVE_VERSION-linux-vulkan-x86_64.tar.gz" "mossinfer-$ARCHIVE_VERSION-windows-vulkan-x86_64.zip"; do
             test -f "$vk" || { echo "missing Vulkan package $vk"; exit 1; }
           done
+          hip="mossinfer-$ARCHIVE_VERSION-linux-hip-x86_64.tar.gz"
+          test -f "$hip" || { echo "missing HIP package $hip"; exit 1; }
       - name: SHA256SUMS
         shell: bash
         working-directory: dist

README.md

@@ -4,7 +4,7 @@ A native C++/CUDA inference engine for the MOSS-Image text-to-image model (5B Di
 
 Models: MOSS-Image Base and SFT, in the official folders and common community formats; other models are refused ([supported models](#supported-models)).
 
-Version 2.1.2. Supported: Linux x86-64 and Windows x64 (MSVC), each with a CUDA 12 and a CUDA 13 build, as a full package and a `-system-cuda` package. Turing (compute capability 7.5, e.g. RTX 2080) is supported by builds configured with `-DMOSSINFER_SM75=ON` (the `windows-msvc-cuda13-sm75` preset packages it): fp16 tensor-core GEMMs replace the bf16 ones, INT8 and INT4 ConvRot files run the INT8 GEMMs, the exact attention and the VAE run sm_75 kernels, all three quality tiers work, and an unset attention resolves to `exact`. See [docs/sm75.md](docs/sm75.md) for what runs, the measured speed on an RTX 2080 Ti and the remaining gaps. Volta (Tesla V100, sm_70) is supported by the CUDA 12 `linux-cuda12-sm70` preset; see [docs/sm70.md](docs/sm70.md) for its kernels, validation and performance. Other GPUs (AMD, Intel, older NVIDIA) can run txt2img at the reference tier through `moss-infer-vulkan`, a separate and slower Vulkan program in its own package ([docs/vulkan.md](docs/vulkan.md)). The full support matrix is in [docs/integration.md](docs/integration.md#support-status).
+Version 2.1.2. Supported: Linux x86-64 and Windows x64 (MSVC), each with a CUDA 12 and a CUDA 13 build, as a full package and a `-system-cuda` package. Turing (compute capability 7.5, e.g. RTX 2080) is supported by builds configured with `-DMOSSINFER_SM75=ON` (the `windows-msvc-cuda13-sm75` preset packages it): fp16 tensor-core GEMMs replace the bf16 ones, INT8 and INT4 ConvRot files run the INT8 GEMMs, the exact attention and the VAE run sm_75 kernels, all three quality tiers work, and an unset attention resolves to `exact`. See [docs/sm75.md](docs/sm75.md) for what runs, the measured speed on an RTX 2080 Ti and the remaining gaps. Volta (Tesla V100, sm_70) is supported by the CUDA 12 `linux-cuda12-sm70` preset; see [docs/sm70.md](docs/sm70.md) for its kernels, validation and performance. Other GPUs (AMD, Intel, older NVIDIA) can run txt2img at the reference tier through `moss-infer-vulkan`, a separate and slower Vulkan program in its own package ([docs/vulkan.md](docs/vulkan.md)); on Linux, `moss-infer-hip`, a separate ROCm 7 program, runs it on AMD GPUs: built for gfx1100, gfx1200 and gfx1201 (Radeon RX 7900 and RX 9000), validated on gfx1200 (RX 9060 XT) ([docs/hip.md](docs/hip.md)). The full support matrix is in [docs/integration.md](docs/integration.md#support-status).
 
 The caller picks the quality tier: `reference` (bf16, exact), `high` (the loaded DiT's GEMMs with 8-bit attention) or `turbo` (a 4-bit DiT). An INT8 DiT with `high` is the everyday choice: 2.5 times the bf16 speed on an RTX 5090, in 40 % less VRAM ([performance](docs/integration.md#performance)). Upgrading from 1.x: [migration](docs/integration.md#migrating-from-1x).
 
@@ -31,6 +31,7 @@ Lowest VRAM against normal use at 1024×1024. `min_memory` streams the DiT and c
 - [docs/sm70.md](docs/sm70.md): Volta (Tesla V100) support: CUDA 12 build, exact INT8 accumulation, validation and performance.
 - [docs/sm75.md](docs/sm75.md): Turing (compute capability 7.5) support: build and package, what runs, measured speed on an RTX 2080 Ti, limitations.
 - [docs/vulkan.md](docs/vulkan.md): `moss-infer-vulkan`, the Vulkan fallback for GPUs the CUDA engine does not support.
+- [docs/hip.md](docs/hip.md): `moss-infer-hip`, the ROCm program for AMD GPUs on Linux: built for gfx1100, gfx1200 and gfx1201, validated on gfx1200 (RX 9060 XT).
 - [docs/c-api.md](docs/c-api.md): the C library, `include/mossinfer/mossinfer.h`.
 - [docs/server-protocol.md](docs/server-protocol.md): `moss-server`, the JSON-lines sidecar process.
 

THIRD_PARTY_NOTICES.md

@@ -18,6 +18,8 @@ the upstream licence file; its `VENDORED.md` lists what was kept and what was ch
 | LyCORIS | Apache-2.0 | Copyright (c) KohakuBlueleaf and contributors | https://github.com/KohakuBlueleaf/LyCORIS | `src/models/lora/` follows its LoHa/LoKr weight definitions and scaling |
 | stb_image | MIT or public domain (Unlicense), at your choice | Copyright (c) 2017 Sean Barrett | https://github.com/nothings/stb (`013ac3beddff3dbffafd5177e7972067cd2b5083`) | `third_party/stb/` (`LICENSE`). Linked into the engine for PNG/JPEG decoding of init images and masks (`src/pipeline/pipeline_image.cpp`), and used by the tests |
 | ggml / llama.cpp (GGUF container, ggml block types) | MIT | Copyright (c) 2023-2026 The ggml authors | https://github.com/ggml-org/ggml, https://github.com/ggml-org/llama.cpp (`gguf-py`) | `src/weights/gguf.*` reads the GGUF format; its dequantization of the Q8_0, Q4_0, Q4_1, Q5_0, Q5_1, Q4_K, Q5_K and Q6_K blocks is derived from ggml's block layouts and dequantization (`dequantize_row_*`, their operation order). `tests/data/gguf/` are outputs of the `gguf` Python package. Licence text below |
+| rocWMMA (C++ headers) | MIT | Copyright (C) 2016-2025 Advanced Micro Devices, Inc. | https://github.com/ROCm/rocWMMA (ROCm 7.2.4) | Not vendored: `hip/mma_leaf.h` includes its headers from the ROCm installation, and their inline code is compiled into `moss-infer-hip` (the AMD builds' 16x16x16 bf16 leaf). Licence text below |
+| HIP (runtime headers) | MIT | Copyright (C) 2015-2025 Advanced Micro Devices, Inc. | https://github.com/ROCm/hip (ROCm 7.2.4) | Not vendored: `hip/` includes its headers from the ROCm installation, and their inline code is compiled into `moss-infer-hip`. The HIP runtime library (`libamdhip64`) is not shipped: ROCm provides it. Licence text below |
 | xxHash (XXH64 algorithm and constants) | BSD-2-Clause | Copyright (c) Yann Collet | https://github.com/Cyan4973/xxHash (specification) | `src/util/xxh64.h` is an independent implementation from the specification |
 | Unicode tables of HF `tokenizers` 0.23.3 | Apache-2.0 (tokenizers); Unicode License v3 (Unicode Character Database) | Copyright (c) Hugging Face, Inc.; Copyright (c) Unicode, Inc. | https://github.com/huggingface/tokenizers | `src/tokenizer/unicode_tables.inc` (generated data, measured from the library by `tools/gen/gen_unicode_tables.py`) |
 
@@ -61,7 +63,8 @@ FlashAttention-2 and SageAttention are vendored directly; see the table above.
 - Apache-2.0: `third_party/sageattention/LICENSE`. comfy-kitchen, torchsde, LyCORIS and HF `tokenizers` use the same licence text,
   https://www.apache.org/licenses/LICENSE-2.0.
 - MIT / Unlicense: `third_party/stb/LICENSE`.
-- BSD-2-Clause (xxHash), MIT (ggml) and Unicode License v3: reproduced below, since they have no file in this repository.
+- BSD-2-Clause (xxHash), MIT (ggml, rocWMMA, HIP) and Unicode License v3: reproduced below, since they have no file in
+  this repository.
 
 ## xxHash licence (BSD-2-Clause)
 
@@ -124,6 +127,55 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 SOFTWARE.
 ```
 
+## rocWMMA licence (MIT)
+
+From ROCm 7.2.4's `share/doc/rocwmma/LICENSE.md` (https://github.com/ROCm/rocWMMA):
+
+```
+Copyright (C) 2016-2025 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated
+documentation files (the "Software"), to deal in the Software without restriction, including without limitation the
+rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit
+persons to whom the Software is furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all copies or substantial portions of the
+Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE
+WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
+COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
+OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+```
+
+## HIP licence (MIT)
+
+From ROCm 7.2.4's `share/doc/hip/LICENSE.md` (https://github.com/ROCm/hip):
+
+```
+MIT License
+
+Copyright (C) Advanced Micro Devices, Inc.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.
+```
+
 ## Unicode License v3
 
 From https://www.unicode.org/license.txt: