Repository navigation
moss-infer-hip: the native AMD (ROCm/HIP) program - #11
Merged
Merged
Conversation
The bf16 exact path of one block, standalone and unknown to the main build. One kernel source builds for gfx1100 and gfx1201 on rocWMMA, and for NVIDIA on a scalar leaf that verify compares with the engine's kernels on the same inputs, synthetic and real. No AMD device has run this. It shows the kernel logic above the leaf matches the engine, and that both AMD targets compile to WMMA with no scratch or spills. It shows nothing about WMMA's execution or AMD speed. The bench's known-answer check is what a first run on a card reports. README has the gates, how they changed after results, the mutants and what stays unverified.
On an RX 9060 XT the known-answer check passes on the first run, so fragment order, __shared__ memory and barriers hold on gfx1200. It is slow: one block at 8484 tokens takes 415 ms against the engine's 31.3 ms on an RTX 4090 Laptop. WMMA on registers reaches 54 to 96 TFLOPS and the real walk 6.3: reading the operands is the limit, and K panels did not lift it.
… against the engine on the card One block at 8484 tokens goes from 415 ms to 104 ms on gfx1200 with real weights. The previous commit blamed reading the operands; that was wrong. Plain reads reach about 270 GB/s, and 8-bit operands on the strided layout were no faster than bf16. The causes were: - strided fragment loads touching 16 cache lines instead of 8; - loads issued before each group of WMMAs instead of all first; - a raster tile order with no cache reuse; - __syncthreads emitting a wait and a cache invalidate in one-wave kernels. Operands are now packed into 16x16 fragments and walked in 2D blocks of about 4 MB, with all loads first. barrier() is a wavefront fence. The padded row-major tail keeps a DS wait before its FLAT read (pad_barrier). On NVIDIA's scalar leaf, the old and new kernels give the same bits on every op, both real blocks, tail shapes and denormal attention. On the card, hip_block_replay compares each op and the whole block with the engine's outputs from verify --real --dump, at full shape on real activations. All outputs are NaN-poisoned before each run. Replay gates each op on scaled and on its lowest token row, against the scalar leaf's figures in meta.txt. It refuses a dump without a manifest. check_device.sh checks exact per-kernel WMMA counts for each arch on ROCm 7.2.4. verify adds op-3 edge rows through the packed and row-major paths. Nothing after the third review's fixes has run on a card. Old against new on WMMA, the fence against __syncthreads, and gfx11 are unrun.
Replay missed errors confined to one fragment or column: one 16x16 attention fragment x1.03, 16 columns of a qkv tile x1.03, or one down column x1.06 passed every replay gate. On a card, a rocWMMA lane-layout bug would have that shape. agreement() now counts outputs above scaled 4, and those finite on one side and Inf on the other. verify --dump records the scalar leaf's count per op in dump format 3. replay gates each op at max(8, 8x that count), and refuses formats 1 and 2. All three mutants fail on the count, at both real blocks. The unmodified code passes. On the scalar leaf the counts are 0 except block-41 up_gate (1) and down (31). Two other accumulation orders give 1 and 29 to 30. The card's earlier run implies 0 wherever its scaled was 4 or less. Its block-41 SiTU and down counts are still unmeasured.
…counts The exact per-kernel WMMA table is ROCm 7.2.4's, and a card's own ROCm may unroll differently. With HIP_BLOCK_WMMA_COUNTS=report, a kernel whose count differs prints "COUNT MISMATCH: <kernel> n vs expected m", and the mismatch is left out of the exit status. A GEMM, SiTU or attention kernel with no WMMA still fails, and so does every other check: - wave and workgroup size 32; - no scratch or spills; - no dynamic stack; - only known kernels; - no undefined symbols. CMake runs its device check with the variable unset, so a build always enforces the counts.
… XT, ROCm 10.1) Old and new kernels give identical bits on WMMA, packed and row-major; the wavefront fence gives __syncthreads' bits; the one-wave probe is clean; the WMMA counts match ROCm 7.2.4's; replay passes against the engine with the first card figures for its row and count limits.
…ols/hip_block A standalone CMake project for AMD (gfx1100, gfx1200, gfx1201, rocWMMA leaf) and for NVIDIA (HIP's NVIDIA headers, scalar leaf): device runtime, safetensors reader, the block kernels from tools/hip_block (gfx1100/1200/1201 assembly unchanged by the move), which now builds against hip/. The DiT (dit.h, dit_kernels) is in progress.
fp32 NHWC activations, TF32 convolutions and fp32 attention linears as the engine's reference decode, on the bf16 WMMA leaf through exact piece products (2 bf16 pieces per TF32 operand, 3 per fp32), GroupNorm with the engine's exact statistics. hip_vae_check: every kernel against host replays (exact-structure cases bit for bit), a non-square decode against a host replay of the whole decoder, and the golden decodes (g_fridge_1024 69.37 dB, g_fridge_1456 69.92 dB, max |diff| 1) on NVIDIA's scalar leaf.
GatedDeltaNet and full-attention layers in the engine's op order, sequences packed along the rows, layers streamed from pinned staging. On the nvidia build every non-GEMM op gives the engine's bits; hip_te_check gates every golden case against the engine's noise-floor test (2.5x, the off-reference factor).
…y stream DeviceBuffer::upload/download were pageable hipMemcpy calls that could return before the DMA landed and were not ordered with non-blocking streams; they now sync the buffer's device before and after. The RoPE table is a device kernel with the engine's expressions. hip_dit_run reads cfg and the uncond skips per case, gates NaN in every step, step 0 on cosine and rms, the final latent as a gross-error gate, and compares bit for bit with --expect. The rms_norm check verifies torch's reduce order through the per-row sums. Budgets below two block slots and runs past the engine's token limit are refused.
# Conflicts: # hip/CMakeLists.txt # hip/check_device.sh
…, gfx1200 and gfx1201 The text encoder, the DiT with CFG and the official Euler step on the host in the engine's op order, and the VAE decode, with the engine's host code for tokens, noise and schedule. Flags, defaults and messages follow moss-infer generate; --vram-budget has the engine's meaning, and every budget that cannot fit is refused before step 0. One fat binary holds every target's code objects; check_device.sh checks each target's. Installed with RUNPATH /opt/rocm/lib (ROCm 7).
# Conflicts: # hip/CMakeLists.txt # hip/README.md # hip/check_device.sh # hip/runtime.cpp # hip/runtime.h
…rom the DiT's own shapes build-hip.yml builds the gfx1100/gfx1200/gfx1201 package in AMD's ROCm 7.2.4 container (pinned by digest), runs check_device on every source and architecture, and checks the package's RUNPATH, linkage and start-up; ci.yml and release.yml call it, and the release requires the HIP archive. Dit::minimum_weight_bytes/block_bytes/ resident_blocks_for come from the shape lists the constructor loads; generate plans the budget from the token counts before the text encoder and again before the DiT loads, with room for allocation rounding and pinned host memory.
… run's own output On the RX 9060 XT the gated attention check failed at 1.44 ULPs against a 1-ULP limit while the ungated check passed at 1.00: the gate multiplies the bf16 attention output and rounds again, so a 1-ULP difference there (the matrix cores' summation order is not the replay's) can reach 2 ULPs of the gated value. The ungated run bounds the attention against the replay; the gated run now checks the gate on top of the ungated run's output, bit-exact on the card (100.000 % same, 0 ULPs).
… golden comparisons
…nt, the card figures' wording in the docs
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Draft: waiting on the AMD card run before it is ready.
moss-infer-hipis a standalone ROCm 7 program for AMD Radeon GPUs, alongsidemoss-infer-vulkan. As with Vulkan, the frontend chooses which program to run, and the engine'ssrc/and top-level CMake are unchanged.hip/: the DiT, the Qwen3.5 text encoder and the FLUX.2 VAE, with GEMMs and attention on RDNA WMMA through rocWMMA. Host code for tokens, noise, schedule, CFG and Euler is the engine's.mossinfer-<version>-linux-hip-x86_64.tar.gz, RUNPATH/opt/rocm/lib.build-hip.ymlbuilds it in AMD's ROCm 7.2.4 container and statically checks every kernel on every architecture. The release requires the archive.docs/hip.md.Verified so far: