feat(rocke): add gfx1250 FP8 scaled WMMA slice

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/arch/data.cpp

@@ -1017,6 +1017,8 @@ static const rocke_ati_mma_frag_row_t rocke_ati_mma_frag[] = {
     {"wmma_gfx1250_f32_16x16x64_fp8_bf8", 8},
     {"wmma_gfx1250_f32_16x16x64_bf8_fp8", 8},
     {"wmma_gfx1250_f32_16x16x64_bf8_bf8", 8},
+    {"wmma_scale_f32_16x16x128_fp8_fp8", 8},
+    {"wmma_scale16_f32_16x16x128_fp8_fp8", 8},
 };
 
 int rocke_arch_mma_c_frag_len(const char* op_id)
@@ -1799,6 +1801,36 @@ static const rocke_mma_op_t k_mma_gfx1250[] = {
      NULL,
      NULL,
      &lm_wmma_gfx12_c},
+    {"wmma_scale",
+     "fp8",
+     "fp8",
+     "fp32",
+     16,
+     16,
+     128,
+     "wmma_scale_f32_16x16x128_fp8_fp8",
+     16,
+     16,
+     8,
+     32,
+     NULL,
+     NULL,
+     &lm_wmma_gfx12_c},
+    {"wmma_scale16",
+     "fp8",
+     "fp8",
+     "fp32",
+     16,
+     16,
+     128,
+     "wmma_scale16_f32_16x16x128_fp8_fp8",
+     16,
+     16,
+     8,
+     32,
+     NULL,
+     NULL,
+     &lm_wmma_gfx12_c},
 };
 
 /* =========================================================================

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp

@@ -49,6 +49,8 @@ static const rocke_mma_hint_row_t ROCKE_MMA_RESULT_HINT[] = {
     {"mfma_f32_16x16x96_fp6", "acc6"},
     {"mfma_f32_16x16x128_fp8", "acc128"},
     {"mfma_scale_f32_16x16x128_f8f6f4", "mxacc"},
+    {"wmma_scale_f32_16x16x128_fp8_fp8", "mxacc"},
+    {"wmma_scale16_f32_16x16x128_fp8_fp8", "mxacc"},
 };
 
 /* Accumulator fragment length for op_id, from the arch SSOT
@@ -651,6 +653,32 @@ rocke_value_t* rocke_b_mfma_scale_f32_16x16x128_f8f6f4(rocke_ir_builder_t* b,
     return rocke_b_mma(b, "mfma_scale_f32_16x16x128_f8f6f4", a, bb, c, extra, 2);
 }
 
+rocke_value_t* rocke_b_wmma_scale_f32_16x16x128_fp8_fp8(rocke_ir_builder_t* b,
+                                                        rocke_value_t* a,
+                                                        rocke_value_t* bb,
+                                                        rocke_value_t* c,
+                                                        rocke_value_t* a_scale,
+                                                        rocke_value_t* b_scale)
+{
+    rocke_value_t* extra[2];
+    extra[0] = a_scale;
+    extra[1] = b_scale;
+    return rocke_b_mma(b, "wmma_scale_f32_16x16x128_fp8_fp8", a, bb, c, extra, 2);
+}
+
+rocke_value_t* rocke_b_wmma_scale16_f32_16x16x128_fp8_fp8(rocke_ir_builder_t* b,
+                                                          rocke_value_t* a,
+                                                          rocke_value_t* bb,
+                                                          rocke_value_t* c,
+                                                          rocke_value_t* a_scale,
+                                                          rocke_value_t* b_scale)
+{
+    rocke_value_t* extra[2];
+    extra[0] = a_scale;
+    extra[1] = b_scale;
+    return rocke_b_mma(b, "wmma_scale16_f32_16x16x128_fp8_fp8", a, bb, c, extra, 2);
+}
+
 /* ===================================================================== */
 /*  register-fragment reshape (P13)                                       */
 /* ===================================================================== */

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp

@@ -76,11 +76,12 @@ const char* const ROCKE_HIP_PROLOGUE
       "_ROCKE_VEC(float, f32x, 1); _ROCKE_VEC(float, f32x, 2); _ROCKE_VEC(float, f32x, 4);\n"
       "_ROCKE_VEC(float, f32x, 8); _ROCKE_VEC(float, f32x, 16);\n"
       "_ROCKE_VEC(int, i32x, 1); _ROCKE_VEC(int, i32x, 2); _ROCKE_VEC(int, i32x, 3);\n"
