preflight: rocm7-0-34-4-dynres lists 908 — no gfx1151 kernel changes

[glennneuber]

This PR adds 908 to the rocm7-0-34-4-dynres patch set and records the gfx1151 check that #375 asked for. It targets #375's branch, not main. I measured it on gfx1151 only, and it touches no other platform's profile.

What

• expectations.toml. rocm7-0-34-4-dynres now lists 908, with a provenance note in the same form as cuda-dynres-903's. The ROCm build applies 908 in both stages, so the patch set must list it. reference_image stays 0.34.3-dynres-5-g29ae523, which is the image the ladders were measured on.
• The fold doc. It gets a 908 bullet under gate 4 on gfx1151 and the new run under gate 5 on gfx1151. The status rows 4 and 5 and the ROCm clause of open item 2 now say that the check is done.
• The run record runs/preflight-rocm7-0344-fold-g5584539.json. It is force-added past the runs/ ignore, as the other cited records are.

Measured: 908 changes no gfx1151 kernel

I built 558453953 with scripts/build_rocm.sh as maxusai-ollama:0.34.3-dynres-22-g5584539-rocm7-gfx1151. The build took 47 s from ccache. All eight patches apply in the CPU stage and the HIP stage. I compared it with the fold image 0.34.3-dynres-5-g29ae523:

• Payload. I took the sha256 of every entry. There are 1863 on each side, with no file on one side only. One file differs: rocm_v7_2/libggml-hip.so.

• gfx1151 device code. I split the library's .hip_fatbin into its offload bundles. There are 138 in each build, and 22 differ, all in .text only. Section sizes, symbol tables, .rodata and the kernel descriptors are identical.

  • The 78 flash_attn_ext_f16 kernels that compile for gfx1151 are byte-identical. They cover D = 64 to 256, and 18 of them are at D = 256.
  • The other 184 flash_attn_ext_f16 entries are 48-byte NO_DEVICE_CODE stubs. On RDNA, every D = 512 variant is one of these. Each stub differs in exactly one byte: the __LINE__ literal that it passes to no_device_code. The literal changes 1833 → 1807 in 90 stubs (the logit-softcap filter) and 1861 → 1835 in 94 stubs (the RDNA WMMA filter). The cause is that 908 deletes 26 lines of fattn-mma-f16.cuh above both.

• Host code in the same library. .text is 80 bytes shorter, and .rodata is 16 bytes longer. After I normalise addresses, the disassembly differs in 277 hunks:

  • 252 are the __LINE__ literals of the two CUDA_CHECK(cudaFuncSetAttribute(…)) calls in the MMA launcher (2085/2098 → 2059/2072);
  • 18 are the inlined Ampere rows: the D = 512 constants {nbatch_V2, nbatch_combine} = 128 ↔ 256, and the compare chain that selects them;
  • 7 are alignment padding.

  The host picks the config table by device. ampere_mma_available(), turing_mma_available() and volta_mma_available() are false on every AMD device, so gfx1151 takes ggml_cuda_fattn_mma_get_config_rdna, which 908 leaves unchanged. The RDNA fallback is fattn_mma_config(32, 1, 0, …), not the Ampere table.

So #375's reading holds. On gfx1151, 908 amounts to a template-parameter rename and line numbers. No kernel changes, and no launch parameter that gfx1151 reads changes.

Preflight on the 908 image: VERDICT PASS, PASS=20 SKIP=12

I ran preflight.py --platform rocm7 --quality against the 908 image, on a container of its own, with 908 in the patch set. Every check reads the same value as run 1 on 0.34.3-dynres-5-g29ae523, apart from the version and the image tag:

• payload_pin reads 161755f29;
• the three token ladders are 5/5;
• the pinned budgets are 3328 → 3270 and 560 → 529;
• think_format passes in 1401, 127 and 273 tokens.

The 12 skips are the same as run 1's.

Verification

• python3 test_verdicts.py passes: 196 tests OK, 6 skipped. test_summarizers.py and test_rescore.py pass too.
• To repeat the kernel check on another HIP target, build both images and extract libggml-hip.so from each. Split .hip_fatbin on __CLANG_OFFLOAD_BUNDLE__ and take each bundle's amdgcn-amd-amdhsa--<gfx> entry. Then compare each FUNC symbol's bytes in .text.

amd-server/rocm-gfx1151

🤖 Generated with Claude Code