feat(hip): RDNA4 kernels and torch.ops.nf12 parity on ROCm by Qubitium · Pull Request #3 · dphnAI/nf12

Qubitium · 2026-11-08T10:44:27Z

Adds source-built RDNA4 ROCm support behind the same nf12 C, Python, and torch APIs as CUDA. Both backends share the packed format, C ABI frontend, and torch operator layer. The HIP backend covers encode/decode, BF16t packing, dense linear, embedding, and MoE, including the single-token fused route. Inference is graph-capturable and uses caller-owned or torch-managed scratch. The default wheel compiles and installs native libraries at package build time, following the CUDA project. NF12_BUILD_NATIVE=OFF ships sources and JIT-builds the libraries on first use.

Native HIP entry points enforce the shared ABI buffer checks. ROCm PyTorch selects HIP automatically, and incompatible torch-backend builds fail at configuration. MoE handles dropped assignments and BF16t weights without reading absent metadata. Public signatures are unchanged.

Validation

R9700 (gfx1201, ROCm SDK 10.1.0 / HIP 7.16.26385), RTX 4090, and RTX 5090; free-threaded Python with GIL disabled and two visible devices per backend.

Suite	Passed	Skipped	Deselected	Seconds
R9700 HIP + shared API	644	12	1	41.58
RTX 4090 shared API	553	16	1	75.81
RTX 5090 shared API (merged with CUDA master)	558	11	1	113.82

Top-level ROCm libraries build with C++17; CUDA-only configuration succeeds, and a mixed HIP/CUDA torch-extension build is rejected.

Tests cover exact NF12/BF16t equality, codec bit patterns, concurrent streams and devices, graph replay, native invalid-buffer rejection, and dropped MoE assignments. The deselected test is the separate reference-tolerance case tracked in PR #7. A ROCm wheel built without isolation contained both native libraries and 78 source files; the NF12_BUILD_NATIVE=OFF wheel contained 78 sources and no .so. Both installed wheels passed an exact BF16 encode/decode round trip, with first-use JIT verified for the source-only wheel. The RTX 5090 run uses PR head 98c8265 merged with CUDA master baa1c419, including its warp-deadlock fix. Hardware validation does not cover gfx1200.

Performance

On R9700, public encoding improves up to 1.170× on large weights; cold dense decode improves 2.1–6.1%, and MoE decode 4.9–32.6%. Single-token fused MoE is 1.08–3.00× faster than the corrected composed path across the measured shapes. Small encode shapes are essentially unchanged. Dense linear is still below BF16 on some shapes; the linked tables include those regressions and warm/cold measurements.

BF16T split-K regression coverage

Commit 9fc6283 adds test_bf16t_twin_matches_nf12_split_k_plans, covering the reachable split-2, split-4, and split-8 dense plans. GPT-QModel PR #3716 carries the same coverage for the vendored kernel and links this PR as the upstream implementation proof.

