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