[clang] [llvm] [AMDGPU] Disable wmma_*_f4 instructions for gfx1250-strict (PR #222434)
Stanislav Mekhanoshin via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 9 12:53:23 PDT 2026
https://github.com/rampitec updated https://github.com/llvm/llvm-project/pull/222434
>From 6d916fab7565848fbc3834b4a6fa8a94b0a9d744 Mon Sep 17 00:00:00 2001
From: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
Date: Wed, 9 Sep 2026 12:24:16 -0700
Subject: [PATCH] [AMDGPU] Disable wmma_*_f4 instructions for gfx1250-strict
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 8 ++---
.../include/clang/Basic/BuiltinsAMDGPUDocs.td | 26 ++++++++++++--
.../builtins-amdgcn-error-gfx1250-strict.cl | 14 ++++++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 5 ++-
llvm/lib/Target/AMDGPU/AMDGPU.td | 9 ++++-
llvm/lib/Target/AMDGPU/VOP3PInstructions.td | 34 ++++++++++++-------
llvm/test/MC/AMDGPU/gfx1250-strict_err.s | 11 ++++++
7 files changed, 87 insertions(+), 20 deletions(-)
create mode 100644 clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-strict.cl
create mode 100644 llvm/test/MC/AMDGPU/gfx1250-strict_err.s
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 4e4e0ee275a16..9eed0f60a6c50 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -1239,16 +1239,16 @@ def __builtin_amdgcn_wmma_f16_16x16x32_f16 : AMDGPUBuiltin<"_ExtVector<8, _Float
let Documentation = [DocWMMA_half_16x16x32_GFX1250];
let ArgNames = ["a_neg", "a", "b_neg", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"];
}
-def __builtin_amdgcn_wmma_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>)", [Const], "wmma-n16-insts,wavefrontsize32"> {
+def __builtin_amdgcn_wmma_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>)", [Const], "wmma-f4-insts,wavefrontsize32"> {
let Documentation = [DocWMMA_f4_GFX1250];
let ArgNames = ["a", "b", "c_mod", "c"];
}
-def __builtin_amdgcn_wmma_scale_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, int, _Constant int, _Constant int, int, _Constant bool, _Constant bool)", [Const], "wmma-n16-insts,wavefrontsize32"> {
+def __builtin_amdgcn_wmma_scale_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, int, _Constant int, _Constant int, int, _Constant bool, _Constant bool)", [Const], "wmma-f4-insts,wavefrontsize32"> {
let Documentation = [DocWMMA_scale_GFX1250];
let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"];
}
-def __builtin_amdgcn_wmma_scale16_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, long int, _Constant int, _Constant int, long int, _Constant bool, _Constant bool)", [Const], "wmma-n16-insts,wavefrontsize32"> {
- let Documentation = [DocWMMA_scale16_GFX1250];
+def __builtin_amdgcn_wmma_scale16_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, long int, _Constant int, _Constant int, long int, _Constant bool, _Constant bool)", [Const], "wmma-f4-insts,wavefrontsize32"> {
+ let Documentation = [DocWMMA_scale16_f4_GFX1250];
let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"];
}
def __builtin_amdgcn_swmmac_f32_16x16x64_bf16 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<32, __bf16>, _ExtVector<8, float>, int, _Constant bool, _Constant bool)", [Const], "swmmac-gfx1250-insts,wavefrontsize32">;
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
index fe61919958f76..bfccd7de6590a 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
@@ -576,7 +576,7 @@ and adds the 32x16 f32 accumulator ``C``.
- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs,
3 = neg(abs)).
-Requires target feature ``wmma-n16-insts`` and wave32.
+Requires target feature ``wmma-f4-insts`` and wave32.
}];
}
@@ -597,7 +597,7 @@ WMMA with per-operand scale factors applied during the computation.
- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that
matrix A or B data can be reused from a previous WMMA instruction.
-Requires target feature ``wmma-n16-insts`` and wave32.
+Requires target feature ``wmma-f4-insts`` and wave32.
}];
}
@@ -675,6 +675,28 @@ Requires target feature ``wmma-n16-insts`` and wave32.
}];
}
+def DocWMMA_scale16_f4_GFX1250 : Documentation {
+ let Category = DocCatWMMA;
+ let Content = [{
+Scaled wave matrix multiply-accumulate with 64-bit (16-bit per element)
+scale exponents, as opposed to the 32-bit exponents used by the non-16
+scale variant.
+
+- For ``..._f8f6f4`` variants, ``matrix_a_fmt`` / ``matrix_b_fmt`` are format
+ selectors for A and B. They must be compile-time constants.
+- ``c_mod``: accumulator modifier.
+- ``matrix_a_scale`` / ``matrix_b_scale``: scale factor selectors.
+ Must be compile-time constants.
+- ``matrix_a_scale_fmt`` / ``matrix_b_scale_fmt``: scale format selectors.
+ Must be compile-time constants.
+- ``matrix_a_scale_exp`` / ``matrix_b_scale_exp``: 64-bit scale exponents.
+- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that
+ matrix A or B data can be reused from a previous WMMA instruction.
+
+Requires target feature ``wmma-f4-insts`` and wave32.
+}];
+}
+
def DocWMMA_f64_16x16x4_f64_GFX1251 : Documentation {
let Category = DocCatWMMA;
let Content = [{
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-strict.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-strict.cl
new file mode 100644
index 0000000000000..a0202d2d76e14
--- /dev/null
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-strict.cl
@@ -0,0 +1,14 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu12.50s-- -verify -S -o - %s
+
+typedef float v16f __attribute__((ext_vector_type(16)));
+typedef float v8f __attribute__((ext_vector_type(8)));
+typedef int v16i __attribute__((ext_vector_type(16)));
+typedef int v8i __attribute__((ext_vector_type(8)));
+
+void test_amdgcn_wmma_f32_wmma_f4(global v16f* out, v16i a, v8i b, v16f c, int scale_src0, int scale_src1)
+{
+ *out = __builtin_amdgcn_wmma_f32_32x16x128_f4(a, b, false, c); // expected-error {{'__builtin_amdgcn_wmma_f32_32x16x128_f4' needs target feature wmma-f4-insts}}
+ *out = __builtin_amdgcn_wmma_scale_f32_32x16x128_f4(a, b, false, c, 1, 0, scale_src0, 2, 0, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale_f32_32x16x128_f4' needs target feature wmma-f4-insts}}
+ *out = __builtin_amdgcn_wmma_scale16_f32_32x16x128_f4(a, b, false, c, 1, 0, scale_src0, 2, 0, scale_src1, 1, 0); // expected-error {{'__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4' needs target feature wmma-f4-insts}}
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 6b15075f75b9c..5769595e2e821 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -4254,10 +4254,13 @@ let TargetFeatures = "wmma-n16-insts" in {
def int_amdgcn_wmma_f32_16x16x128_f8f6f4 : AMDGPUWmmaIntrinsicModsC_MatrixFMT;
def int_amdgcn_wmma_scale_f32_16x16x128_f8f6f4 : AMDGPUWmmaScaleIntrinsicModsC<llvm_i32_ty>;
def int_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4 : AMDGPUWmmaScaleIntrinsicModsC<llvm_i64_ty>;
+} // End TargetFeatures = "wmma-n16-insts"
+
+let TargetFeatures = "wmma-f4-insts" in {
def int_amdgcn_wmma_f32_32x16x128_f4 : AMDGPUWmmaIntrinsicF4ModsC<llvm_anyint_ty, llvm_anyint_ty, llvm_anyfloat_ty>;
def int_amdgcn_wmma_scale_f32_32x16x128_f4 : AMDGPUWmmaScaleF4IntrinsicModsC<llvm_i32_ty>;
def int_amdgcn_wmma_scale16_f32_32x16x128_f4 : AMDGPUWmmaScaleF4IntrinsicModsC<llvm_i64_ty>;
-} // End TargetFeatures = "wmma-n16-insts"
+} // End TargetFeatures = "wmma-f4-insts"
}
class AMDGPUSWmmacIntrinsicABIdx<LLVMType A, LLVMType B, LLVMType CD, LLVMType Index> :
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 3502e65a3f6e6..7bb8cae31cc87 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -960,6 +960,10 @@ defm WMMAN16Insts : AMDGPUSubtargetFeature<"wmma-n16-insts",
"Has WMMA instructions with N = 16"
>;
+defm WMMAF4Insts : AMDGPUSubtargetFeature<"wmma-f4-insts",
+ "Has WMMA_*_F4 instructions"
+>;
+
defm SWMMACGfx1200Insts : AMDGPUSubtargetFeature<"swmmac-gfx1200-insts",
"Has GFX1200 SWMMAC instructions"
>;
@@ -2448,6 +2452,7 @@ def FeatureISAVersion12_50_STRICT : FeatureSet<
!listconcat(FeatureISAVersion12_50_Common.Features,
[FeatureGFX1250_STRICT,
FeatureAddressableLocalMemorySize327680,
+ FeatureWMMAN16Insts,
FeatureVOP3PX2IncrementsVaVdstTwice,
FeatureSetregVGPRMSBFixup,
FeatureCubeInsts,
@@ -2476,6 +2481,7 @@ def FeatureISAVersion12_50 : FeatureSet<
!listconcat(FeatureISAVersion12_50_Common.Features,
[FeatureAddressableLocalMemorySize327680,
FeatureWMMAN16Insts,
+ FeatureWMMAF4Insts,
FeatureVOP3PX2IncrementsVaVdstTwice,
FeatureSetregVGPRMSBFixup,
FeatureCubeInsts,
@@ -2503,6 +2509,7 @@ def FeatureISAVersion12_51 : FeatureSet<
!listconcat(FeatureISAVersion12_50_Common.Features,
[FeatureAddressableLocalMemorySize327680,
FeatureWMMAN16Insts,
+ FeatureWMMAF4Insts,
FeatureFullRate64Ops,
FeatureVOP3PX2IncrementsVaVdstTwice,
FeatureDPALU_DPP,
@@ -3230,7 +3237,7 @@ def AMDGPUFrontendVisibleFeatures {
FeatureSadInsts, FeatureSetPrioIncWgInst, FeatureSmemPrefetchInsts,
FeatureTanhInsts, FeatureTensorCvtLutInsts, FeatureTransposeLoadF4F6Insts,
FeatureVMemToLDSLoad, FeatureVmemPrefInsts, FeatureWMMA128bInsts,
- FeatureWMMA256bInsts, FeatureWMMAN16Insts, FeatureXF32Insts,
+ FeatureWMMA256bInsts, FeatureWMMAN16Insts, FeatureWMMAF4Insts, FeatureXF32Insts,
FeatureWavefrontSize32, FeatureWavefrontSize64, FeatureSupportsWGP,
FeatureSupportsWave32, FeatureFastFMAF32, FeatureFastDenormalF32,
FeatureSupportsXNACK, FeatureSupportsSRAMECC, FeatureXNACKOnOffModes,
diff --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
index 8c516245faff7..a29e6e82c12eb 100644
--- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -2160,15 +2160,18 @@ let WaveSizePredicate = isWave32 in {
defm V_WMMA_F32_16X16X128_FP8_BF8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x128_fp8_bf8", F32_FP8BF8X128_WMMA_w32, "_w32">;
defm V_WMMA_F32_16X16X128_BF8_FP8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x128_bf8_fp8", F32_FP8BF8X128_WMMA_w32, "_w32">;
defm V_WMMA_F32_16X16X128_BF8_BF8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x128_bf8_bf8", F32_FP8BF8X128_WMMA_w32, "_w32">;
- defm V_WMMA_F32_32X16X128_F4_w32 : WMMAInstGFX12<"v_wmma_f32_32x16x128_f4", F32_32X16X128_F4_WMMA_w32, "_w32">;
defm V_WMMA_F32_16X16X128_F8F6F4 : WMMAInst_SrcFormats_mc<"v_wmma_f32_16x16x128_f8f6f4", "F32_16X16X128_F8F6F4">;
defm V_WMMA_SCALE_F32_16X16X128_F8F6F4 : WMMAInst_SrcFormats_mc<"v_wmma_scale_f32_16x16x128_f8f6f4", "F32_16X16X128_F8F6F4_SCALE">;
defm V_WMMA_SCALE16_F32_16X16X128_F8F6F4 : WMMAInst_SrcFormats_mc<"v_wmma_scale16_f32_16x16x128_f8f6f4", "F32_16X16X128_F8F6F4_SCALE16">;
+ } // End SubtargetPredicate = HasWMMAN16Insts
+
+ let SubtargetPredicate = HasWMMAF4Insts in {
+ defm V_WMMA_F32_32X16X128_F4_w32 : WMMAInstGFX12<"v_wmma_f32_32x16x128_f4", F32_32X16X128_F4_WMMA_w32, "_w32">;
defm V_WMMA_SCALE_F32_32X16X128_F4_w32 : WMMAInstGFX12<"v_wmma_scale_f32_32x16x128_f4", F32_32X16X128_F4_SCALE_w32, "_w32">;
defm V_WMMA_SCALE16_F32_32X16X128_F4_w32 : WMMAInstGFX12<"v_wmma_scale16_f32_32x16x128_f4", F32_32X16X128_F4_SCALE16_w32, "_w32">;
- } // End SubtargetPredicate = HasWMMAN16Insts
+ } // End SubtargetPredicate = HasWMMAF4Insts
let SubtargetPredicate = HasSWMMACGfx1250Insts in {
defm V_SWMMAC_F32_16X16X64_BF16_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x64_bf16", F32_BF16X64_SWMMAC_w32, "_w32">;
@@ -2359,9 +2362,6 @@ let SubtargetPredicate = HasWMMAN16Insts in {
defm : WMMAPat<"V_WMMA_F32_16X16X128_FP8_BF8_w32", int_amdgcn_wmma_f32_16x16x128_fp8_bf8, F32_FP8BF8X128_WMMA_w32>;
defm : WMMAPat<"V_WMMA_F32_16X16X128_BF8_FP8_w32", int_amdgcn_wmma_f32_16x16x128_bf8_fp8, F32_FP8BF8X128_WMMA_w32>;
defm : WMMAPat<"V_WMMA_F32_16X16X128_BF8_BF8_w32", int_amdgcn_wmma_f32_16x16x128_bf8_bf8, F32_FP8BF8X128_WMMA_w32>;
- defm : WMMAPat<"V_WMMA_F32_32X16X128_F4_w32", int_amdgcn_wmma_f32_32x16x128_f4, F32_32X16X128_F4_WMMA_w32>;
- defm : WMMAPat<"V_WMMA_SCALE_F32_32X16X128_F4_w32", int_amdgcn_wmma_scale_f32_32x16x128_f4, F32_32X16X128_F4_SCALE_w32>;
- defm : WMMAPat<"V_WMMA_SCALE16_F32_32X16X128_F4_w32", int_amdgcn_wmma_scale16_f32_32x16x128_f4, F32_32X16X128_F4_SCALE16_w32>;
foreach I = ["f8_f8", "f8_f6", "f8_f4", "f6_f8", "f6_f6", "f6_f4", "f4_f8", "f4_f6", "f4_f4"] in {
defm : WMMAPat<"V_WMMA_F32_16X16X128_F8F6F4_" # I # "_w32", int_amdgcn_wmma_f32_16x16x128_f8f6f4, !cast<VOP3PWMMA_Profile>("F32_16X16X128_F8F6F4_" # I # "_w32")>;
@@ -2370,6 +2370,12 @@ let SubtargetPredicate = HasWMMAN16Insts in {
}
} // End SubtargetPredicate = HasWMMAN16Insts
+let SubtargetPredicate = HasWMMAF4Insts in {
+ defm : WMMAPat<"V_WMMA_F32_32X16X128_F4_w32", int_amdgcn_wmma_f32_32x16x128_f4, F32_32X16X128_F4_WMMA_w32>;
+ defm : WMMAPat<"V_WMMA_SCALE_F32_32X16X128_F4_w32", int_amdgcn_wmma_scale_f32_32x16x128_f4, F32_32X16X128_F4_SCALE_w32>;
+ defm : WMMAPat<"V_WMMA_SCALE16_F32_32X16X128_F4_w32", int_amdgcn_wmma_scale16_f32_32x16x128_f4, F32_32X16X128_F4_SCALE16_w32>;
+} // End SubtargetPredicate = HasWMMAF4Insts
+
let SubtargetPredicate = HasSWMMACGfx1250Insts in {
def : SWMMACPat<V_SWMMAC_F32_16X16X64_BF16_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x64_bf16, F32_BF16X64_SWMMAC_w32>;
def : SWMMACPat<V_SWMMAC_BF16_16X16X64_BF16_w32_twoaddr, int_amdgcn_swmmac_bf16_16x16x64_bf16, BF16_BF16X64_SWMMAC_w32>;
@@ -2681,14 +2687,18 @@ defm V_WMMA_F16_16X16X128_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x086, F16_FP8B
defm V_WMMA_F16_16X16X128_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x087, F16_FP8BF8X128_WMMA_w32>;
defm V_WMMA_F32_32X16X128_F4_w32 : VOP3P_Real_WMMA_gfx1250 <0x088, F32_32X16X128_F4_WMMA_w32>;
-let WaveSizePredicate = isWave32, SubtargetPredicate = HasWMMAN16Insts, DecoderNamespace = "GFX1250" in {
-defm V_WMMA_F32_16X16X128_F8F6F4 : VOP3P_Real_WMMA_SrcFormats <"gfx1250", 0x033, "F32_16X16X128_F8F6F4">;
-defm V_WMMA_SCALE_F32_16X16X128_F8F6F4 : VOP3PX2_Real_ScaledWMMA_SrcFormats <"gfx1250", 0x033, 0x35, "F32_16X16X128_F8F6F4_SCALE">;
-defm V_WMMA_SCALE16_F32_16X16X128_F8F6F4 : VOP3PX2_Real_ScaledWMMA_SrcFormats <"gfx1250", 0x033, 0x3a, "F32_16X16X128_F8F6F4_SCALE16">;
+let WaveSizePredicate = isWave32, DecoderNamespace = "GFX1250" in {
+ let SubtargetPredicate = HasWMMAN16Insts in {
+ defm V_WMMA_F32_16X16X128_F8F6F4 : VOP3P_Real_WMMA_SrcFormats <"gfx1250", 0x033, "F32_16X16X128_F8F6F4">;
+ defm V_WMMA_SCALE_F32_16X16X128_F8F6F4 : VOP3PX2_Real_ScaledWMMA_SrcFormats <"gfx1250", 0x033, 0x35, "F32_16X16X128_F8F6F4_SCALE">;
+ defm V_WMMA_SCALE16_F32_16X16X128_F8F6F4 : VOP3PX2_Real_ScaledWMMA_SrcFormats <"gfx1250", 0x033, 0x3a, "F32_16X16X128_F8F6F4_SCALE16">;
+ }
-defm V_WMMA_SCALE_F32_32X16X128_F4_w32 : VOP3PX2_Real_ScaledWMMA_F4 <"gfx1250", 0x088, 0x35, F32_32X16X128_F4_SCALE_w32>;
-defm V_WMMA_SCALE16_F32_32X16X128_F4_w32 : VOP3PX2_Real_ScaledWMMA_F4 <"gfx1250", 0x088, 0x3a, F32_32X16X128_F4_SCALE16_w32>;
-} // End WaveSizePredicate = isWave32, SubtargetPredicate = HasWMMAN16Insts, DecoderNamespace = "GFX1250"
+ let SubtargetPredicate = HasWMMAF4Insts in {
+ defm V_WMMA_SCALE_F32_32X16X128_F4_w32 : VOP3PX2_Real_ScaledWMMA_F4 <"gfx1250", 0x088, 0x35, F32_32X16X128_F4_SCALE_w32>;
+ defm V_WMMA_SCALE16_F32_32X16X128_F4_w32 : VOP3PX2_Real_ScaledWMMA_F4 <"gfx1250", 0x088, 0x3a, F32_32X16X128_F4_SCALE16_w32>;
+ }
+} // End WaveSizePredicate = isWave32, DecoderNamespace = "GFX1250"
let SubtargetPredicate = HasSWMMACGfx1250Insts in {
defm V_SWMMAC_F32_16X16X64_F16_w32 : VOP3P_Real_WMMA_gfx1250 <0x065, F32_F16X64_SWMMAC_w32>;
diff --git a/llvm/test/MC/AMDGPU/gfx1250-strict_err.s b/llvm/test/MC/AMDGPU/gfx1250-strict_err.s
new file mode 100644
index 0000000000000..52683e859ceb4
--- /dev/null
+++ b/llvm/test/MC/AMDGPU/gfx1250-strict_err.s
@@ -0,0 +1,11 @@
+// NOTE: Assertions have been autogenerated by utils/update_mc_test_checks.py UTC_ARGS: --version 6
+// RUN: not llvm-mc -triple=amdgpu12.50s -filetype=null %s 2>&1 | FileCheck --check-prefixes=GFX1250-ERR --implicit-check-not=error: -strict-whitespace %s
+
+v_wmma_f32_32x16x128_f4 v[0:15], v[8:23], v[0:7], v[0:15]
+// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_wmma_f32_32x16x128_f4
+
+v_wmma_scale_f32_32x16x128_f4 v[0:15], v[8:23], v[0:7], v[0:15], s0, s0
+// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_wmma_scale_f32_32x16x128_f4
+
+v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[8:23], v[0:7], v[0:15], s[0:1], s[0:1]
+// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_wmma_scale16_f32_32x16x128_f4
More information about the cfe-commits
mailing list