Adds csrc/hip/, a HIP port of the NF12 codec and fused linear kernels for AMD RDNA4 (gfx1200/gfx1201), consuming and producing the identical tiled byte stream as the CUDA kernels so checkpoints stay portable. * format_hip.cuh -- SPEC section 4 geometry, the AMD v_perm_b32 selector conversion, decode_bulk/decode_group, tiling and wmma fragment helpers. Restates the geometry from include/nf12/format.cuh and static_asserts that they match, so a format change breaks the build instead of diverging. * decode.hip -- nf12_hip_decode (bit-exact NF12 -> BF16), nf12_hip_linear_fused (small-M fused decode+GEMM for m <= 16), nf12_hip_selftest. * encode.hip -- native analyze/count/scan/write encoder with a hand-written block-wide exclusive scan, since ROCm ships no CUB. Byte-identical to the reference codec. * wmma.hip -- nf12_hip_linear_wmma, the matrix-core fused path for 17 <= m <= 128 using v_wmma_f32_16x16x16_bf16_w32_gfx12. Decodes each 16x64 super-atom once into shared memory and reuses it for every token tile, so the packed weight is read exactly once and the dense weight never exists. Build is additive and off by default: `-DNF12_ENABLE_HIP=ON` adds the target without touching any CUDA target, and `cmake -S csrc/hip -B build-hip` builds libnf12_hip.so standalone on a host with no CUDA toolkit. Verified on a Radeon AI PRO R9700 (gfx1201): * decode reproduces every input BF16 bit, including +/-0, subnormals, Inf and NaN payloads; * encode matches the reference packed bytes and escape table on 30 cases spanning 16x64 .. 11008x4096; * fused linear: 0 violations of |err| <= 2e-3 + 2e-3*|ref| over 21 shapes, rms_rel ~1.66e-3 (the FP64 oracle's own BF16 rounding floor); * linear wins 1.30-1.44x over a dense BF16 GEMM at 11008x4096 for m=17..96, running at ~640 GB/s against a measured 633 GB/s HBM read ceiling. Known gaps, documented in csrc/hip/README.md: the library is not yet wired into python/nf12 or the torch extension, m > 128 still needs a larger-tile kernel, and only gfx1200/gfx1201 are validated (gfx1250 has no wmma-128b-insts and falls back). Porting notes recorded there: v_perm_b32 selector semantics are reversed versus __byte_perm, hip_bfloat16 degrades to int in a hipified TU, __gfx1201__ is device-pass only so the arch test must be a runtime gcnArchName check, and __builtin_amdgcn_fdot2 is FP16 rather than BF16.

Makes the HIP path reachable rather than a subtree you have to build by hand. CMake: * CUDA is no longer a required project language. `check_language(CUDA)` enables it only when a compiler exists, so an ROCm-only host configures and builds instead of failing at project(). The CUDA targets stay behind NF12_HAVE_CUDA; NF12_ENABLE_HIP stays additive and off by default, and configuring with neither is a clear fatal error. * nf12_hip lands its library in the build root, next to libnf12.so. * csrc/hip/CMakeLists.txt only declares its own project() when built standalone, so add_subdirectory() does not reset the parent scope. JIT (`nf12._jit`): * Detects a ROCm build and builds libnf12_hip.so instead of the CUDA pair, using the visible gcnArchName targets and amdclang++ (CMake >= 4.2 rejects the hipcc wrapper). Compiler discovery honours CMAKE_HIP_COMPILER, amdclang++ on PATH, then ROCM_PATH / HIP_PATH / /opt/rocm. * The cache key includes the backend, so CUDA and HIP builds of the same source cannot collide. Loader (`nf12._library`): * `library_path()` resolves libnf12_hip.so on ROCm and still defaults to libnf12.so on CUDA. * `load()` raises a clear error on ROCm instead of trying to dlopen a torch extension that does not exist there yet. Adds include/nf12/nf12_hip.h: the C ABI actually implemented, with hipStream_t and hipError_t. Every declared symbol was link-checked against the built libnf12_hip.so from a plain C compiler. Verified on this host: * CUDA-only top-level build still produces libnf12.so unchanged; * ROCm-only top-level build (no CUDA compiler) produces libnf12_hip.so; * standalone csrc/hip build produces libnf12_hip.so; * on ROCm, nf12._jit.library_directory() builds and caches the library, nf12.library_path() resolves to it, and the built library launches the selftest kernel (sentinel 0x4E463132) with wmma_supported(64) == 1. Docs: top-level README gains an RDNA4 (ROCm) section; csrc/hip/README.md now documents the real C signatures (the previous block had the wrong argument order) and the JIT integration.

Name both upstream sources in the RDNA4 README and the top-level README, so the attribution is visible without opening the SPDX headers. No code change.

* bench/hip_linear.py reuses bench/common.py's k x n grid so the HIP numbers are directly comparable with the CUDA results, and reports warm/cold ratios plus the FP64 accuracy gate. * tests/hip/test_hip.py drives the C ABI: selftest, byte-identical encode, bit-exact decode (specials included), and fused/wmma linear against an FP64 oracle. * Replace the ad-hoc 11008x4096 table in csrc/hip/README.md with the canonical grid.

* m=16 was the worst point (0.42-1.18x of BF16) because it ran the v_dot2 small-M kernel. A 16-token call is exactly one WMMA tile, so route m >= 16 to the matrix-core path and keep the fused kernel for m <= 15. m=16 improves to 1.20-1.28x on the large canonical shapes. * Fix a latent multi-block bug: origin = m0 + tile_row[t] double-counted m0, so the path only worked while m_blocks == 1. Compute the absolute tile origin inline instead. * Extend the matrix-core route to m <= 256 (two M blocks); above that the decode-then-rocBLAS fallback wins on traffic. * Refresh the README tables and the HIP tests (m=16 now covers WMMA).

At kMt=8 the 8x8 FP32 accumulator costs 64 VGPRs and the kernel lands at 167 VGPRs / 9 waves per SIMD, which under-fills the machine at m=128. When K is long (k_blocks >= 128) and the launch is short of CTAs (n_blocks <= 64), run the 4-tile kernel over two M blocks instead. Measured on gfx1201 (warm / cold, vs dense BF16 F.linear): 8192x8192 m=128 0.68 -> 0.89x 14336x4096 m=128 0.79 -> 0.83x Short-K and wide-N shapes keep the 8-tile kernel and do not regress. m_blocks is now derived from the tile count the dispatcher actually launches (min(tiles, kMaxTiles)), not the unclamped one. At m=256 full_tiles=16 but only the 8-tile kernel is instantiated, so the old expression produced m_blocks=1 and silently dropped tokens 128..255 (caught by the FP64 oracle). Also drop the ModelCloud SPDX line and note the Qubitium adaptation.

Regenerate the warm/cold table from a fresh sequential run on one R9700 and document the shape-aware M split. Also drop the ModelCloud SPDX line from the HIP sources (keeping the AlpinDale/contributors line) and add an adaptation note for Qubitium.

Every split wrote its FP32 partial plane with bias_v already folded in, and wmma_reduce_kernel sums the planes, so the bias was multiplied by kSplits: 4x at m<=64 and 2x at m=128. Add it to the blockIdx.z == 0 plane only, which keeps the FP32-accumulate-then-round-once semantics. Caught by the validator's bias section once m=16 started using this kernel (N=1024, K=1024, m=16: max_abs 0.89 -> 6.1e-05). Regression test covers m = 16, 32, 64, 128 and asserts the delta equals the bias exactly once.

The decoded super-atom is 16x64 BF16, a 128-byte shared row, so every lane in a row group lands on the same LDS banks: the store was 8-way conflicted and the B-fragment load 16-way. Eight BF16 of padding (144-byte rows, 36 banks) turns that into 4-way / 2-way. It costs 4 KiB of LDS per CTA, which drops the resident CTAs per CU from 4 to 3, so the padded stride is a template variant selected only when the launch fits in three CTAs per CU (n_blocks * m_blocks <= 96). Split-K is now chosen from CTA count, K length and M instead of CTA count alone. A split reads a DISJOINT k slice, so the packed weight is never read twice; the cost is the per-split FP32 plane and the reduce pass. Very narrow N (896 rows -> 7 CTAs) splits 8 ways; long-K shapes split twice at small M and once at m=128; 4096x4096 does not split below m=128. R9700, warm, versus dense BF16 (m=16 / m=64): 896x4864 0.62 -> 0.83 0.50 -> 0.69 2048x8192 0.66 -> 0.83 0.59 -> 0.65 4096x4096 0.48 -> 0.54 0.45 -> 0.48 25600x5120 1.15 -> 1.33 0.98 -> 1.27 8192x8192 m=128 0.68 -> 0.90 No shape regresses; 0 accuracy violations, 22/22 hip tests pass.

Three fixes to the fused matrix-core path, aimed at the weights whose dense BF16 form fits in the 64 MiB Infinity Cache (4096x4096, 2048x8192, 896x4864). * The store no longer runs quad_transpose. It is a 4x4 dword transpose inside each 16-byte chunk (16 ds_bpermute_b32 + 32 selects per super-atom), and every ds_bpermute occupies the same LDS crossbar as the store and the B-fragment load. It is self-inverse, so undoing it on the address instead of the data is equivalent: dword d of lane l's chunk (h,q) belongs to dword (lane&3) of chunk (d+4q) in row (lane>>2)+8h. That is 16 STS.32 whose bank is 4*(lane>>2)+(lane&3)+const, i.e. all 32 banks exactly once. * The next super-atom's loads are issued before the mma and decoded after it, so the L3 latency is covered by the matrix-core work instead of stalling the decode. A 4096x4096 launch at m=16 is only two warps per SIMD. * Split-K is picked per shape from measured optima rather than a CTA-count rule: narrow N splits 8 ways, very wide N splits 2 ways at eight token tiles, everything else stays unsplit. R9700, warm, versus dense BF16 (m=16 / m=64), before -> after: 4096x4096 0.48 -> 0.81 0.45 -> 0.65 2048x8192 0.66 -> 1.03 0.59 -> 0.78 896x4864 0.62 -> 1.14 0.50 -> 0.84 25600x5120 1.15 -> 1.35 0.98 -> 1.21 8192x8192 m=16 1.21 -> 1.29 0 accuracy violations, 22/22 hip tests pass.

The split-K reducer was launched with (total+255)/256 blocks, and on gfx1201 exactly 256 blocks is a pathological 14.2 us case against 3.3 us at 64, 128 or 512. Cap the grid at 128. pick_splits also split shapes whose K is too short to give each split a real slice. 896x4864 (14 k-blocks) was taking 4 ways and paying four prologues for three super-atoms each, costing 9.7 -> 13.5 us; it now stays unsplit, while 4096x4096 (64 k-blocks) keeps 4 ways at m=16 (27.3 -> 25.7 us) and 2 above it. m=16 warm ratios against dense BF16: 4096x4096 0.88, 2048x8192 1.02, 896x4864 1.14 (cold 1.09 / 1.23 / 1.09). Refresh both README performance tables and correct the stale k-block figures in the pick_splits comment block.

One full bench run at bf0d676 across the canonical k x n grid and all six token counts, replacing numbers gathered across several kernel revisions. m=16 warm/cold: 4096x4096 0.89/1.10, 2048x8192 1.03/1.23, 896x4864 1.25/1.25. Wide shapes are 1.07-1.35x at m=16..32 and 0.96-1.24x at m=64; m=128 is 0.75-1.20x.

The m <= 15 fused kernel took m as a runtime argument, so acc[16][2] was indexed dynamically -- which keeps it out of VGPRs -- and the warp-partial plane was always sized for 16 tokens. Making the token count a template parameter fixes both and unrolls the token loop. At m=1 that raises occupancy enough to hurt: the kernel then runs eight CTAs per CU and 256 concurrent packed streams, which thrashes the 64 MiB LLC. Over-allocating the partial plane to 32 KiB throttles it to two CTAs per CU. Same-process A/B, old -> new: 4096x4096 m=1 28.6 -> 20.5 us (1.40x), m=8 32.6 -> 24.2 (1.35x); 2048x8192 m=1 19.4 -> 18.5 (1.05x), m=8 32.0 -> 23.9 (1.34x); 896x4864 m=8 15.4 -> 11.6 (1.33x). Against dense BF16 that is m=1 1.20x / 1.38x / 1.69x and m=8 0.94x / 0.97x / 0.97x, up from 0.84 / 1.33 / 1.63 and 0.67 / 0.80 / 0.83. m=16 still routes to the matrix-core kernel, which is 1.2-2.3x faster there on every measured shape.

Mirrors dphnAI/nf12#3. The m <= 15 fused kernel took m as a runtime argument, so acc[16][2] was indexed dynamically -- which keeps it out of VGPRs -- and the warp-partial plane was always sized for 16 tokens. Making the token count a template parameter fixes both and unrolls the token loop. At m=1 that raises occupancy enough to hurt: eight CTAs per CU and 256 concurrent packed streams thrash the 64 MiB LLC. Over-allocating the partial plane to 32 KiB throttles it to two CTAs per CU. Same-process A/B, old -> new: 4096x4096 m=1 28.6 -> 20.5 us (1.40x), m=8 32.6 -> 24.2 (1.35x); 2048x8192 m=1 19.4 -> 18.5 (1.05x), m=8 32.0 -> 23.9 (1.34x); 896x4864 m=8 15.4 -> 11.6 (1.33x).

…3645) Mirrors dphnAI/nf12#3. The m <= 15 fused kernel took m as a runtime argument, so acc[16][2] was indexed dynamically -- which keeps it out of VGPRs -- and the warp-partial plane was always sized for 16 tokens. Making the token count a template parameter fixes both and unrolls the token loop. At m=1 that raises occupancy enough to hurt: eight CTAs per CU and 256 concurrent packed streams thrash the 64 MiB LLC. Over-allocating the partial plane to 32 KiB throttles it to two CTAs per CU. Same-process A/B, old -> new: 4096x4096 m=1 28.6 -> 20.5 us (1.40x), m=8 32.6 -> 24.2 (1.35x); 2048x8192 m=1 19.4 -> 18.5 (1.05x), m=8 32.0 -> 23.9 (1.34x); 896x4864 m=8 15.4 -> 11.6 (1.33x).

