feat(rocke): enable mixed scaled GEMM on shared operand storage

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/arch_specs.json

@@ -144,7 +144,7 @@
       ]
     },
     "gfx1250": {
-      "_comment": "gfx1250-class CDNA multi-chip device on the GFX12 programming model (wave32, WMMA, no MFMA). Primary fp16/bf16 WMMA atom is 16x16x32 (NOT gfx1201's 16x16x16). ABI empirically confirmed on ROCm 7.13 clang 23 (gfx1250 device target): __builtin_amdgcn_wmma_f32_16x16x32_f16 -> llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1,<16 x half>,i1,<16 x half>,i16,<8 x float>,i1,i1); bf16 uses v16bf16 (<16 x bfloat>) directly, not the gfx11 <16 x i16> bitcast. Datalayout/triple byte-identical to gfx950/gfx1201. This catalog ships validated fp16/bf16 and fp8/bf8 WMMA atoms plus homogeneous FP8xFP8 and BF8xBF8 SCALE/SCALE16 with E8M0 scales; FP4/FP6/BF6 block-scaled forms remain intentionally omitted. Lane maps are a HYPOTHESIS verified by examples/gfx1250/wmma_probe.py. MEMORY: gfx1250 HAS async global->LDS DMA (its own GFX12 family: global_load_async_to_lds_b128 + s_wait_asynccnt -- NOT the gfx9 buffer_load_lds/global_load_lds, which do not select), so has_async_global_lds=true (GPU-validated via the DTLA path). It also HAS a hardware transpose-LDS read, ds_load_tr16_b128 (GPU-validated via compute_pv_dstr). The has_async_lds and has_ds_read_tr flags stay false on purpose: in this catalog they denote the gfx9/gfx950 ds_read_*_tr_* / async-LDS *ABI* used to admit the gfx950 MFMA tiled builder (attention_arch.py), which gfx1250 does NOT have -- its async/transpose use distinct GFX12 opcodes gated directly in the ISA backend, not via these MFMA-pipeline bits. has_tdm=false: TDM not verified on this part.",
+      "_comment": "gfx1250-class CDNA multi-chip device on the GFX12 programming model (wave32, WMMA, no MFMA). Primary fp16/bf16 WMMA atom is 16x16x32 (NOT gfx1201's 16x16x16). ABI empirically confirmed on ROCm 7.13 clang 23 (gfx1250 device target): __builtin_amdgcn_wmma_f32_16x16x32_f16 -> llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1,<16 x half>,i1,<16 x half>,i16,<8 x float>,i1,i1); bf16 uses v16bf16 (<16 x bfloat>) directly, not the gfx11 <16 x i16> bitcast. Datalayout/triple byte-identical to gfx950/gfx1201. This catalog ships validated fp16/bf16 and fp8/bf8 WMMA atoms plus homogeneous and mixed FP8/BF8/FP6/BF6/FP4 SCALE/SCALE16 with E8M0 scales. Lane maps are a HYPOTHESIS verified by examples/gfx1250/wmma_probe.py. MEMORY: gfx1250 HAS async global->LDS DMA (its own GFX12 family: global_load_async_to_lds_b128 + s_wait_asynccnt -- NOT the gfx9 buffer_load_lds/global_load_lds, which do not select), so has_async_global_lds=true (GPU-validated via the DTLA path). It also HAS a hardware transpose-LDS read, ds_load_tr16_b128 (GPU-validated via compute_pv_dstr). The has_async_lds and has_ds_read_tr flags stay false on purpose: in this catalog they denote the gfx9/gfx950 ds_read_*_tr_* / async-LDS *ABI* used to admit the gfx950 MFMA tiled builder (attention_arch.py), which gfx1250 does NOT have -- its async/transpose use distinct GFX12 opcodes gated directly in the ISA backend, not via these MFMA-pipeline bits. has_tdm=false: TDM not verified on this part.",
       "family": "cdna",
       "target_family": "gfx12_cdna",
       "matrix_path": "wmma",
@@ -188,7 +188,47 @@
         {"family": "wmma_scaled", "a": "bf6", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_bf6_scale_e8m0_e8m0_k16"},
         {"family": "wmma_scaled", "a": "fp4", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_fp4_scale_e8m0_e8m0_k32"},
         {"family": "wmma_scaled", "a": "fp4", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_fp4_scale_e8m0_e8m0_k16"},
-        {"family": "wmma_scaled", "a": "bf8", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_bf8_scale_e8m0_e8m0_k16"}
+        {"family": "wmma_scaled", "a": "bf8", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_bf8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp8", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp8_bf8_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp8", "b": "fp6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp8_fp6_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp8", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp8_bf6_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp8", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp8_fp4_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "bf8", "b": "fp8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_fp8_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "bf8", "b": "fp6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_fp6_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "bf8", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_bf6_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "bf8", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_fp4_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp6", "b": "fp8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp6_fp8_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp6", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp6_bf8_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp6", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp6_bf6_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp6", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp6_fp4_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "bf6", "b": "fp8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_fp8_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "bf6", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_bf8_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "bf6", "b": "fp6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_fp6_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "bf6", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_fp4_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp4", "b": "fp8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_fp8_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp4", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_bf8_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp4", "b": "fp6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_fp6_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp4", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 32, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_bf6_scale_e8m0_e8m0_k32"},
+        {"family": "wmma_scaled", "a": "fp8", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp8_bf8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp8", "b": "fp6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp8_fp6_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp8", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp8_bf6_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp8", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp8_fp4_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "bf8", "b": "fp8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_fp8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "bf8", "b": "fp6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_fp6_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "bf8", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_bf6_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "bf8", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf8_fp4_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp6", "b": "fp8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp6_fp8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp6", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp6_bf8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp6", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp6_bf6_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp6", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp6_fp4_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "bf6", "b": "fp8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_fp8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "bf6", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_bf8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "bf6", "b": "fp6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_fp6_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "bf6", "b": "fp4", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_bf6_fp4_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp4", "b": "fp8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_fp8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp4", "b": "bf8", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_bf8_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp4", "b": "fp6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_fp6_scale_e8m0_e8m0_k16"},
+        {"family": "wmma_scaled", "a": "fp4", "b": "bf6", "c": "fp32", "a_scale_dtype": "e8m0", "b_scale_dtype": "e8m0", "scale_block_k": 16, "m": 16, "n": 16, "k": 128, "op_id": "wmma_gfx1250_f32_16x16x128_fp4_bf6_scale_e8m0_e8m0_k16"}
       ]
     },
     "gfx11-generic": {

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/target.py

@@ -1006,6 +1006,20 @@ class _FragInfo:
 }
 
 
+# Mixed operands keep the same padded ABI and scale-coordinate maps.
+_MMA_FRAGMENT_INFO.update(
+    {
+        f"wmma_gfx1250_f32_16x16x128_{a}_{b}_scale_e8m0_e8m0_k{block_k}": _MMA_FRAGMENT_INFO[
+            f"wmma_gfx1250_f32_16x16x128_fp8_fp8_scale_e8m0_e8m0_k{block_k}"
+        ]
+        for block_k in (32, 16)
+        for a in ("fp8", "bf8", "fp6", "bf6", "fp4")
+        for b in ("fp8", "bf8", "fp6", "bf6", "fp4")
+        if a != b
+    }
+)
+
+
 def _frag_info(op_id: str) -> _FragInfo:
     info = _MMA_FRAGMENT_INFO.get(op_id)
     if info is None:

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx1250/gemm/README.md

@@ -14,6 +14,7 @@ LLVM 23 ROCm libraries, NumPy, and `ml_dtypes`. Torch is optional.
 | `mxfp8_gemm` | FP8 E4M3 (`fp8`) or BF8 E5M2 (`bf8`) on both operands | E8M0 |
 | `mxfp6_gemm` | Packed FP6 E2M3 or E3M2 on both operands | E8M0 |
 | `mxfp4_gemm` | Packed FP4 E2M1 on both operands | E8M0 |
+| [`mixed_scaled_gemm`](mixed_scaled_gemm.py) | Different A/B formats from FP8, BF8, FP6, BF6, and FP4 | E8M0 |
 
 ```sh
 ROCKE_LLVM_FLAVOR=llvm23 python -m rocke.examples.gfx1250.gemm.mxfp8_gemm
@@ -54,6 +55,17 @@ or `--dtype both` explicitly. Both operands use the selected format with E8M0
 scales. `mxfp6_gemm` similarly defaults to both homogeneous FP6 encodings;
 use `--dtype fp6` or `--dtype bf6` to select one. Mixed matrix formats are outside these examples.
 
+[`mixed_scaled_gemm`](mixed_scaled_gemm.py) defaults to FP8 E4M3 x FP4 E2M1.
+Use `--dtype-a` and `--dtype-b` to select either operand independently; equal
+normalized formats are rejected with guidance to the homogeneous examples.
+All 20 ordered mixed pairs have SCALE and SCALE16 catalog entries. The builder
+uses each operand's atom-selected layout and shared packed-storage helpers.
+
+```sh
+ROCKE_LLVM_FLAVOR=llvm23 python -m rocke.examples.gfx1250.gemm.mixed_scaled_gemm --case all
+ROCKE_LLVM_FLAVOR=llvm23 python -m rocke.examples.gfx1250.gemm.mixed_scaled_gemm --dtype-a fp6 --dtype-b fp4 --matrix-path wmma_scale16 --compile-route hip
+```
+
 The target-independent aliases in `core/dtypes.py` normalize logical format
 names. Recognition does not imply that a target supports an atom; use its
 catalog to query supported operand formats and shapes. Each example defines

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx1250/gemm/block_scaled_gemm_verify.py

@@ -3,7 +3,7 @@
 """Launch gfx1250 block-scaled GEMM and compare with an independent reference.
 
 The default invocation keeps the K=64 FP8/BF8 WMMA + FP32-scale verifier.
-Native ``--matrix-path wmma_scale`` / ``wmma_scale16`` use homogeneous FP8, FP6, or FP4 and
+Native ``--matrix-path wmma_scale`` / ``wmma_scale16`` use independent FP8/BF8, FP6/BF6, or FP4 operands and
 E8M0 scales with K=32 / K=16 groups. Native fixtures cover K=128 or 256 and use
 bounded dyadic values, permitting exact comparison after output-type rounding.
 Numerical equality treats positive and negative zero as equal; output zero-sign
@@ -364,7 +364,7 @@ def run_cases(
                 dtype_c=spec.dtype_c,
             )
             label = (
-                f"{spec.resolved_matrix_path()}/{spec.dtype_a}/{compile_route}/{case} "
+                f"{spec.resolved_matrix_path()}/{spec.dtype_a}x{spec.dtype_b}/{compile_route}/{case} "
                 f"{spec.M}x{spec.N}x{spec.K} bk{spec.block_k}"
             )
             got = _launch(rt, fn, spec, inputs)

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx1250/gemm/mixed_scaled_gemm.py

@@ -0,0 +1,50 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+"""GEMM with different A/B matrix formats and E8M0 scales on gfx1250."""
+
+from __future__ import annotations
+
+from ....core.dtypes import normalize_dtype
+from ....instances.gfx1250.block_scaled_gemm import BlockScaledGemmSpec, _LOWBIT_DTYPES
+from ._scaled_gemm_example import argument_parser, verify
+
+FORMATS = tuple(sorted(_LOWBIT_DTYPES))
+
+
+def make_spec(args) -> BlockScaledGemmSpec:
+    if normalize_dtype(args.dtype_a) == normalize_dtype(args.dtype_b):
+        raise ValueError(
+            "choose different A/B formats or use a homogeneous mxfp example"
+        )
+    return BlockScaledGemmSpec(
+        name="mixed_scaled_gemm",
+        M=args.m,
+        N=args.n,
+        K=args.k,
+        dtype_a=args.dtype_a,
+        dtype_b=args.dtype_b,
+        dtype_c=args.output_dtype,
+        scale_dtype="e8m0",
+        matrix_path=args.matrix_path,
+        block_k=16 if args.matrix_path == "wmma_scale16" else 32,
+    )
+
+
+def main(argv: list[str] | None = None) -> int:
+    parser = argument_parser(__doc__)
+    parser.add_argument(
+        "--dtype-a", type=normalize_dtype, choices=FORMATS, default="fp8e4m3"
+    )
+    parser.add_argument(
+        "--dtype-b", type=normalize_dtype, choices=FORMATS, default="fp4e2m1"
+    )
+    args = parser.parse_args(argv)
+    try:
+        spec = make_spec(args)
+    except ValueError as exc:
+        parser.error(str(exc))
+    return verify(spec, args)
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md

@@ -44,7 +44,7 @@ These emit generic AMDGPU IR; arch only sets the comgr target triple.
 
 ## gfx1250 native scaled GEMM
 
-`block_scaled_gemm` supports homogeneous FP8 E4M3, BF8 E5M2, FP6 E2M3, FP6 E3M2,
+`block_scaled_gemm` supports homogeneous and mixed FP8 E4M3, BF8 E5M2, FP6 E2M3, FP6 E3M2,
 and FP4 E2M1 through SCALE and SCALE16 with LLVM 23 and E8M0 scales.
 See the [packed FP6 input contract](../examples/gfx1250/gemm/FP6.md).
 

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/gfx1250/block_scaled_gemm.py

@@ -179,8 +179,6 @@ def is_valid_spec(spec: BlockScaledGemmSpec, arch: str = "gfx1250") -> Tuple[boo
         )
     matrix_path = spec.resolved_matrix_path()
     native_scale = matrix_path in ("wmma_scale", "wmma_scale16")
-    if native_scale and _canon_lowbit(spec.dtype_a) != _canon_lowbit(spec.dtype_b):
-        return False, "native gfx1250 SCALE/SCALE16 requires matching matrix formats"
     family = matrix_path if native_scale else "wmma"
     atom_k = _WMMA_SCALE_K if native_scale else _WMMA_K
     if not native_scale and not target.mma.has_shape(
@@ -238,7 +236,8 @@ def is_valid_spec(spec: BlockScaledGemmSpec, arch: str = "gfx1250") -> Tuple[boo
     if native_scale:
         return (
             True,
-            f"ok: gfx1250 K=128 native {matrix_path} {_canon_lowbit(spec.dtype_a)} GEMM",
+            f"ok: gfx1250 K=128 native {matrix_path} "
+            f"{_canon_lowbit(spec.dtype_a)}x{_canon_lowbit(spec.dtype_b)} GEMM",
         )
     return True, "ok: gfx1250 K=64 FP8/BF8 WMMA block-scaled GEMM"