-      "_ROCKE_VEC(int, i32x, 4); _ROCKE_VEC(int, i32x, 8);\n"
+      "_ROCKE_VEC(int, i32x, 4); _ROCKE_VEC(int, i32x, 8); _ROCKE_VEC(int, i32x, 16);\n"
       "_ROCKE_VEC(int16_t, i16x, 1); _ROCKE_VEC(int16_t, i16x, 2);\n"
       "_ROCKE_VEC(int16_t, i16x, 4); _ROCKE_VEC(int16_t, i16x, 8);\n"
       "_ROCKE_VEC(int8_t, i8x, 1); _ROCKE_VEC(int8_t, i8x, 2);\n"
       "_ROCKE_VEC(int8_t, i8x, 4); _ROCKE_VEC(int8_t, i8x, 8); _ROCKE_VEC(int8_t, i8x, 16);\n"
+      "_ROCKE_VEC(int8_t, i8x, 32); _ROCKE_VEC(int8_t, i8x, 48); _ROCKE_VEC(int8_t, i8x, 64);\n"
       "_ROCKE_VEC(bool, boolx, 2); _ROCKE_VEC(bool, boolx, 4); _ROCKE_VEC(bool, boolx, 8);\n"
       "_ROCKE_VEC(bool, boolx, 16);\n"
       "#undef _ROCKE_VEC\n"

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_cast.cpp

@@ -32,16 +32,41 @@ static const rocke_value_t* h_res(const rocke_op_t* op)
 
 /* ----------------------------- arith: int width casts ---------------------- */
 
-/* def _op_arith_zext(self, op):
- *     (v,) = op.operands
- *     self._emit(f"{_type_to_hip(op.result.type)} {_name(op.result)} = "
- *                f"({_type_to_hip(op.result.type)}){_name(v)};") */
+static const char* h_unsigned_int_type(const rocke_type_t* t)
+{
+    const char* name = t ? t->name : NULL;
+    if(!name)
+        return NULL;
+    if(__builtin_strcmp(name, "i1") == 0)
+        return "bool";
+    if(__builtin_strcmp(name, "i8") == 0)
+        return "uint8_t";
+    if(__builtin_strcmp(name, "i16") == 0)
+        return "uint16_t";
+    if(__builtin_strcmp(name, "i32") == 0)
+        return "uint32_t";
+    if(__builtin_strcmp(name, "i64") == 0)
+        return "uint64_t";
+    return NULL;
+}
+
+/* Zero extension must cast through the unsigned spelling of the source type.
+ * A direct int8_t -> int/int64_t cast sign-extends bytes 0x80..0xff. */
 static rocke_status_t rocke_h_op_arith_zext(rocke_h_lowerer_t* lw, const rocke_op_t* op)
 {
     const rocke_value_t* v = op->operands[0];
     const rocke_value_t* r = h_res(op);
     const char* t = rocke_h_type_to_hip(lw, r->type);
-    rocke_h_emitf(lw, "%s %s = (%s)%s;", t, rocke_h_name(lw, r), t, rocke_h_name(lw, v));
+    const char* source_unsigned = h_unsigned_int_type(v->type);
+    if(!source_unsigned)
+    {
+        return rocke_h_fail(lw,
+                            ROCKE_ERR_NOTIMPL,
+                            "HIP zext requires an integer scalar source, got '%s'",
+                            v->type ? v->type->name : "(null)");
+    }
+    rocke_h_emitf(
+        lw, "%s %s = (%s)(%s)%s;", t, rocke_h_name(lw, r), t, source_unsigned, rocke_h_name(lw, v));
     return lw->status;
 }
 

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp

@@ -75,6 +75,38 @@ static const char* h_elem_scalar(const rocke_type_t* t)
  * op.result (and the WMMA gate keys off the op_id *string* it is passed, not
  * op.attrs), so a synthetic op aliasing the same operands/results/regions and
  * reusing the original attrs map reproduces the Python emission exactly. */