Qubitium · 2026-11-08T16:17:46Z

@AlpinDale PR ready to review. RDNA4 is fast enough (still has headroom). I will get MI300X up as well.

Two knobs were hand-tuned and are now computed from the grid. Double buffering costs 18 KiB of shared memory per CTA (36 against 18 single-buffered) out of 128 KiB per CU. At kMt <= 2 the VGPR budget allows well over three CTAs, so that footprint is what caps occupancy; from kMt = 3 the ~135 VGPRs already cap a CU at three and dropping the buffer buys nothing (25600x5120 m=128: 447 us buffered against 479 us not). Where LDS does bind, dropping it only helps in a narrow band -- the launch must be starved enough that one more resident CTA matters, but not so starved that the buffer's own latency hiding wins. 4096x4096 m=16 is four waves against a cap of three: 28.4 -> 23.5 us. The wide shapes are 14-28 waves deep and only lose (4096x28672 m=16: 1.14 -> 1.08x). Split-K multiplies CTAs without duplicating a read, but only pays while the extra CTAs are resident. Four ways is now used up to m=32 (4096x4096: 30.1 -> 27.1 us) and not beyond (m=64: 38.2 -> 43.3, m=128: 62.9 -> 79.8). 4096x4096 m=16 goes 0.88 -> 1.02x warm (1.12x cold) and m=32 0.76 -> 0.81x. Everything else is unchanged within noise.

