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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
