[llvm-branch-commits] [llvm] [AMDGPU] Guard more intrinsics with target features (PR #203956)
Shilei Tian via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Jun 22 09:27:40 PDT 2026
https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/203956
>From 14693610cb5ae5563e67e7a51312d43a726c7516 Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
Date: Mon, 15 Jun 2026 12:09:29 -0400
Subject: [PATCH] [AMDGPU] Guard more intrinsics with target features
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 17 +++++-
.../AMDGPU/AMDGPUInstructionSelector.cpp | 52 +------------------
.../lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp | 24 ---------
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 42 ---------------
.../inst-select-amdgcn.exp.compr.mir | 3 +-
.../GlobalISel/llvm.amdgcn.intersect_ray.ll | 4 +-
.../AMDGPU/llvm.amdgcn.dual_intersect_ray.ll | 6 +--
.../CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll | 3 +-
.../CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll | 2 +-
.../AMDGPU/llvm.amdgcn.intersect_ray.ll | 6 ++-
.../CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll | 2 +-
.../CodeGen/AMDGPU/llvm.amdgcn.permlane.ll | 2 +-
.../AMDGPU/llvm.amdgcn.permlane16.swap.ll | 2 +-
.../test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll | 2 +-
.../test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll | 2 +-
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll | 2 +-
llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll | 5 ++
.../CodeGen/AMDGPU/unsupported-av-load.ll | 8 +--
.../CodeGen/AMDGPU/unsupported-av-store.ll | 8 +--
19 files changed, 48 insertions(+), 144 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index a484efb0c1e36..eb42a4da74898 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -318,6 +318,7 @@ def int_amdgcn_s_barrier_join : ClangBuiltin<"__builtin_amdgcn_s_barrier_join">,
// void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %barrier)
// The %barrier argument must be uniform, otherwise behavior is undefined.
+let TargetFeatures = "s-wakeup-barrier-inst" in
def int_amdgcn_s_wakeup_barrier : ClangBuiltin<"__builtin_amdgcn_s_wakeup_barrier">,
Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn,
IntrNoCallback, IntrNoFree]>;
@@ -475,6 +476,7 @@ def int_amdgcn_fmul_legacy : ClangBuiltin<"__builtin_amdgcn_fmul_legacy">,
// intended for use on subtargets that have the v_fma_legacy_f32 and/or
// v_fmac_legacy_f32 instructions. (Note that v_fma_legacy_f16 is unrelated and
// has a completely different kind of legacy behaviour.)
+let TargetFeatures = "gfx10-3-insts" in
def int_amdgcn_fma_legacy :
PureIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
[Commutative]
@@ -586,6 +588,7 @@ def int_amdgcn_fmad_ftz :
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>]
>;
+let TargetFeatures = "tanh-insts" in
def int_amdgcn_tanh : PureIntrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>]
>;
@@ -915,8 +918,6 @@ class AMDGPUAVStore : Intrinsic <
[SDNPMemOperand, SDNPMayStore]
>;
-def int_amdgcn_av_store_b128 : AMDGPUAVStore;
-
class AMDGPUAVLoad : Intrinsic <
[llvm_v4i32_ty],
[llvm_anyptr_ty, // Pointer to load from (flat or global)
@@ -927,7 +928,10 @@ class AMDGPUAVLoad : Intrinsic <
[SDNPMemOperand, SDNPMayLoad]
>;
+let TargetFeatures = "flat-global-insts" in {
+def int_amdgcn_av_store_b128 : AMDGPUAVStore;
def int_amdgcn_av_load_b128 : AMDGPUAVLoad;
+} // End TargetFeatures = "flat-global-insts"
} // TargetPrefix = "amdgcn"
@@ -2159,6 +2163,7 @@ def int_amdgcn_exp_row : DefaultAttrsIntrinsic <[], [
>;
// exp with compr bit set. Not supported on GFX11+.
+let TargetFeatures = "-gfx11-insts" in
def int_amdgcn_exp_compr : DefaultAttrsIntrinsic <[], [
llvm_i32_ty, // tgt,
llvm_i32_ty, // en
@@ -2881,6 +2886,7 @@ def int_amdgcn_wait_asyncmark :
// GFX10 Intrinsics
//===----------------------------------------------------------------------===//
+let TargetFeatures = "permlane16-insts" in {
// llvm.amdgcn.permlane16 <old> <src0> <src1> <src2> <fi> <bound_control>
def int_amdgcn_permlane16 :
Intrinsic<[llvm_any_ty],
@@ -2894,10 +2900,12 @@ def int_amdgcn_permlanex16 :
[LLVMMatchType<0>, LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
[IntrNoMem, IntrConvergent, IntrWillReturn,
ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>, IntrNoCallback, IntrNoFree]>;
+} // End TargetFeatures = "permlane16-insts"
// llvm.amdgcn.mov.dpp8 <src> <sel>
// <sel> is a 32-bit constant whose high 8 bits must be zero which selects
// the lanes to read from.
+let TargetFeatures = "dpp8" in
def int_amdgcn_mov_dpp8 :
Intrinsic<[llvm_any_ty],
[LLVMMatchType<0>, llvm_i32_ty],
@@ -2920,6 +2928,7 @@ class AMDGPUAtomicRtn<LLVMType vt, LLVMType pt = llvm_anyptr_ty> : Intrinsic <
// <ray_dir>, <ray_inv_dir>, <texture_descr>
// <node_ptr> is i32 or i64.
// <ray_dir> and <ray_inv_dir> are both v3f16 or both v3f32.
+let TargetFeatures = "gfx10_a-encoding" in
def int_amdgcn_image_bvh_intersect_ray :
DefaultAttrsIntrinsic<[llvm_v4i32_ty],
[llvm_anyint_ty, llvm_float_ty, llvm_v3f32_ty, llvm_anyvector_ty,
@@ -3063,6 +3072,7 @@ def int_amdgcn_ds_bvh_stack_push8_pop1_rtn : IntDSBVHStackRtn<vdst = llvm_i32_t
def int_amdgcn_ds_bvh_stack_push8_pop2_rtn : IntDSBVHStackRtn<vdst = llvm_i64_ty,
data1 = llvm_v8i32_ty>;
+let TargetFeatures = "bvh-dual-bvh-8-insts" in {
// <vdata>, <ray_origin>, <ray_dir>
// llvm.amdgcn.image.bvh.dual.intersect.ray <node_ptr>, <ray_extent>,
// <instance_mask>, <ray_origin>,
@@ -3084,6 +3094,7 @@ def int_amdgcn_image_bvh8_intersect_ray :
[llvm_i64_ty, llvm_float_ty, llvm_i8_ty, llvm_v3f32_ty,
llvm_v3f32_ty, llvm_i32_ty, llvm_v4i32_ty],
[IntrReadMem, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
+} // End TargetFeatures = "bvh-dual-bvh-8-insts"
// llvm.amdgcn.permlane16.var <old> <src0> <src1> <fi> <bound_control>
def int_amdgcn_permlane16_var : ClangBuiltin<"__builtin_amdgcn_permlane16_var">,
@@ -3403,6 +3414,7 @@ def int_amdgcn_udot4 :
// a[i in 0. . . 3] = (%a_sign ? a.i8[i] : promoteToSigned(a.u8[i]));
// b[i in 0. . . 3] = (%b_sign ? b.i8[i] : promoteToSigned(b.u8[i]));
// %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c
+let TargetFeatures = "dot8-insts" in
def int_amdgcn_sudot4 :
ClangBuiltin<"__builtin_amdgcn_sudot4">,
PureIntrinsic<
@@ -3456,6 +3468,7 @@ def int_amdgcn_udot8 :
// b[i in 0. . . 7] = (%b_sign ? b.i4[i] : promoteToSigned(b.u4[i]));
// %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] +
// %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c
+ let TargetFeatures = "dot8-insts" in
def int_amdgcn_sudot8 :
ClangBuiltin<"__builtin_amdgcn_sudot8">,
PureIntrinsic<
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
index c85d7124f62bb..d52237f1dc6e4 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
@@ -72,12 +72,6 @@ static Register getWaveAddress(const MachineInstr *Def) {
: Register();
}
-static void diagnoseUnsupportedIntrinsic(const MachineInstr &I) {
- const Function &F = I.getMF()->getFunction();
- F.getContext().diagnose(DiagnosticInfoUnsupported(
- F, "intrinsic not supported on subtarget", I.getDebugLoc(), DS_Error));
-}
-
bool AMDGPUInstructionSelector::isVCC(Register Reg,
const MachineRegisterInfo &MRI) const {
// The verifier is oblivious to s1 being a valid value for wavesize registers.
@@ -1283,38 +1277,6 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC(MachineInstr &I) const {
return selectPermlaneSwapIntrin(I, IntrinsicID);
case Intrinsic::amdgcn_wave_shuffle:
return selectWaveShuffleIntrin(I);
- case Intrinsic::amdgcn_fma_legacy:
- if (!STI.hasFmaLegacy32Insts()) {
- diagnoseUnsupportedIntrinsic(I);
- return false;
- }
- return selectImpl(I, *CoverageInfo);
- case Intrinsic::amdgcn_sudot4:
- case Intrinsic::amdgcn_sudot8:
- if (!STI.hasDot8Insts()) {
- diagnoseUnsupportedIntrinsic(I);
- return false;
- }
- return selectImpl(I, *CoverageInfo);
- case Intrinsic::amdgcn_permlane16:
- case Intrinsic::amdgcn_permlanex16:
- if (!STI.hasPermlane16Insts()) {
- diagnoseUnsupportedIntrinsic(I);
- return false;
- }
- return selectImpl(I, *CoverageInfo);
- case Intrinsic::amdgcn_mov_dpp8:
- if (!STI.hasDPP8()) {
- diagnoseUnsupportedIntrinsic(I);
- return false;
- }
- return selectImpl(I, *CoverageInfo);
- case Intrinsic::amdgcn_tanh:
- if (!STI.hasTanhInsts()) {
- diagnoseUnsupportedIntrinsic(I);
- return false;
- }
- return selectImpl(I, *CoverageInfo);
default:
return selectImpl(I, *CoverageInfo);
}
@@ -2497,12 +2459,6 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS(
if (!Subtarget->hasAsyncMark())
return false;
break;
- case Intrinsic::amdgcn_exp_compr:
- if (!STI.hasCompressedExport()) {
- diagnoseUnsupportedIntrinsic(I);
- return false;
- }
- break;
case Intrinsic::amdgcn_ds_bvh_stack_rtn:
case Intrinsic::amdgcn_ds_bvh_stack_push4_pop1_rtn:
case Intrinsic::amdgcn_ds_bvh_stack_push8_pop1_rtn:
@@ -2527,13 +2483,7 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS(
case Intrinsic::amdgcn_s_barrier_init:
case Intrinsic::amdgcn_s_barrier_signal_var:
return selectNamedBarrierInit(I, IntrinsicID);
- case Intrinsic::amdgcn_s_wakeup_barrier: {
- if (!STI.hasSWakeupBarrier()) {
- diagnoseUnsupportedIntrinsic(I);
- return false;
- }
- return selectNamedBarrierInst(I, IntrinsicID);
- }
+ case Intrinsic::amdgcn_s_wakeup_barrier:
case Intrinsic::amdgcn_s_barrier_join:
case Intrinsic::amdgcn_s_get_named_barrier_state:
return selectNamedBarrierInst(I, IntrinsicID);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
index 65ae77cec0846..84428ee046863 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
@@ -7814,13 +7814,6 @@ bool AMDGPULegalizerInfo::legalizeBVHIntersectRayIntrinsic(
Register RayInvDir = MI.getOperand(6).getReg();
Register TDescr = MI.getOperand(7).getReg();
- if (!ST.hasGFX10_AEncoding()) {
- Function &Fn = B.getMF().getFunction();
- Fn.getContext().diagnose(DiagnosticInfoUnsupported(
- Fn, "intrinsic not supported on subtarget", MI.getDebugLoc()));
- return false;
- }
-
const bool IsGFX11 = AMDGPU::isGFX11(ST);
const bool IsGFX11Plus = AMDGPU::isGFX11Plus(ST);
const bool IsGFX12Plus = AMDGPU::isGFX12Plus(ST);
@@ -7966,13 +7959,6 @@ bool AMDGPULegalizerInfo::legalizeBVHDualOrBVH8IntersectRayIntrinsic(
Register Offsets = MI.getOperand(9).getReg();
Register TDescr = MI.getOperand(10).getReg();
- if (!ST.hasBVHDualAndBVH8Insts()) {
- Function &Fn = B.getMF().getFunction();
- Fn.getContext().diagnose(DiagnosticInfoUnsupported(
- Fn, "intrinsic not supported on subtarget", MI.getDebugLoc()));
- return false;
- }
-
bool IsBVH8 = cast<GIntrinsic>(MI).getIntrinsicID() ==
Intrinsic::amdgcn_image_bvh8_intersect_ray;
const unsigned NumVDataDwords = 10;
@@ -8605,16 +8591,6 @@ bool AMDGPULegalizerInfo::legalizeIntrinsic(LegalizerHelper &Helper,
return true;
case Intrinsic::amdgcn_av_load_b128:
case Intrinsic::amdgcn_av_store_b128: {
- const GCNSubtarget &ST = B.getMF().getSubtarget<GCNSubtarget>();
- if (!ST.hasFlatGlobalInsts()) {
- const char *Name = IntrID == Intrinsic::amdgcn_av_load_b128
- ? "llvm.amdgcn.av.load.b128"
- : "llvm.amdgcn.av.store.b128";
- Function &Fn = B.getMF().getFunction();
- Fn.getContext().diagnose(DiagnosticInfoUnsupported(
- Fn, Twine(Name) + " not supported on subtarget", MI.getDebugLoc()));
- return false;
- }
assert(MI.hasOneMemOperand() && "Expected IRTranslator to set MemOp!");
if (IntrID == Intrinsic::amdgcn_av_load_b128)
B.buildLoad(MI.getOperand(0), MI.getOperand(2), **MI.memoperands_begin());
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 0267f696a8a93..6f90b4b5d3b4d 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -7944,10 +7944,6 @@ static SDValue lowerLaneOp(const SITargetLowering &TLI, SDNode *N,
MVT IntVT = MVT::getIntegerVT(ValSize);
const GCNSubtarget *ST = TLI.getSubtarget();
- if ((IsPermLane16 && !ST->hasPermlane16Insts()) ||
- (IID == Intrinsic::amdgcn_mov_dpp8 && !ST->hasDPP8()))
- return emitRemovedIntrinsicError(DAG, SL, VT);
-
unsigned SplitSize = 32;
if (IID == Intrinsic::amdgcn_update_dpp && (ValSize % 64 == 0) &&
ST->hasDPALU_DPP() &&
@@ -10857,17 +10853,9 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
return emitRemovedIntrinsicError(DAG, DL, VT);
return DAG.getNode(AMDGPUISD::RCP_LEGACY, DL, VT, Op.getOperand(1));
case Intrinsic::amdgcn_fma_legacy:
- if (!Subtarget->hasFmaLegacy32Insts())
- return emitRemovedIntrinsicError(DAG, DL, VT);
- return SDValue();
case Intrinsic::amdgcn_sudot4:
case Intrinsic::amdgcn_sudot8:
- if (!Subtarget->hasDot8Insts())
- return emitRemovedIntrinsicError(DAG, DL, VT);
- return SDValue();
case Intrinsic::amdgcn_tanh:
- if (!Subtarget->hasTanhInsts())
- return emitRemovedIntrinsicError(DAG, DL, VT);
return SDValue();
case Intrinsic::amdgcn_rsq_clamp: {
if (Subtarget->getGeneration() < AMDGPUSubtarget::VOLCANIC_ISLANDS)
@@ -11737,11 +11725,6 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op,
assert(NodePtr.getValueType() == MVT::i64);
assert(RayDir.getValueType() == MVT::v3f32);
- if (!Subtarget->hasBVHDualAndBVH8Insts()) {
- emitRemovedIntrinsicError(DAG, DL, Op.getValueType());
- return SDValue();
- }
-
bool IsBVH8 = IntrID == Intrinsic::amdgcn_image_bvh8_intersect_ray;
const unsigned NumVDataDwords = 10;
const unsigned NumVAddrDwords = IsBVH8 ? 11 : 12;
@@ -11782,11 +11765,6 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op,
assert(RayDir.getValueType() == MVT::v3f16 ||
RayDir.getValueType() == MVT::v3f32);
- if (!Subtarget->hasGFX10_AEncoding()) {
- emitRemovedIntrinsicError(DAG, DL, Op.getValueType());
- return SDValue();
- }
-
const bool IsGFX11 = AMDGPU::isGFX11(*Subtarget);
const bool IsGFX11Plus = AMDGPU::isGFX11Plus(*Subtarget);
const bool IsGFX12Plus = AMDGPU::isGFX12Plus(*Subtarget);
@@ -11980,14 +11958,6 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op,
Chain, Ptr, MII->getMemOperand());
}
case Intrinsic::amdgcn_av_load_b128: {
- if (!Subtarget->hasFlatGlobalInsts()) {
- DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
- DAG.getMachineFunction().getFunction(),
- "llvm.amdgcn.av.load.b128 not supported on subtarget",
- DL.getDebugLoc()));
- return DAG.getMergeValues(
- {DAG.getPOISON(Op->getValueType(0)), Op->getOperand(0)}, DL);
- }
MemIntrinsicSDNode *MII = cast<MemIntrinsicSDNode>(Op);
SDValue Chain = Op->getOperand(0);
SDValue Ptr = Op->getOperand(2);
@@ -12177,11 +12147,6 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
switch (IntrinsicID) {
case Intrinsic::amdgcn_exp_compr: {
- if (!Subtarget->hasCompressedExport()) {
- DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
- DAG.getMachineFunction().getFunction(),
- "intrinsic not supported on subtarget", DL.getDebugLoc()));
- }
SDValue Src0 = Op.getOperand(4);
SDValue Src1 = Op.getOperand(5);
// Hack around illegal type on SI by directly selecting it.
@@ -12703,13 +12668,6 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
Ptr, MII->getMemOperand());
}
case Intrinsic::amdgcn_av_store_b128: {
- if (!Subtarget->hasFlatGlobalInsts()) {
- DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
- DAG.getMachineFunction().getFunction(),
- "llvm.amdgcn.av.store.b128 not supported on subtarget",
- DL.getDebugLoc()));
- return Op->getOperand(0); // return the input chain
- }
MemIntrinsicSDNode *MII = cast<MemIntrinsicSDNode>(Op);
SDValue Chain = Op->getOperand(0);
SDValue Ptr = Op->getOperand(2);
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.exp.compr.mir b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.exp.compr.mir
index de41bdf674508..16428405fcb4d 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.exp.compr.mir
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.exp.compr.mir
@@ -1,7 +1,6 @@
# RUN: llc -mtriple=amdgcn -run-pass=instruction-select -verify-machineinstrs %s -o - | FileCheck %s
-# RUN: not llc -mtriple=amdgcn -mcpu=gfx1100 -run-pass=instruction-select -verify-machineinstrs -global-isel-abort=0 %s -o - 2>&1 | FileCheck --check-prefix=ERR %s
-# ERR: error: <unknown>:0:0: in function exp0 void (): intrinsic not supported on subtarget
+# Note: gfx11+ rejection of llvm.amdgcn.exp.compr is now enforced at IR translation (see llvm.amdgcn.exp.compr.ll), not during instruction selection.
---
name: exp0
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll
index 3c2a65084381d..3d0a7882e11b1 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll
@@ -3,7 +3,7 @@
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx1013 < %s | FileCheck -check-prefixes=GCN,GFX10,GFX1013 %s
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx1100 -mattr=+real-true16 < %s | FileCheck -check-prefixes=GCN,GFX11,GFX11-TRUE16 %s
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx1100 -mattr=-real-true16 < %s | FileCheck -check-prefixes=GCN,GFX11,GFX11-FAKE16 %s
-; RUN: not llc -global-isel -mtriple=amdgcn -mcpu=gfx1012 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx1012 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s
; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr)
; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr)
@@ -22,7 +22,7 @@ define amdgpu_ps <4 x float> @image_bvh_intersect_ray(i32 %node_ptr, float %ray_
; GCN-NEXT: image_bvh_intersect_ray v[0:3], v[0:10], s[0:3]
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: ; return to shader part epilog
-; ERR: in function image_bvh_intersect_ray{{.*}}intrinsic not supported on subtarget
+; ERR: in function image_bvh_intersect_ray{{.*}}requires target feature 'gfx10_a-encoding'
%v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
%r = bitcast <4 x i32> %v to <4 x float>
ret <4 x float> %r
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dual_intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dual_intersect_ray.ll
index 195e72ecaf336..50031f05cf94a 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dual_intersect_ray.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dual_intersect_ray.ll
@@ -1,12 +1,12 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
-; RUN: not --crash llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 < %s | FileCheck -check-prefixes=GFX12-SDAG %s
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 < %s | FileCheck -check-prefixes=GFX12-GISEL %s
declare {<10 x i32>, <3 x float>, <3 x float>} @llvm.amdgcn.image.bvh.dual.intersect.ray(i64, float, i8, <3 x float>, <3 x float>, <2 x i32>, <4 x i32>)
-; ERR: in function image_bvh_dual_intersect_ray{{.*}}intrinsic not supported on subtarget
+; ERR: in function image_bvh_dual_intersect_ray{{.*}}requires target feature 'bvh-dual-bvh-8-insts'
define amdgpu_ps <10 x float> @image_bvh_dual_intersect_ray(i64 %node_ptr, float %ray_extent, float %ray_origin_x, float %ray_origin_y, float %ray_origin_z, float %ray_dir_x, float %ray_dir_y, float %ray_dir_z, <2 x i32> %offsets, <4 x i32> inreg %tdescr, ptr addrspace(1) %origin, ptr addrspace(1) %dir) {
; GFX12-SDAG-LABEL: image_bvh_dual_intersect_ray:
; GFX12-SDAG: ; %bb.0: ; %main_body
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll
index b126631e16dd1..c2b8b9cbf4983 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll
@@ -3,8 +3,9 @@
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -strict-whitespace -check-prefix=GCN %s
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -strict-whitespace -check-prefix=GCN %s
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -strict-whitespace -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -strict-whitespace -check-prefix=ERR %s
-; ERR: error: <unknown>:0:0: in function test_export_compr_zeroes_v2f16 void (): intrinsic not supported on subtarget
+; ERR: error: <unknown>:0:0: in function test_export_compr_zeroes_v2f16 void (): llvm.amdgcn.exp.compr requires target feature '-gfx11-insts'
declare void @llvm.amdgcn.exp.compr.v2f16(i32, i32, <2 x half>, <2 x half>, i1, i1) #0
declare void @llvm.amdgcn.exp.compr.v2i16(i32, i32, <2 x i16>, <2 x i16>, i1, i1) #0
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll
index c1ad188671f64..42879d7600290 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll
@@ -13,7 +13,7 @@
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; ERR: error: <unknown>:0:0: in function v_fma float (float, float, float): intrinsic not supported on subtarget
+; ERR: error: <unknown>:0:0: in function v_fma float (float, float, float): llvm.amdgcn.fma.legacy requires target feature 'gfx10-3-insts'
define float @v_fma(float %a, float %b, float %c) {
; GFX10-LABEL: v_fma:
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll
index e6f2bb49c5726..105272b098ce4 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll
@@ -3,7 +3,6 @@
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1013 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX10,GFX10-GISEL,GFX1013-GISEL %s
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1030 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX10,GFX10-SDAG,GFX1030-SDAG %s
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1030 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX10,GFX10-GISEL,GFX1030-GISEL %s
-; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx1012 < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 -mattr=+real-true16 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX11-SDAG %s
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1100 -mattr=+real-true16 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX11-GISEL %s
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 -mattr=-real-true16 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX11-SDAG %s
@@ -13,6 +12,9 @@
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 -mattr=+real-true16 < %s | FileCheck -check-prefixes=GFX12,GFX12-GISEL %s
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 -mattr=-real-true16 < %s | FileCheck -check-prefixes=GFX12,GFX12-GISEL %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1012 < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx1012 < %s 2>&1 | FileCheck -check-prefix=ERR %s
+
; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr)
; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr)
; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(ulong node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr)
@@ -23,7 +25,7 @@ declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32, float, <3
declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64, float, <3 x float>, <3 x float>, <3 x float>, <4 x i32>)
declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64, float, <3 x float>, <3 x half>, <3 x half>, <4 x i32>)
-; ERR: in function image_bvh_intersect_ray{{.*}}intrinsic not supported on subtarget
+; ERR: in function image_bvh_intersect_ray{{.*}}requires target feature 'gfx10_a-encoding'
; Arguments are flattened to represent the actual VGPR_A layout, so we have no
; extra moves in the generated kernel.
define amdgpu_ps <4 x float> @image_bvh_intersect_ray(i32 %node_ptr, float %ray_extent, float %ray_origin_x, float %ray_origin_y, float %ray_origin_z, float %ray_dir_x, float %ray_dir_y, float %ray_dir_z, float %ray_inv_dir_x, float %ray_inv_dir_y, float %ray_inv_dir_z, <4 x i32> inreg %tdescr) {
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll
index cf8c681df6d5a..4d8f40cb98846 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll
@@ -8,7 +8,7 @@
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx900 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; ERR: error: <unknown>:0:0: in function dpp8_test void (ptr addrspace(1), i32): intrinsic not supported on subtarget
+; ERR: error: <unknown>:0:0: in function dpp8_test void (ptr addrspace(1), i32): llvm.amdgcn.mov.dpp8 requires target feature 'dpp8'
; GFX10PLUS-LABEL: {{^}}dpp8_test:
; GFX10PLUS: v_mov_b32_e32 [[SRC:v[0-9]+]], s{{[0-9]+}}
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane.ll
index 1e6b33bd45a04..d2dfe3af18354 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane.ll
@@ -11,7 +11,7 @@
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx900 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; ERR: error: <unknown>:0:0: in function v_permlane16_b32_vss_i32 void (ptr addrspace(1), i32, i32, i32): intrinsic not supported on subtarget
+; ERR: error: <unknown>:0:0: in function v_permlane16_b32_vss_i32 void (ptr addrspace(1), i32, i32, i32): llvm.amdgcn.permlane16 requires target feature 'permlane16-insts'
declare i32 @llvm.amdgcn.permlane16(i32, i32, i32, i32, i1, i1)
declare i32 @llvm.amdgcn.permlanex16(i32, i32, i32, i32, i1, i1)
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.swap.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.swap.ll
index 7bf1d0b6be982..56bec2e39e423 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.swap.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.swap.ll
@@ -5,7 +5,7 @@
; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s | FileCheck -check-prefix=GFX1250 %s
; RUN: not llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR %s
-; RUN: not llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR %s
; ERR: error: {{.*}}: in function v_permlane16_swap_b32_vv {{.*}}: llvm.amdgcn.permlane16.swap requires target feature 'permlane16-swap'
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll
index 1a9cc7615a518..64877691bfff1 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll
@@ -8,7 +8,7 @@
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; ERR: error: <unknown>:0:0: in function test_llvm_amdgcn_sudot4_uu i32 (i32, i32, i32): intrinsic not supported on subtarget
+; ERR: error: <unknown>:0:0: in function test_llvm_amdgcn_sudot4_uu i32 (i32, i32, i32): llvm.amdgcn.sudot4 requires target feature 'dot8-insts'
declare i32 @llvm.amdgcn.sudot4(i1 %asign, i32 %a, i1 %bsign, i32 %b, i32 %c, i1 %clamp)
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll
index 4c3dd27ae0300..53005220632b3 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll
@@ -8,7 +8,7 @@
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; ERR: error: <unknown>:0:0: in function test_llvm_amdgcn_sudot8_uu i32 (i32, i32, i32): intrinsic not supported on subtarget
+; ERR: error: <unknown>:0:0: in function test_llvm_amdgcn_sudot8_uu i32 (i32, i32, i32): llvm.amdgcn.sudot8 requires target feature 'dot8-insts'
declare i32 @llvm.amdgcn.sudot8(i1 %asign, i32 %a, i1 %bsign, i32 %b, i32 %c, i1 %clamp)
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll
index 9d6c5ea82c8ea..a8dd0381756c0 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll
@@ -11,7 +11,7 @@
; RUN: not llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR %s
-; ERR: error: <unknown>:0:0: in function tanh_f32 void (ptr addrspace(1), float): intrinsic not supported on subtarget
+; ERR: error: <unknown>:0:0: in function tanh_f32 void (ptr addrspace(1), float): llvm.amdgcn.tanh requires target feature 'tanh-insts'
; FIXME: GlobalISel does not work with bf16
diff --git a/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll b/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll
index bc995f4738af8..e3c6f7a082494 100644
--- a/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll
+++ b/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll
@@ -2,6 +2,11 @@
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX1250-SDAG %s
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX1250-GISEL %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx1200 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
+
+; ERR: error: <unknown>:0:0: in function kernel1 void (ptr addrspace(1), ptr addrspace(3)): llvm.amdgcn.s.wakeup.barrier requires target feature 's-wakeup-barrier-inst'
+
@bar = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison
define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) #0 {
diff --git a/llvm/test/CodeGen/AMDGPU/unsupported-av-load.ll b/llvm/test/CodeGen/AMDGPU/unsupported-av-load.ll
index 942a384184924..fcf05dbe8e1f0 100644
--- a/llvm/test/CodeGen/AMDGPU/unsupported-av-load.ll
+++ b/llvm/test/CodeGen/AMDGPU/unsupported-av-load.ll
@@ -2,12 +2,12 @@
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx705 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx810 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx602 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx705 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx810 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx602 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx705 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx810 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
define <4 x i32> @av_load_b128(ptr addrspace(1) %addr) {
-; ERR: error: {{.*}}: in function av_load_b128 {{.*}}: llvm.amdgcn.av.load.b128 not supported on subtarget
+; ERR: error: {{.*}}: in function av_load_b128 {{.*}}: llvm.amdgcn.av.load.b128 requires target feature 'flat-global-insts'
entry:
%data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !0)
ret <4 x i32> %data
diff --git a/llvm/test/CodeGen/AMDGPU/unsupported-av-store.ll b/llvm/test/CodeGen/AMDGPU/unsupported-av-store.ll
index 4bd8945fea0b6..fef4840e815a2 100644
--- a/llvm/test/CodeGen/AMDGPU/unsupported-av-store.ll
+++ b/llvm/test/CodeGen/AMDGPU/unsupported-av-store.ll
@@ -2,12 +2,12 @@
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx705 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx810 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx602 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx705 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
-; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx810 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx602 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx705 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx810 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s
define void @av_store_b128(ptr addrspace(1) %addr, <4 x i32> %data) {
-; ERR: error: {{.*}}: in function av_store_b128 {{.*}}: llvm.amdgcn.av.store.b128 not supported on subtarget
+; ERR: error: {{.*}}: in function av_store_b128 {{.*}}: llvm.amdgcn.av.store.b128 requires target feature 'flat-global-insts'
entry:
call void @llvm.amdgcn.av.store.b128.p1(ptr addrspace(1) %addr, <4 x i32> %data, metadata !0)
ret void
More information about the llvm-branch-commits
mailing list