Mirrors dphnAI/nf12#3. Double buffering costs 18 KiB of shared memory per CTA out of 128 KiB per CU. At kMt <= 2 the VGPR budget allows well over three CTAs, so that footprint is what caps occupancy; from kMt = 3 the ~135 VGPRs already cap a CU at three and dropping the buffer buys nothing. Where LDS does bind, dropping it only helps in a narrow band: 4096x4096 m=16 is four waves against a cap of three and goes 28.4 -> 23.5 us, while the 14-28-wave wide shapes only lose. Split-K now goes four ways up to m=32 (4096x4096: 30.1 -> 27.1 us) and not beyond (m=64: 38.2 -> 43.3, m=128: 62.9 -> 79.8), because the accumulator already caps a CU at three CTAs. 4096x4096 m=16 goes 0.887 -> 1.004x against dense BF16 in the validator.

…pe (#3648) Mirrors dphnAI/nf12#3. Double buffering costs 18 KiB of shared memory per CTA out of 128 KiB per CU. At kMt <= 2 the VGPR budget allows well over three CTAs, so that footprint is what caps occupancy; from kMt = 3 the ~135 VGPRs already cap a CU at three and dropping the buffer buys nothing. Where LDS does bind, dropping it only helps in a narrow band: 4096x4096 m=16 is four waves against a cap of three and goes 28.4 -> 23.5 us, while the 14-28-wave wide shapes only lose. Split-K now goes four ways up to m=32 (4096x4096: 30.1 -> 27.1 us) and not beyond (m=64: 38.2 -> 43.3, m=128: 62.9 -> 79.8), because the accumulator already caps a CU at three CTAs. 4096x4096 m=16 goes 0.887 -> 1.004x against dense BF16 in the validator.

Each header stated the same three facts twice: an SPDX copyright line and then 'Copyright (c) 2026 AlpinDale and the nf12 contributors', an SPDX licence identifier and then 'Licensed under the Apache License, Version 2.0', and the dphnAI/nf12 URL both inside the reference-kernel block and one line above it. The decode, encode and format headers also repeated the 'adapted for GPT-QModel by Qubitium' sentence that the line above already makes. Collapse all of that into one three-line credits block that names each party once. Net -23 lines, no content lost. Also fix the gfx1200/gfx1200/gfx1250 typo in format_hip.cuh and retitle the encode section from 'Why the port is not a copy' to 'Port notes', since the paragraph above it already says the port is a transliteration.

…3650) Mirrors dphnAI/nf12#3. Each header stated the same three facts twice: an SPDX copyright line and then 'Copyright (c) 2026 AlpinDale and the nf12 contributors', an SPDX licence identifier and then 'Licensed under the Apache License, Version 2.0', and the dphnAI/nf12 URL both inside the reference-kernel block and one line above it. Three files also repeated the 'written for GPT-QModel (ModelCloud.ai)' sentence that the SPDX line above already makes. Collapse all of that into one three-line credits block that names each party once. Net -28 lines, no content lost. Also fix the gfx1200/gfx1200/gfx1250 typo and retitle the encode section from 'Why the port is not a copy' to 'Port notes'.

