[clang] [llvm] [AMDGPU] v_cvt_scale_pk8_* are block32 in strict mode (PR #227426)
Stanislav Mekhanoshin via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 30 11:53:36 PDT 2026
https://github.com/rampitec updated https://github.com/llvm/llvm-project/pull/227426
>From 58547c5b1f4fd744bb628a7dd1aba7b829823070 Mon Sep 17 00:00:00 2001
From: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
Date: Tue, 29 Sep 2026 11:29:34 -0700
Subject: [PATCH] [AMDGPU] v_cvt_scale_pk8_* arr block32 in strict mode
Reenable these instructions in strict mode.
Fixes: LCOMPILER-2841
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 18 ++++++-------
.../builtins-amdgcn-error-gfx1250-strict.cl | 9 -------
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 19 ++++++-------
llvm/lib/Target/AMDGPU/VOP3Instructions.td | 27 ++++++++++---------
llvm/test/MC/AMDGPU/gfx1250-strict_err.s | 27 -------------------
5 files changed, 33 insertions(+), 67 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index c2506743a698f..ea02122a5452a 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -1065,15 +1065,15 @@ def __builtin_amdgcn_cvt_pk_fp8_f16 : AMDGPUBuiltin<"short(_ExtVector<2, _Float1
def __builtin_amdgcn_cvt_pk_bf8_f16 : AMDGPUBuiltin<"short(_ExtVector<2, _Float16>)", [Const], "gfx1250-insts">;
def __builtin_amdgcn_cvt_sr_fp8_f16 : AMDGPUBuiltin<"int(_Float16, int, unsigned int, _Constant int)", [Const], "gfx1250-insts">;
def __builtin_amdgcn_cvt_sr_bf8_f16 : AMDGPUBuiltin<"int(_Float16, int, unsigned int, _Constant int)", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_f16_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_bf16_fp8 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_f16_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_bf16_bf8 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_f16_fp4 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(unsigned int, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_bf16_fp4 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(unsigned int, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_f32_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_f32_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
-def __builtin_amdgcn_cvt_scale_pk8_f32_fp4 : AMDGPUBuiltin<"_ExtVector<8, float>(unsigned int, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_f16_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_bf16_fp8 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_f16_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_bf16_bf8 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_f16_fp4 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(unsigned int, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_bf16_fp4 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(unsigned int, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_f32_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_f32_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cvt_scale_pk8_f32_fp4 : AMDGPUBuiltin<"_ExtVector<8, float>(unsigned int, unsigned int, _Constant unsigned int)", [Const], "gfx1250-insts">;
def __builtin_amdgcn_cvt_scale_pk16_f16_fp6 : AMDGPUBuiltin<"_ExtVector<16, _Float16>(_ExtVector<3, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
def __builtin_amdgcn_cvt_scale_pk16_bf16_fp6 : AMDGPUBuiltin<"_ExtVector<16, __bf16>(_ExtVector<3, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
def __builtin_amdgcn_cvt_scale_pk16_f16_bf6 : AMDGPUBuiltin<"_ExtVector<16, _Float16>(_ExtVector<3, unsigned int>, unsigned int, _Constant unsigned int)", [Const], "block16-cvt-scale-insts">;
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-strict.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-strict.cl
index e8e2392ab8a59..1dc90a4db9e4a 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-strict.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-strict.cl
@@ -14,15 +14,6 @@ void test(global v16f* out, v16i a, v8i b, v16f c, int scale_src0, int scale_src
*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}}
- (void)__builtin_amdgcn_cvt_scale_pk8_f16_fp8(src2, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_f16_fp8' needs target feature block16-cvt-scale-insts}}
- (void)__builtin_amdgcn_cvt_scale_pk8_bf16_fp8(src2, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_bf16_fp8' needs target feature block16-cvt-scale-insts}}
- (void)__builtin_amdgcn_cvt_scale_pk8_f16_bf8(src2, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_f16_bf8' needs target feature block16-cvt-scale-insts}}
- (void)__builtin_amdgcn_cvt_scale_pk8_bf16_bf8(src2, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_bf16_bf8' needs target feature block16-cvt-scale-insts}}
- (void)__builtin_amdgcn_cvt_scale_pk8_f16_fp4(src1, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_f16_fp4' needs target feature block16-cvt-scale-insts}}
- (void)__builtin_amdgcn_cvt_scale_pk8_bf16_fp4(src1, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_bf16_fp4' needs target feature block16-cvt-scale-insts}}
- (void)__builtin_amdgcn_cvt_scale_pk8_f32_fp8(src2, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_f32_fp8' needs target feature block16-cvt-scale-insts}}
- (void)__builtin_amdgcn_cvt_scale_pk8_f32_bf8(src2, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_f32_bf8' needs target feature block16-cvt-scale-insts}}
- (void)__builtin_amdgcn_cvt_scale_pk8_f32_fp4(src1, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk8_f32_fp4' needs target feature block16-cvt-scale-insts}}
(void)__builtin_amdgcn_cvt_scale_pk16_f16_fp6(src3, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk16_f16_fp6' needs target feature block16-cvt-scale-insts}}
(void)__builtin_amdgcn_cvt_scale_pk16_bf16_fp6(src3, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk16_bf16_fp6' needs target feature block16-cvt-scale-insts}}
(void)__builtin_amdgcn_cvt_scale_pk16_f16_bf6(src3, scale, 0); // expected-error {{'__builtin_amdgcn_cvt_scale_pk16_f16_bf6' needs target feature block16-cvt-scale-insts}}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 82e60e5da872e..f283acee34b04 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -689,16 +689,17 @@ class AMDGPUCvtScaleF32SRIntrinsic<LLVMType DstTy, LLVMType Src0Ty, string name>
[DstTy], [Src0Ty, llvm_i32_ty, llvm_float_ty]
>, ClangBuiltin<"__builtin_amdgcn_"#name>;
+def int_amdgcn_cvt_scale_pk8_f16_fp8 : AMDGPUCvtScaleIntrinsic<llvm_v8f16_ty, llvm_v2i32_ty, "cvt_scale_pk8_f16_fp8">;
+def int_amdgcn_cvt_scale_pk8_bf16_fp8 : AMDGPUCvtScaleIntrinsic<llvm_v8bf16_ty, llvm_v2i32_ty, "cvt_scale_pk8_bf16_fp8">;
+def int_amdgcn_cvt_scale_pk8_f16_bf8 : AMDGPUCvtScaleIntrinsic<llvm_v8f16_ty, llvm_v2i32_ty, "cvt_scale_pk8_f16_bf8">;
+def int_amdgcn_cvt_scale_pk8_bf16_bf8 : AMDGPUCvtScaleIntrinsic<llvm_v8bf16_ty, llvm_v2i32_ty, "cvt_scale_pk8_bf16_bf8">;
+def int_amdgcn_cvt_scale_pk8_f16_fp4 : AMDGPUCvtScaleIntrinsic<llvm_v8f16_ty, llvm_i32_ty, "cvt_scale_pk8_f16_fp4">;
+def int_amdgcn_cvt_scale_pk8_bf16_fp4 : AMDGPUCvtScaleIntrinsic<llvm_v8bf16_ty, llvm_i32_ty, "cvt_scale_pk8_bf16_fp4">;
+def int_amdgcn_cvt_scale_pk8_f32_fp8 : AMDGPUCvtScaleIntrinsic<llvm_v8f32_ty, llvm_v2i32_ty, "cvt_scale_pk8_f32_fp8">;
+def int_amdgcn_cvt_scale_pk8_f32_bf8 : AMDGPUCvtScaleIntrinsic<llvm_v8f32_ty, llvm_v2i32_ty, "cvt_scale_pk8_f32_bf8">;
+def int_amdgcn_cvt_scale_pk8_f32_fp4 : AMDGPUCvtScaleIntrinsic<llvm_v8f32_ty, llvm_i32_ty, "cvt_scale_pk8_f32_fp4">;
+
let TargetFeatures = "block16-cvt-scale-insts" in {
- def int_amdgcn_cvt_scale_pk8_f16_fp8 : AMDGPUCvtScaleIntrinsic<llvm_v8f16_ty, llvm_v2i32_ty, "cvt_scale_pk8_f16_fp8">;
- def int_amdgcn_cvt_scale_pk8_bf16_fp8 : AMDGPUCvtScaleIntrinsic<llvm_v8bf16_ty, llvm_v2i32_ty, "cvt_scale_pk8_bf16_fp8">;
- def int_amdgcn_cvt_scale_pk8_f16_bf8 : AMDGPUCvtScaleIntrinsic<llvm_v8f16_ty, llvm_v2i32_ty, "cvt_scale_pk8_f16_bf8">;
- def int_amdgcn_cvt_scale_pk8_bf16_bf8 : AMDGPUCvtScaleIntrinsic<llvm_v8bf16_ty, llvm_v2i32_ty, "cvt_scale_pk8_bf16_bf8">;
- def int_amdgcn_cvt_scale_pk8_f16_fp4 : AMDGPUCvtScaleIntrinsic<llvm_v8f16_ty, llvm_i32_ty, "cvt_scale_pk8_f16_fp4">;
- def int_amdgcn_cvt_scale_pk8_bf16_fp4 : AMDGPUCvtScaleIntrinsic<llvm_v8bf16_ty, llvm_i32_ty, "cvt_scale_pk8_bf16_fp4">;
- def int_amdgcn_cvt_scale_pk8_f32_fp8 : AMDGPUCvtScaleIntrinsic<llvm_v8f32_ty, llvm_v2i32_ty, "cvt_scale_pk8_f32_fp8">;
- def int_amdgcn_cvt_scale_pk8_f32_bf8 : AMDGPUCvtScaleIntrinsic<llvm_v8f32_ty, llvm_v2i32_ty, "cvt_scale_pk8_f32_bf8">;
- def int_amdgcn_cvt_scale_pk8_f32_fp4 : AMDGPUCvtScaleIntrinsic<llvm_v8f32_ty, llvm_i32_ty, "cvt_scale_pk8_f32_fp4">;
def int_amdgcn_cvt_scale_pk16_f16_bf6 : AMDGPUCvtScaleIntrinsic<llvm_v16f16_ty, llvm_v3i32_ty, "cvt_scale_pk16_f16_bf6">;
def int_amdgcn_cvt_scale_pk16_bf16_bf6 : AMDGPUCvtScaleIntrinsic<llvm_v16bf16_ty, llvm_v3i32_ty, "cvt_scale_pk16_bf16_bf6">;
def int_amdgcn_cvt_scale_pk16_f16_fp6 : AMDGPUCvtScaleIntrinsic<llvm_v16f16_ty, llvm_v3i32_ty, "cvt_scale_pk16_f16_fp6">;
diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
index 9bb5ee7b6e29e..39a8a2bd3aac3 100644
--- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
@@ -1872,14 +1872,6 @@ let SubtargetPredicate = isGFX1250Plus in {
let SubtargetPredicate = HasBlock16ConversionScaleInsts in {
let Constraints = "@earlyclobber $vdst" in {
- let WaveSizePredicate = isWave32, GFX1250BlockingCycles = 4, HasGFX1250BlockingCycles = true in {
- defm V_CVT_SCALE_PK8_F16_FP8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f16_fp8", VOP_V8F16_V2I32_I32, int_amdgcn_cvt_scale_pk8_f16_fp8>;
- defm V_CVT_SCALE_PK8_BF16_FP8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_bf16_fp8", VOP_V8BF16_V2I32_I32, int_amdgcn_cvt_scale_pk8_bf16_fp8>;
- defm V_CVT_SCALE_PK8_F16_BF8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f16_bf8", VOP_V8F16_V2I32_I32, int_amdgcn_cvt_scale_pk8_f16_bf8>;
- defm V_CVT_SCALE_PK8_BF16_BF8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_bf16_bf8", VOP_V8BF16_V2I32_I32, int_amdgcn_cvt_scale_pk8_bf16_bf8>;
- defm V_CVT_SCALE_PK8_F32_FP8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f32_fp8", VOP_V8F32_V2I32_I32, int_amdgcn_cvt_scale_pk8_f32_fp8>;
- defm V_CVT_SCALE_PK8_F32_BF8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f32_bf8", VOP_V8F32_V2I32_I32, int_amdgcn_cvt_scale_pk8_f32_bf8>;
- }
let GFX1250BlockingCycles = 8, HasGFX1250BlockingCycles = true in {
defm V_CVT_SCALE_PK16_F16_FP6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk16_f16_fp6", VOP_V16F16_V3I32_I32, int_amdgcn_cvt_scale_pk16_f16_fp6>;
defm V_CVT_SCALE_PK16_BF16_FP6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk16_bf16_fp6", VOP_V16BF16_V3I32_I32, int_amdgcn_cvt_scale_pk16_bf16_fp6>;
@@ -1889,13 +1881,22 @@ let SubtargetPredicate = isGFX1250Plus in {
defm V_CVT_SCALE_PK16_F32_BF6 : VOP3CvtScaleSelInst<"v_cvt_scale_pk16_f32_bf6", VOP_V16F32_V3I32_I32, int_amdgcn_cvt_scale_pk16_f32_bf6>;
}
} // End Constraints = "@earlyclobber $vdst"
+ } // End SubtargetPredicate = HasBlock16ConversionScaleInsts
- let WaveSizePredicate = isWave32, GFX1250BlockingCycles = 4, HasGFX1250BlockingCycles = true in {
- defm V_CVT_SCALE_PK8_F16_FP4 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f16_fp4", VOP_V8F16_I32_I32, int_amdgcn_cvt_scale_pk8_f16_fp4>;
- defm V_CVT_SCALE_PK8_BF16_FP4 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_bf16_fp4", VOP_V8BF16_I32_I32, int_amdgcn_cvt_scale_pk8_bf16_fp4>;
- defm V_CVT_SCALE_PK8_F32_FP4 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f32_fp4", VOP_V8F32_I32_I32, int_amdgcn_cvt_scale_pk8_f32_fp4>;
+ let WaveSizePredicate = isWave32, GFX1250BlockingCycles = 4, HasGFX1250BlockingCycles = true in {
+ let Constraints = "@earlyclobber $vdst" in {
+ defm V_CVT_SCALE_PK8_F16_FP8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f16_fp8", VOP_V8F16_V2I32_I32, int_amdgcn_cvt_scale_pk8_f16_fp8>;
+ defm V_CVT_SCALE_PK8_BF16_FP8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_bf16_fp8", VOP_V8BF16_V2I32_I32, int_amdgcn_cvt_scale_pk8_bf16_fp8>;
+ defm V_CVT_SCALE_PK8_F16_BF8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f16_bf8", VOP_V8F16_V2I32_I32, int_amdgcn_cvt_scale_pk8_f16_bf8>;
+ defm V_CVT_SCALE_PK8_BF16_BF8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_bf16_bf8", VOP_V8BF16_V2I32_I32, int_amdgcn_cvt_scale_pk8_bf16_bf8>;
+ defm V_CVT_SCALE_PK8_F32_FP8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f32_fp8", VOP_V8F32_V2I32_I32, int_amdgcn_cvt_scale_pk8_f32_fp8>;
+ defm V_CVT_SCALE_PK8_F32_BF8 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f32_bf8", VOP_V8F32_V2I32_I32, int_amdgcn_cvt_scale_pk8_f32_bf8>;
}
- } // End SubtargetPredicate = HasBlock16ConversionScaleInsts
+
+ defm V_CVT_SCALE_PK8_F16_FP4 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f16_fp4", VOP_V8F16_I32_I32, int_amdgcn_cvt_scale_pk8_f16_fp4>;
+ defm V_CVT_SCALE_PK8_BF16_FP4 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_bf16_fp4", VOP_V8BF16_I32_I32, int_amdgcn_cvt_scale_pk8_bf16_fp4>;
+ defm V_CVT_SCALE_PK8_F32_FP4 : VOP3CvtScaleSelInst<"v_cvt_scale_pk8_f32_fp4", VOP_V8F32_I32_I32, int_amdgcn_cvt_scale_pk8_f32_fp4>;
+ } // End WaveSizePredicate = isWave32, GFX1250BlockingCycles = 4, HasGFX1250BlockingCycles = true
} // End ReadsModeReg = 0
let Constraints = "@earlyclobber $vdst" in {
diff --git a/llvm/test/MC/AMDGPU/gfx1250-strict_err.s b/llvm/test/MC/AMDGPU/gfx1250-strict_err.s
index c394380db7c11..12e20a8ffbdc7 100644
--- a/llvm/test/MC/AMDGPU/gfx1250-strict_err.s
+++ b/llvm/test/MC/AMDGPU/gfx1250-strict_err.s
@@ -10,24 +10,6 @@ v_wmma_scale_f32_32x16x128_f4 v[0:15], v[8:23], v[0:7], v[0:15], s0, s0
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
-v_cvt_scale_pk8_f16_fp8 v[10:13], v[20:21], v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_f16_fp8
-
-v_cvt_scale_pk8_bf16_fp8 v[10:13], v[20:21], v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_bf16_fp8
-
-v_cvt_scale_pk8_f16_bf8 v[10:13], v[20:21], v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_f16_bf8
-
-v_cvt_scale_pk8_bf16_bf8 v[10:13], v[20:21], v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_bf16_bf8
-
-v_cvt_scale_pk8_f32_fp8 v[10:17], v[20:21], v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_f32_fp8
-
-v_cvt_scale_pk8_f32_bf8 v[10:17], v[20:21], v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_f32_bf8
-
v_cvt_scale_pk16_f16_fp6 v[10:17], v[20:22], v8
// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk16_f16_fp6
@@ -45,12 +27,3 @@ v_cvt_scale_pk16_f32_fp6 v[10:25], v[20:22], v8
v_cvt_scale_pk16_f32_bf6 v[10:25], v[20:22], v8
// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk16_f32_bf6
-
-v_cvt_scale_pk8_f16_fp4 v[10:13], v20, v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_f16_fp4
-
-v_cvt_scale_pk8_bf16_fp4 v[10:13], v20, v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_bf16_fp4
-
-v_cvt_scale_pk8_f32_fp4 v[10:17], v20, v8
-// GFX1250-ERR: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx1250-strict): v_cvt_scale_pk8_f32_fp4
More information about the llvm-commits
mailing list