+static rocke_status_t
+    h_emit_gfx1250_scaled_wmma(rocke_h_lowerer_t* lw, const rocke_op_t* op, bool scale16)
+{
+    const char* op_id
+        = scale16 ? "wmma_scale16_f32_16x16x128_fp8_fp8" : "wmma_scale_f32_16x16x128_fp8_fp8";
+    const char* builtin = scale16 ? "__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4"
+                                  : "__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4";
+    if(!lw->arch.gfx || __builtin_strcmp(lw->arch.gfx, "gfx1250") != 0)
+    {
+        return rocke_h_fail(lw,
+                            ROCKE_ERR_NOTIMPL,
+                            "WMMA op '%s' is not available on %s",
+                            op_id,
+                            lw->arch.gfx ? lw->arch.gfx : "(unknown)");
+    }
+    if(op->num_operands != 5 || op->num_results != 1)
+    {
+        return rocke_h_fail(lw, ROCKE_ERR_VALUE, "%s expects 5 operands and 1 result", op_id);
+    }
+    rocke_h_emitf(lw,
+                  "f32x8 %s = %s(0, %s, 0, %s, (int16_t)0, %s, "
+                  "0, 0, %s, 0, 0, %s, false, false);",
+                  rocke_h_name(lw, op->results[0]),
+                  builtin,
+                  rocke_h_name(lw, op->operands[0]),
+                  rocke_h_name(lw, op->operands[1]),
+                  rocke_h_name(lw, op->operands[2]),
+                  rocke_h_name(lw, op->operands[3]),
+                  rocke_h_name(lw, op->operands[4]));
+    return lw->status;
+}
+
 static rocke_status_t rocke_h_op_tile_mma(rocke_h_lowerer_t* lw, const rocke_op_t* op)
 {
     const char* op_id;
@@ -91,6 +123,14 @@ static rocke_status_t rocke_h_op_tile_mma(rocke_h_lowerer_t* lw, const rocke_op_
     {
         return rocke_h_fail(lw, ROCKE_ERR_KEY, "tile.mma: missing 'op_id' attr");
     }
+    if(__builtin_strcmp(op_id, "wmma_scale_f32_16x16x128_fp8_fp8") == 0)
+    {
+        return h_emit_gfx1250_scaled_wmma(lw, op, false);
+    }
+    if(__builtin_strcmp(op_id, "wmma_scale16_f32_16x16x128_fp8_fp8") == 0)
+    {
+        return h_emit_gfx1250_scaled_wmma(lw, op, true);
+    }
     snprintf(dotted, sizeof(dotted), "tile.%s", op_id);
     legacy_opcode = rocke_opcode_from_name(dotted);
     fn = rocke_h_dispatch(legacy_opcode);

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp

@@ -180,6 +180,14 @@ const rocke_ll_decl_t ROCKE_LL_INTRINSIC_DECLS[] = {
     {"wmma.gfx1250.f32.16x16x64.bf8.bf8",
      "declare <8 x float> @llvm.amdgcn.wmma.f32.16x16x64.bf8.bf8.v8f32.v8i32(<8 x i32>, <8 x i32>, "
      "i16 immarg, <8 x float>, i1 immarg, i1 immarg)"},
+    {"wmma.scale.gfx1250.f32.16x16x128.fp8.fp8",
+     "declare <8 x float> @llvm.amdgcn.wmma.scale.f32.16x16x128.f8f6f4.v8f32.v16i32.v16i32(i32 "
+     "immarg, <16 x i32>, i32 immarg, <16 x i32>, i16 immarg, <8 x float>, i32 immarg, i32 immarg, "
+     "i32, i32 immarg, i32 immarg, i32, i1 immarg, i1 immarg)"},
+    {"wmma.scale16.gfx1250.f32.16x16x128.fp8.fp8",
+     "declare <8 x float> @llvm.amdgcn.wmma.scale16.f32.16x16x128.f8f6f4.v8f32.v16i32.v16i32(i32 "
+     "immarg, <16 x i32>, i32 immarg, <16 x i32>, i16 immarg, <8 x float>, i32 immarg, i32 immarg, "
+     "i64, i32 immarg, i32 immarg, i64, i1 immarg, i1 immarg)"},
     {"mfma.f32.16x16x16f16",
      "declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, "
      "i32 immarg, i32 immarg, i32 immarg)"},

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mma.cpp

@@ -35,6 +35,7 @@ namespace ckc
 static void _op_tile_wmma_f32_16x16x16_f16(rocke_lower_t* L, const rocke_op_t* op);
 static void _op_tile_wmma_f32_16x16x16_bf16(rocke_lower_t* L, const rocke_op_t* op);
 static void _emit_wmma(rocke_lower_t* L, const rocke_op_t* op, const char* op_id);
+static void _emit_wmma_scale(rocke_lower_t* L, const rocke_op_t* op, bool scale16);
 static void _op_tile_mma(rocke_lower_t* L, const rocke_op_t* op);
 static void _op_tile_mfma_f32_16x16x32_fp8(rocke_lower_t* L, const rocke_op_t* op);
 static void _op_tile_mfma_f32_16x16x32_bf8(rocke_lower_t* L, const rocke_op_t* op);
@@ -133,6 +134,14 @@ static void _op_tile_mma(rocke_lower_t* L, const rocke_op_t* op)
          * reject (Python ISABackend.emit_wmma raises NotImplementedError). The
          * gfx12-specific op_ids ("wmma_gfx12_*") can only occur on RDNA4. */
     }
+    else if(strcmp(op_id, "wmma_scale_f32_16x16x128_fp8_fp8") == 0)
+    {
+        _emit_wmma_scale(L, op, false);
+    }
+    else if(strcmp(op_id, "wmma_scale16_f32_16x16x128_fp8_fp8") == 0)
+    {
+        _emit_wmma_scale(L, op, true);
+    }
     else if(strncmp(op_id, "wmma_", 5) == 0)
     {
         if(L->backend && L->backend->kind == ROCKE_LL_ISA_RDNA)
@@ -164,6 +173,84 @@ static void _op_tile_mma(rocke_lower_t* L, const rocke_op_t* op)
     }
 }
 
+/* gfx1250 native MX FP8 WMMA. ROCm 7.13 introduced the LLVM 23 ABI:
+ * matrix operands are <16 x i32>; SCALE carries four packed E8M0 bytes in
+ * each i32 scale operand and SCALE16 carries eight in i64. */
+static void _emit_wmma_scale(rocke_lower_t* L, const rocke_op_t* op, bool scale16)
+{
+    const char* intrinsic;
+    const char* decl_key;
+    const char* scale_ty;
+    const char* op_name;
+
+    if(!rocke_ll_live(L))
+    {
+        return;
+    }
+    op_name = scale16 ? "tile.wmma_scale16_f32_16x16x128_fp8_fp8"
+                      : "tile.wmma_scale_f32_16x16x128_fp8_fp8";
+    if(!L->backend || strcmp(L->backend->gfx, "gfx1250") != 0)
+    {
+        rocke_ll_fail(L,
+                      ROCKE_ERR_NOTIMPL,
+                      "%s is only available on gfx1250 (got %s)",
+                      op_name,
+                      L->backend ? L->backend->gfx : "(unknown)");
+    }
+    if(L->flavor != ROCKE_LLVM_FLAVOR_LLVM23)
+    {
+        rocke_ll_fail(L,
+                      ROCKE_ERR_NOTIMPL,
+                      "%s requires llvm23 (ROCm 7.13+), got %s",
+                      op_name,
+                      rocke_llvm_flavor_name(L->flavor));
+    }
+    if(op->num_operands != 5)
+    {
+        rocke_ll_fail(
+            L, ROCKE_ERR_VALUE, "%s expects 5 operands, got %d", op_name, op->num_operands);
+    }
+
+    scale_ty = scale16 ? "i64" : "i32";
+    if(strcmp(op->operands[3]->type->name, scale_ty) != 0
+       || strcmp(op->operands[4]->type->name, scale_ty) != 0)
+    {
+        rocke_ll_fail(L,
+                      ROCKE_ERR_VALUE,
+                      "%s expects %s scale operands, got %s/%s",
+                      op_name,
+                      scale_ty,
+                      op->operands[3]->type->name,
+                      op->operands[4]->type->name);
+    }
+
+    if(scale16)
+    {
+        decl_key = "wmma.scale16.gfx1250.f32.16x16x128.fp8.fp8";
+        intrinsic = "llvm.amdgcn.wmma.scale16.f32.16x16x128.f8f6f4.v8f32.v16i32.v16i32";
+    }
+    else
+    {
+        decl_key = "wmma.scale.gfx1250.f32.16x16x128.fp8.fp8";
+        intrinsic = "llvm.amdgcn.wmma.scale.f32.16x16x128.f8f6f4.v8f32.v16i32.v16i32";
+    }
+    rocke_ll_need(L, decl_key);
+    rocke_ll_emitf(L,
+                   "  %s = call <8 x float> @%s("
+                   "i32 0, <16 x i32> %s, i32 0, <16 x i32> %s, "
+                   "i16 0, <8 x float> %s, i32 0, i32 0, %s %s, "
+                   "i32 0, i32 0, %s %s, i1 false, i1 false)",
+                   mma_result_name(L, op),
+                   intrinsic,
+                   rocke_ll_operand(L, op->operands[0]),
+                   rocke_ll_operand(L, op->operands[1]),
+                   rocke_ll_operand(L, op->operands[2]),
+                   scale_ty,
+                   rocke_ll_operand(L, op->operands[3]),
+                   scale_ty,
+                   rocke_ll_operand(L, op->operands[4]));
+}
+
 /* ====================================================================== */
 /* WMMA on CDNA: faithful rejection (Python ISABackend.emit_wmma raises)   */
 /* ====================================================================== */

dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h

@@ -1036,6 +1036,18 @@ rocke_value_t* rocke_b_mfma_scale_f32_16x16x128_f8f6f4(rocke_ir_builder_t* b,
                                                        rocke_value_t* c,
                                                        rocke_value_t* a_scale,
                                                        rocke_value_t* b_scale);
+rocke_value_t* rocke_b_wmma_scale_f32_16x16x128_fp8_fp8(rocke_ir_builder_t* b,
+                                                        rocke_value_t* a,
+                                                        rocke_value_t* bb,
+                                                        rocke_value_t* c,
+                                                        rocke_value_t* a_scale,
+                                                        rocke_value_t* b_scale);
+rocke_value_t* rocke_b_wmma_scale16_f32_16x16x128_fp8_fp8(rocke_ir_builder_t* b,
+                                                          rocke_value_t* a,
+                                                          rocke_value_t* bb,
+                                                          rocke_value_t* c,
+                                                          rocke_value_t* a_scale,
+                                                          rocke_value_t* b_scale);
 rocke_value_t* rocke_b_wmma_f32_16x16x16_f16(rocke_ir_builder_t* b,
                                              rocke_value_t* a,
                                              rocke_value_t* bb,

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

@@ -142,7 +142,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 only the validated fp16/bf16 and fp8/bf8 WMMA atoms; block-scaled F4/block16 forms are 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 the FP8xFP8 SCALE/SCALE16 prerequisite slice; 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.",
       "family": "cdna",
       "target_family": "gfx12_cdna",
       "matrix_path": "wmma",
@@ -176,7 +176,9 @@
         {"family": "wmma", "a": "fp8", "b": "fp8", "c": "fp32", "m": 16, "n": 16, "k": 64, "op_id": "wmma_gfx1250_f32_16x16x64_fp8_fp8"},
         {"family": "wmma", "a": "fp8", "b": "bf8", "c": "fp32", "m": 16, "n": 16, "k": 64, "op_id": "wmma_gfx1250_f32_16x16x64_fp8_bf8"},
         {"family": "wmma", "a": "bf8", "b": "fp8", "c": "fp32", "m": 16, "n": 16, "k": 64, "op_id": "wmma_gfx1250_f32_16x16x64_bf8_fp8"},
-        {"family": "wmma", "a": "bf8", "b": "bf8", "c": "fp32", "m": 16, "n": 16, "k": 64, "op_id": "wmma_gfx1250_f32_16x16x64_bf8_bf8"}
+        {"family": "wmma", "a": "bf8", "b": "bf8", "c": "fp32", "m": 16, "n": 16, "k": 64, "op_id": "wmma_gfx1250_f32_16x16x64_bf8_bf8"},
+        {"family": "wmma_scale", "a": "fp8", "b": "fp8", "c": "fp32", "m": 16, "n": 16, "k": 128, "op_id": "wmma_scale_f32_16x16x128_fp8_fp8"},
+        {"family": "wmma_scale16", "a": "fp8", "b": "fp8", "c": "fp32", "m": 16, "n": 16, "k": 128, "op_id": "wmma_scale16_f32_16x16x128_fp8_fp8"}
       ]
     },
     "gfx11-generic": {