AlpinDale

Thanks for the contribution! I left a few comments.

Also please compare against BF16 instead of FP64, as our spec only guarantees losslessness against a BF16 baseline (see: §6.2 of the spec).

The CTA-wide activation staging is re-read once per column stripe, and at m=64 that traffic (33.5 MB) is larger than the packed weight (25.2 MB). It is the main thing diluting the format's 1.33x byte advantage, so widening the stripe group is worth more than it looks. Eight stripes per CTA was gated on n_tiles >= 512. Measured, it also wins at n=4096 (256 tiles) once the token tile is at least four tiles tall: m=64 26.1 -> 25.8 us, m=128 36.5 -> 35.5 us, and nf12/bf16t 1.07 -> 1.09 and 1.06 -> 1.08. A one- or two-tile CTA has too little MFMA per barrier to amortise the wider staging block, so it keeps four. Grid, cold L2, nf12 vs the bf16t twin: shape m=1 m=16 m=32 m=64 m=96 m=128 4096x4096 1.067 1.017 1.052 1.098 1.074 1.076 4096x14336 1.109 1.064 1.113 1.100 1.077 1.079 14336x4096 1.115 1.058 1.079 1.148 1.112 1.079 tests/api + tests/hip: 591 passed, 27 skipped. gfx1200/gfx1201 device code unchanged (218 dphnAI#3 functions identical, 52 additive).

