feat(ck-tile): enable multi-D GEMM bridge on gfx1250 (MI400)

[ozturkosu]

Summary

Enables the merged multi-D GEMM bridge (#9308) on gfx1250 (MI400).

Stacked on #10904 (users/muozturk/ck/gemm-universal-gfx1250-enable), which supplies the shared gfx1250 dispatcher enablement. This PR adds only the multi_d-specific piece on top.

Does the merged multi_d bridge support gfx1250? No (as merged).

Two reasons the merged #9308 does not run on gfx1250:

1. Shared arch gate — gemm_utils._SUPPORTED_ARCHES, the wave-combo fallback tables, and arch_specs_generated.py did not include gfx1250. This is fixed by the base PR feat(ck-tile): enable gemm_universal GEMM bridge on gfx1250 (MI400) #10904.
2. MFMA-only CI config — the merged multi_d CI config (gemm_multi_d/configs/default_ci_config.json) pins warp_tile = 32x32x16, an MFMA shape. gfx1250 has no MFMA units; it runs the RDNA4 WMMA path whose fp16/bf16 warp tile is 16x16x32. On gfx1250 the MFMA kernels report as unsupported (status -2/-1) and never run.

Changes

• gemm_multi_d/configs/default_ci_config_gfx1250.json (new): mirrors the multi_d CI config but pins the WMMA warp tile (16x16x32) and keeps the multi_d_config D-fusion block (elementwise_ops x num_d_tensors).
• dispatcher/tests/test_multi_d_bridge.py: new TestMultiDGfx1250Config asserting the gfx1250 config exists, uses WMMA (16x16x32, not MFMA 32x32x16), and retains the multi_d block. The pre-existing test_configs_are_valid_sweeps also now covers the new config.
• gemm_multi_d_full_benchmark.py: add gfx1250 to the --arch help text.

No C/C++ changes, so no clang-format needed. The merged multi_d bridge is fp16-only by design, so this enablement targets fp16.

MI400 validation (gfx1250, ROCm 10.1)

Built the dispatcher static lib for -DGPU_TARGETS=gfx1250 and ran:

python3 gemm_multi_d_full_benchmark.py \
    gemm_multi_d/configs/default_ci_config_gfx1250.json \
    --arch gfx1250 --dtype fp16 --layout rcrr --verify

• 64 kernels compiled for gfx1250 (WMMA 16x16x32).
• 80 measurements VERIFY against the fp32 numpy reference (max_rel ~5e-4), evenly across all 4 default problems (1024^3, 2048^3, 4096^3, 512x512x8192) — 20 verified rows per problem.
• Coverage spans compv3/compv4/mem pipelines, both MultiDAdd/MultiDMultiply, and both num_d_tensors 1/2.
• Throughput 26-69 TFLOPs (median 31).
• Remaining status -1 kernels are unsupported trait combos gracefully surfaced by the bridge (not correctness failures); every problem still has passing kernels.

Notes

• Depends on feat(ck-tile): enable gemm_universal GEMM bridge on gfx1250 (MI400) #10904; merge that first (this PR targets its branch).
• arch_specs_generated.py is auto-generated from arch_specs.json, which has no gfx1250 entry. Per the precedent in feat(ck-tile): enable gemm_universal GEMM bridge on gfx1250 (MI400) #10904, the generated file is hand-edited for gfx1250 until the source spec is updated.

Test plan

• CPU unit tests pass: pytest dispatcher/tests/test_multi_d_bridge.py (13 passed)
• MI400 gfx1250 sweep with --verify: 80/80 measured kernels correct
• CI (stacked on feat(ck-tile): enable gemm_universal GEMM bridge on gfx1250 (MI400) #10904)

[ozturkosu]

gfx1250 (MI400) GEMM bridge enablement — status across the stacked PR series

Variant	Merged PR	Enablement PR	MI400 (gfx1250) result	Status
gemm_universal (foundation)	#8997 / #8998	#10904	640/640 fp16+bf16 VERIFY; fp8 up to ~625 TFLOP/s	✅ full (fp16/bf16/fp8)
grouped	#9000	#10914	64/64 fp16 OK (~178–217 TFLOP/s)	✅ full
multi-D	#9308	#10915	80 measurements VERIFY (max_rel ~5e-4)	✅ full
multi-ABD	#9305	#10917	32/32 fp16 VERIFY	✅ full (fp16-only, upstream design)
microscaling (mx)	#9329	#10916	codegen fp8+fp4 + CPU tests pass; GPU blocked by gfx950-only C++	⚠️ partial (C++ follow-up)
stream-K	#9028	#10919	codegen+compile OK; kernel hangs at runtime on RDNA4	⚠️ partial (C++/kernel follow-up)

All enablement PRs are stacked on #10904 (shared gfx1250 arch enablement), Python/JSON-only, each with a CPU test. Validated on MI400 / gfx1250, ROCm 10.1.

Remaining C++ follow-ups: stream-K cross-workgroup coherency handshake (CDNA atomics → RDNA4/wave32); mx gfx1250 scale-preshuffle (ck_tile already ships preShuffleScaleBuffer_gfx1250).

[ozturkosu]

Superseded by #10921, which combines the shared gfx1250 enablement (gemm_universal) with the grouped, multi-D, and multi-ABD variant configs+tests into a single PR against develop. Branch retained (not deleted) as the base of the remaining partial PRs.