Analyze, count and write each stream the weight once. The 64 MiB Infinity Cache only retains the tail of the previous pass, so reading every pass front-to-back made the next pass start on the evicted end. Counting back-to-front starts where analyze ended and finishes on the head, which is where the write pass starts. Counts are still stored by flat unit id, so scan order, escape order and packed bytes are unchanged. Public nf12.encode wall clock, 8 warmup / 31 alternating samples, median: 1.165x on 4096x14336, 1.170x on 14336x4096, 1.130x on 8192x8192, 1.074-1.078x on 4096x28672 and 5120x25600, 1.017x on 8192x57344, neutral at 4096x4096 and below. Kernel-only encode_gpu 1.212x/1.224x. A 512 MiB L3 flush before each call still shows 1.152x, so the gain is not a warm-benchmark artifact.

Dense decode and bf16t unpack already ran at the streaming ceiling, so the remaining cost was occupancy: two super-atoms per warp took 48 VGPRs and halved the CTA count. One super-atom per warp drops that to 32 VGPRs and doubles CTAs; the arithmetic is untouched. moe_decode was genuinely under-parallelised. Its grid was (row blocks, experts) and each warp walked k_blocks/kWarpsPerBlock super-atoms serially, reaching 435-496 GB/s. blockIdx.x now packs (row block, K slice) so each CTA owns a disjoint contiguous run of whole super-atoms and each warp owns one. Every output element is still written exactly once; there is no reduction. Cold, interleaved A/B on gfx1201: - decode 4096x4096 105.16 -> 102.03 us, 14336x4096 360.28 -> 350.63, 8192x8192 411.66 -> 403.27, 896x4864 30.14 -> 28.42 (+2.1-6.1%); unpack_bf16t improves similarly. - moe_decode E=4 14336x4096 1891.6 -> 1426.6 us (+32.6%), E=4 4096x4096 485.4 -> 405.6 (+19.7%), E=8 4096x4096 947.6 -> 818.5 (+15.8%), E=4 2048x8192 424.2 -> 404.3 (+4.9%); now 574-581 GB/s. WMMA m=16..256 is unchanged (median ratio 1.000, min 0.965, max 1.040). Verified bit-exact for all 65536 BF16 patterns through encode/decode, both forced exp_bases, pack_bf16t/unpack, and MoE decode at K=4096 with slices; full suite 600 passed, 12 skipped.

Refresh the MI300X dense table for the CTA-count split rule (m=16 21.2 -> 19.0 us, m=32 24.0 -> 21.6 us) and state the two baselines separately: nf12 beats the bf16t twin by 1.04-1.15x on every PR dphnAI#3 grid shape, but loses to torch.matmul above m=1 (0.56-0.86x). Explain the loss in bytes rather than hand-waving: split-K is what fills 304 CUs from 256 column tiles, and at m=128/n=4096 the eight-split FP32 partial plane is 16.8 MB written plus 16.8 MB read, so nf12 moves 59.8 MB where rocBLAS moves 34.5 MB. Dropping to kSplits=1 restores the format's 1.32 byte ceiling but costs 4x absolute time (108.5 us vs 25.9 us at m=64).

A split whose range is shorter than two 64-column K blocks spends more on its FP32 partial plane (m*n*4 bytes written and read) than the extra CTAs buy back. This is the floor the OPUS gfx942 launch plan applies as min_iters; it binds on short-K shapes, where 896x4864 gives 14 K blocks so eight splits is 1.75 iterations each. Measured on MI300X, 896x4864, cold L2, graph timing, median of five runs of 100 samples: m nf12 us bf16t us ratio before 32 10.88 11.91 1.094 1.064 64 16.21 17.75 1.095 1.020 128 23.63 25.17 1.065 1.057 Restricted to a CTA that batches token tiles: a one-tile CTA has so little work per split that the extra CTAs are what fills the machine, and halving them measured 12% slower at m = 8 on the same shape. Every other cell of the PR dphnAI#3 grid is unchanged, because its k_blocks/kSplits is already >= 2.

k = 896 is 14 super-atoms, so the sixteen-warp launch gives every warp a single K block: the plane prefetch has nothing to overlap and the load latency is paid in full on a kernel whose decode is already the dominant work. Eight warps give two blocks each and the prefetch overlaps again. Only a one- or two-tile launch takes it. From three tiles up the decode is amortised over more MFMAs and the wider CTA measured faster, so the choice is gated on both the K depth and the token-tile count. 896x4864, cold L2, graph timing, median of three runs of 100 samples, ratio against the bf16t twin: m 8 warps 16 warps 1 1.017 0.990 2 1.029 1.023 4 1.012 1.011 5 1.006 1.007 8 1.002 1.002 Every other shape in the PR dphnAI#3 grid has 32 or more K blocks, so its launch is unchanged.

Fifty-two of fifty-four cells are at or above parity, peak 1.194. The two on the line are m=16 on the two smallest weights, where a CTA holds one token tile and the decode latency is the same order as the bytes it saves. Also corrects the aiter item 6 verdict: the OPUS min_iters guard is not a no-op, it binds at k=896 and is now adopted.

Per-symbol comparison of the gfx1200/gfx1201 offload bundles against PR dphnAI#3 (24ee25d) by encoded instruction words: 218 PR dphnAI#3 functions, 0 changed, 0 removed, 88 added, all of them CDNA3-only kernels. Also widens the quoted nf12/bf16t range to 1.05-1.17x, which is what the grid now measures.

Measured at m = 2, 4, 96, 256 on three shapes: nothing below parity (min 1.013), and the peak is 1.511x at 5120x25600 m=256 because m >= 129 leaves the fused route for decode + rocBLAS. That is where the 1.3 figure lives; the fused band caps at 1.19-1.24 because it carries the split-K partial plane and the decode in the same launch.

Qubitium · 2026-11-09T11:30:05Z

@AlpinDale branch passing unit tests on r9700, 4090 and 5090. Will pause now for review.

Qubitium · 2026-11-09T12:37:04Z

@AlpinDale I rechecked the current PR head against the requested changes:

wmma_arch_cached(device) now caches the capability per device ordinal (cefa071), including the device-specific query path.
csrc/hip/CMakeLists.txt is included in source-only wheel packaging.
nf12_hip now has an install target.
_violations() rejects non-finite outputs, so NaNs cannot pass the tolerance check.
HIP JIT cache identity includes the ROCm compiler, toolkit, and HIPFLAGS.
The shared torch op layer now builds for HIP, and the README reflects NF12_BUILD_TORCH=ON plus the source-only JIT fallback.
Commit 9fc6283 adds BF16T split-2/4/8 coverage; GPT-QModel PR #3716 carries the corresponding vendored regression test.

Qubitium added 3 commits November 8, 2026 10:44

docs(hip): add explicit AlpinDale and Together AI credits

c0d4bdb

Name both upstream sources in the RDNA4 README and the top-level README, so the attribution is visible without opening the SPDX headers. No code change.

Qubitium force-pushed the feat/rdna4-hip branch from cfd1958 to c0d4bdb Compare November 8, 2026 11:27

Qubitium added 4 commits November 8, 2026 12:34

Qubitium mentioned this pull request Nov 8, 2026

perf(nf12): route m=16 to RDNA4 matrix cores, fix m=128 and split-K bias ModelCloud/GPTQModel#3633

Merged

Qubitium added 3 commits November 8, 2026 13:13

Qubitium marked this pull request as ready for review November 8, 2026 14:12

Qubitium mentioned this pull request Nov 8, 2026

perf(nf12): cap the RDNA4 split-K reduce grid and retune the tail ModelCloud/GPTQModel#3644

Merged

Qubitium added 2 commits November 8, 2026 14:59

Qubitium mentioned this pull request Nov 8, 2026

perf(nf12): specialize the RDNA4 fused small-M kernel on the token count ModelCloud/GPTQModel#3645

Merged

AlpinDale self-requested a review November 8, 2026 16:30

Qubitium mentioned this pull request Nov 8, 2026

perf(nf12): derive the RDNA4 buffer and split-K choices from the launch shape ModelCloud/GPTQModel#3648

Merged

Qubitium mentioned this pull request Nov 8, 2026

docs(nf12): de-duplicate the credits block in the HIP file headers ModelCloud/GPTQModel#3650

Merged

AlpinDale requested changes Nov 8, 2026

View reviewed changes

Comment thread csrc/hip/wmma.hip Resolved

Comment thread CMakeLists.txt Resolved

Comment thread csrc/hip/CMakeLists.txt Resolved

Comment thread tests/hip/test_hip.py Resolved

Comment thread python/nf12/_jit.py Resolved

Qubitium added 3 commits November 9, 2026 08:04

docs: record the encode pass-order and decode parallelism measurements

96dd23e

Qubitium added 7 commits November 9, 2026 09:12

fix(hip): validate native codec and linear buffers before launch

338c345

fix(build): select HIP by default for ROCm torch

b0d6e9b

fix(moe): zero dropped slots before composed combine

15a41a8

docs: refresh ROCm validation and fused MoE timings

38a0bd5

fix(build): discover HIP compiler for native ROCm wheels

8026762

docs(hip): label dense BF16 accuracy baseline correctly

bee58d8

docs(api): clarify encode stream completion on CUDA and HIP

412dd22

Qubitium added 4 commits November 9, 2026 09:54

fix(hip): validate native MoE align buffers before launch

6e6f185

fix(build): require C++17 for top-level HIP targets

e94280b

fix(build): reject HIP torch extension with non-ROCm torch

561b0fd

docs(build): show ROCm wheel build without isolation

8e1835a

Qubitium added 2 commits November 9, 2026 10:02

fix(hip): bound native MoE align workspace sizing

40761fe

docs(validation): record scoped clean-review test results

7cf9c9b

Qubitium added 3 commits November 9, 2026 11:13

fix(hip): skip NF12 metadata reads for BF16t MoE

5250c61

docs(validation): refresh full CUDA and ROCm review results

98c8265

docs(validation): record RTX 5090 merge integration checks

d34407a

test(hip): cover BF16t split-K twin plans

9fc6283

Qubitium mentioned this pull request Nov 9, 2026

fix(nf12): run BF16T through split-K dispatcher ModelCloud/GPTQModel#3716

Merged

Navigation Menu

Repository navigation

feat(hip): RDNA4 kernels and torch.ops.nf12 parity on ROCm - #3

feat(hip): RDNA4 kernels and torch.ops.nf12 parity on ROCm#3
Qubitium wants to merge 119 commits into
dphnAI:masterfrom
Qubitium:feat/rdna4-hip

Qubitium commented Nov 8, 2026 •

edited

Loading

Uh oh!

Qubitium commented Nov 8, 2026

Uh oh!

AlpinDale left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Qubitium commented Nov 9, 2026 •

edited

Loading

Uh oh!

Qubitium commented Nov 9, 2026 •

edited

Loading

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

Qubitium commented Nov 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Validation

Performance

BF16T split-K regression coverage

Uh oh!

Qubitium commented Nov 8, 2026

Uh oh!

AlpinDale left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Qubitium commented Nov 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Qubitium commented Nov 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Qubitium commented Nov 8, 2026 •

edited

Loading

Qubitium commented Nov 9, 2026 •

edited

Loading

Qubitium commented Nov 9, 2026 •

edited

Loading