forked from OSchip/llvm-project
AMDGPU: Make getTgtMemIntrinsic table-driven for resource-based intrinsics
Summary: Avoids having to list all intrinsics manually. This is in preparation for the new dimension-aware image intrinsics, which I'd rather not have to list here by hand. Change-Id: If7ced04998397ef68c4cb8f7de66b5050fb767e5 Reviewers: arsenm, rampitec, b-sumner Subscribers: kzhuravl, wdng, mgorny, yaxunl, dstuttard, tpr, llvm-commits, t-tye Differential Revision: https://reviews.llvm.org/D44937 llvm-svn: 328938
This commit is contained in:
Nicolai Haehnle
parent
398c0b6701
commit
5d0d30304c
|
|
@ -17,6 +17,13 @@ class AMDGPUReadPreloadRegisterIntrinsic
|
|||
class AMDGPUReadPreloadRegisterIntrinsicNamed<string name>
|
||||
: Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, GCCBuiltin<name>;
|
||||
|
||||
// Used to tag image and resource intrinsics with information used to generate
|
||||
// mem operands.
|
||||
class AMDGPURsrcIntrinsic<int rsrcarg, bit isimage = 0> {
|
||||
int RsrcArg = rsrcarg;
|
||||
bit IsImage = isimage;
|
||||
}
|
||||
|
||||
let TargetPrefix = "r600" in {
|
||||
|
||||
multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz {
|
||||
|
|
@ -330,6 +337,8 @@ def int_amdgcn_ds_fadd : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fadd">;
|
|||
def int_amdgcn_ds_fmin : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmin">;
|
||||
def int_amdgcn_ds_fmax : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmax">;
|
||||
|
||||
defset list<AMDGPURsrcIntrinsic> AMDGPUImageIntrinsics = {
|
||||
|
||||
class AMDGPUImageLoad<bit NoMem = 0> : Intrinsic <
|
||||
[llvm_anyfloat_ty], // vdata(VGPR)
|
||||
[llvm_anyint_ty, // vaddr(VGPR)
|
||||
|
|
@ -340,7 +349,8 @@ class AMDGPUImageLoad<bit NoMem = 0> : Intrinsic <
|
|||
llvm_i1_ty, // lwe(imm)
|
||||
llvm_i1_ty], // da(imm)
|
||||
!if(NoMem, [IntrNoMem], [IntrReadMem]), "",
|
||||
!if(NoMem, [], [SDNPMemOperand])>;
|
||||
!if(NoMem, [], [SDNPMemOperand])>,
|
||||
AMDGPURsrcIntrinsic<1, 1>;
|
||||
|
||||
def int_amdgcn_image_load : AMDGPUImageLoad;
|
||||
def int_amdgcn_image_load_mip : AMDGPUImageLoad;
|
||||
|
|
@ -356,7 +366,8 @@ class AMDGPUImageStore : Intrinsic <
|
|||
llvm_i1_ty, // slc(imm)
|
||||
llvm_i1_ty, // lwe(imm)
|
||||
llvm_i1_ty], // da(imm)
|
||||
[IntrWriteMem], "", [SDNPMemOperand]>;
|
||||
[IntrWriteMem], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<2, 1>;
|
||||
|
||||
def int_amdgcn_image_store : AMDGPUImageStore;
|
||||
def int_amdgcn_image_store_mip : AMDGPUImageStore;
|
||||
|
|
@ -373,7 +384,8 @@ class AMDGPUImageSample<bit NoMem = 0> : Intrinsic <
|
|||
llvm_i1_ty, // lwe(imm)
|
||||
llvm_i1_ty], // da(imm)
|
||||
!if(NoMem, [IntrNoMem], [IntrReadMem]), "",
|
||||
!if(NoMem, [], [SDNPMemOperand])>;
|
||||
!if(NoMem, [], [SDNPMemOperand])>,
|
||||
AMDGPURsrcIntrinsic<1, 1>;
|
||||
|
||||
// Basic sample
|
||||
def int_amdgcn_image_sample : AMDGPUImageSample;
|
||||
|
|
@ -465,7 +477,8 @@ class AMDGPUImageAtomic : Intrinsic <
|
|||
llvm_i1_ty, // r128(imm)
|
||||
llvm_i1_ty, // da(imm)
|
||||
llvm_i1_ty], // slc(imm)
|
||||
[], "", [SDNPMemOperand]>;
|
||||
[], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<2, 1>;
|
||||
|
||||
def int_amdgcn_image_atomic_swap : AMDGPUImageAtomic;
|
||||
def int_amdgcn_image_atomic_add : AMDGPUImageAtomic;
|
||||
|
|
@ -488,7 +501,12 @@ def int_amdgcn_image_atomic_cmpswap : Intrinsic <
|
|||
llvm_i1_ty, // r128(imm)
|
||||
llvm_i1_ty, // da(imm)
|
||||
llvm_i1_ty], // slc(imm)
|
||||
[], "", [SDNPMemOperand]>;
|
||||
[], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<3, 1>;
|
||||
|
||||
} // defset AMDGPUImageIntrinsics
|
||||
|
||||
defset list<AMDGPURsrcIntrinsic> AMDGPUBufferIntrinsics = {
|
||||
|
||||
class AMDGPUBufferLoad : Intrinsic <
|
||||
[llvm_anyfloat_ty],
|
||||
|
|
@ -497,7 +515,8 @@ class AMDGPUBufferLoad : Intrinsic <
|
|||
llvm_i32_ty, // offset(SGPR/VGPR/imm)
|
||||
llvm_i1_ty, // glc(imm)
|
||||
llvm_i1_ty], // slc(imm)
|
||||
[IntrReadMem], "", [SDNPMemOperand]>;
|
||||
[IntrReadMem], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<0>;
|
||||
def int_amdgcn_buffer_load_format : AMDGPUBufferLoad;
|
||||
def int_amdgcn_buffer_load : AMDGPUBufferLoad;
|
||||
|
||||
|
|
@ -509,7 +528,8 @@ class AMDGPUBufferStore : Intrinsic <
|
|||
llvm_i32_ty, // offset(SGPR/VGPR/imm)
|
||||
llvm_i1_ty, // glc(imm)
|
||||
llvm_i1_ty], // slc(imm)
|
||||
[IntrWriteMem], "", [SDNPMemOperand]>;
|
||||
[IntrWriteMem], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<1>;
|
||||
def int_amdgcn_buffer_store_format : AMDGPUBufferStore;
|
||||
def int_amdgcn_buffer_store : AMDGPUBufferStore;
|
||||
|
||||
|
|
@ -524,7 +544,8 @@ def int_amdgcn_tbuffer_load : Intrinsic <
|
|||
llvm_i32_ty, // nfmt(imm)
|
||||
llvm_i1_ty, // glc(imm)
|
||||
llvm_i1_ty], // slc(imm)
|
||||
[IntrReadMem], "", [SDNPMemOperand]>;
|
||||
[IntrReadMem], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<0>;
|
||||
|
||||
def int_amdgcn_tbuffer_store : Intrinsic <
|
||||
[],
|
||||
|
|
@ -538,7 +559,8 @@ def int_amdgcn_tbuffer_store : Intrinsic <
|
|||
llvm_i32_ty, // nfmt(imm)
|
||||
llvm_i1_ty, // glc(imm)
|
||||
llvm_i1_ty], // slc(imm)
|
||||
[IntrWriteMem], "", [SDNPMemOperand]>;
|
||||
[IntrWriteMem], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<1>;
|
||||
|
||||
class AMDGPUBufferAtomic : Intrinsic <
|
||||
[llvm_i32_ty],
|
||||
|
|
@ -547,7 +569,8 @@ class AMDGPUBufferAtomic : Intrinsic <
|
|||
llvm_i32_ty, // vindex(VGPR)
|
||||
llvm_i32_ty, // offset(SGPR/VGPR/imm)
|
||||
llvm_i1_ty], // slc(imm)
|
||||
[], "", [SDNPMemOperand]>;
|
||||
[], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<1, 0>;
|
||||
def int_amdgcn_buffer_atomic_swap : AMDGPUBufferAtomic;
|
||||
def int_amdgcn_buffer_atomic_add : AMDGPUBufferAtomic;
|
||||
def int_amdgcn_buffer_atomic_sub : AMDGPUBufferAtomic;
|
||||
|
|
@ -566,7 +589,10 @@ def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
|
|||
llvm_i32_ty, // vindex(VGPR)
|
||||
llvm_i32_ty, // offset(SGPR/VGPR/imm)
|
||||
llvm_i1_ty], // slc(imm)
|
||||
[], "", [SDNPMemOperand]>;
|
||||
[], "", [SDNPMemOperand]>,
|
||||
AMDGPURsrcIntrinsic<2, 0>;
|
||||
|
||||
} // defset AMDGPUBufferIntrinsics
|
||||
|
||||
// Uses that do not set the done bit should set IntrWriteMem on the
|
||||
// call site.
|
||||
|
|
|
|||
|
|
@ -801,3 +801,4 @@ include "AMDGPURegisterInfo.td"
|
|||
include "AMDGPURegisterBanks.td"
|
||||
include "AMDGPUInstructions.td"
|
||||
include "AMDGPUCallingConv.td"
|
||||
include "AMDGPUSearchableTables.td"
|
||||
|
|
|
|||
|
|
@ -25,6 +25,13 @@ using namespace llvm;
|
|||
#define GET_INSTRINFO_CTOR_DTOR
|
||||
#include "AMDGPUGenInstrInfo.inc"
|
||||
|
||||
namespace llvm {
|
||||
namespace AMDGPU {
|
||||
#define GET_RSRCINTRINSIC_IMPL
|
||||
#include "AMDGPUGenSearchableTables.inc"
|
||||
}
|
||||
}
|
||||
|
||||
// Pin the vtable to this file.
|
||||
void AMDGPUInstrInfo::anchor() {}
|
||||
|
||||
|
|
|
|||
|
|
@ -53,6 +53,17 @@ public:
|
|||
|
||||
static bool isUniformMMO(const MachineMemOperand *MMO);
|
||||
};
|
||||
|
||||
namespace AMDGPU {
|
||||
|
||||
struct RsrcIntrinsic {
|
||||
unsigned Intr;
|
||||
uint8_t RsrcArg;
|
||||
bool IsImage;
|
||||
};
|
||||
const RsrcIntrinsic *lookupRsrcIntrinsicByIntr(unsigned Intr);
|
||||
|
||||
} // end AMDGPU namespace
|
||||
} // End llvm namespace
|
||||
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -0,0 +1,28 @@
|
|||
//===-- AMDGPUSearchableTables.td - ------------------------*- tablegen -*-===//
|
||||
//
|
||||
// The LLVM Compiler Infrastructure
|
||||
//
|
||||
// This file is distributed under the University of Illinois Open Source
|
||||
// License. See LICENSE.TXT for details.
|
||||
//
|
||||
//===----------------------------------------------------------------------===//
|
||||
|
||||
include "llvm/TableGen/SearchableTable.td"
|
||||
|
||||
//===----------------------------------------------------------------------===//
|
||||
// Resource intrinsics table.
|
||||
//===----------------------------------------------------------------------===//
|
||||
|
||||
class RsrcIntrinsic<AMDGPURsrcIntrinsic intr> : SearchableTable {
|
||||
let SearchableFields = ["Intr"];
|
||||
let EnumNameField = ?;
|
||||
|
||||
Intrinsic Intr = !cast<Intrinsic>(intr);
|
||||
bits<8> RsrcArg = intr.RsrcArg;
|
||||
bit IsImage = intr.IsImage;
|
||||
}
|
||||
|
||||
foreach intr = !listconcat(AMDGPUBufferIntrinsics,
|
||||
AMDGPUImageIntrinsics) in {
|
||||
def : RsrcIntrinsic<!cast<AMDGPURsrcIntrinsic>(intr)>;
|
||||
}
|
||||
|
|
@ -13,6 +13,7 @@ tablegen(LLVM AMDGPUGenAsmMatcher.inc -gen-asm-matcher)
|
|||
tablegen(LLVM AMDGPUGenDisassemblerTables.inc -gen-disassembler)
|
||||
tablegen(LLVM AMDGPUGenMCPseudoLowering.inc -gen-pseudo-lowering)
|
||||
tablegen(LLVM AMDGPUGenRegisterBank.inc -gen-register-bank)
|
||||
tablegen(LLVM AMDGPUGenSearchableTables.inc -gen-searchable-tables)
|
||||
add_public_tablegen_target(AMDGPUCommonTableGen)
|
||||
|
||||
add_llvm_target(AMDGPUCodeGen
|
||||
|
|
|
|||
|
|
@ -570,6 +570,49 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info,
|
|||
const CallInst &CI,
|
||||
MachineFunction &MF,
|
||||
unsigned IntrID) const {
|
||||
if (const AMDGPU::RsrcIntrinsic *RsrcIntr =
|
||||
AMDGPU::lookupRsrcIntrinsicByIntr(IntrID)) {
|
||||
AttributeList Attr = Intrinsic::getAttributes(CI.getContext(),
|
||||
(Intrinsic::ID)IntrID);
|
||||
if (Attr.hasFnAttribute(Attribute::ReadNone))
|
||||
return false;
|
||||
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
|
||||
if (RsrcIntr->IsImage) {
|
||||
Info.ptrVal = MFI->getImagePSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(RsrcIntr->RsrcArg));
|
||||
Info.align = 0;
|
||||
} else {
|
||||
Info.ptrVal = MFI->getBufferPSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(RsrcIntr->RsrcArg));
|
||||
}
|
||||
|
||||
Info.flags = MachineMemOperand::MODereferenceable;
|
||||
if (Attr.hasFnAttribute(Attribute::ReadOnly)) {
|
||||
Info.opc = ISD::INTRINSIC_W_CHAIN;
|
||||
Info.memVT = MVT::getVT(CI.getType());
|
||||
Info.flags |= MachineMemOperand::MOLoad;
|
||||
} else if (Attr.hasFnAttribute(Attribute::WriteOnly)) {
|
||||
Info.opc = ISD::INTRINSIC_VOID;
|
||||
Info.memVT = MVT::getVT(CI.getArgOperand(0)->getType());
|
||||
Info.flags |= MachineMemOperand::MOStore;
|
||||
} else {
|
||||
// Atomic
|
||||
Info.opc = ISD::INTRINSIC_W_CHAIN;
|
||||
Info.memVT = MVT::getVT(CI.getType());
|
||||
Info.flags = MachineMemOperand::MOLoad |
|
||||
MachineMemOperand::MOStore |
|
||||
MachineMemOperand::MODereferenceable;
|
||||
|
||||
// XXX - Should this be volatile without known ordering?
|
||||
Info.flags |= MachineMemOperand::MOVolatile;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
switch (IntrID) {
|
||||
case Intrinsic::amdgcn_atomic_inc:
|
||||
case Intrinsic::amdgcn_atomic_dec:
|
||||
|
|
@ -589,220 +632,6 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info,
|
|||
return true;
|
||||
}
|
||||
|
||||
// Image load.
|
||||
case Intrinsic::amdgcn_image_load:
|
||||
case Intrinsic::amdgcn_image_load_mip:
|
||||
|
||||
// Sample.
|
||||
case Intrinsic::amdgcn_image_sample:
|
||||
case Intrinsic::amdgcn_image_sample_cl:
|
||||
case Intrinsic::amdgcn_image_sample_d:
|
||||
case Intrinsic::amdgcn_image_sample_d_cl:
|
||||
case Intrinsic::amdgcn_image_sample_l:
|
||||
case Intrinsic::amdgcn_image_sample_b:
|
||||
case Intrinsic::amdgcn_image_sample_b_cl:
|
||||
case Intrinsic::amdgcn_image_sample_lz:
|
||||
case Intrinsic::amdgcn_image_sample_cd:
|
||||
case Intrinsic::amdgcn_image_sample_cd_cl:
|
||||
|
||||
// Sample with comparison.
|
||||
case Intrinsic::amdgcn_image_sample_c:
|
||||
case Intrinsic::amdgcn_image_sample_c_cl:
|
||||
case Intrinsic::amdgcn_image_sample_c_d:
|
||||
case Intrinsic::amdgcn_image_sample_c_d_cl:
|
||||
case Intrinsic::amdgcn_image_sample_c_l:
|
||||
case Intrinsic::amdgcn_image_sample_c_b:
|
||||
case Intrinsic::amdgcn_image_sample_c_b_cl:
|
||||
case Intrinsic::amdgcn_image_sample_c_lz:
|
||||
case Intrinsic::amdgcn_image_sample_c_cd:
|
||||
case Intrinsic::amdgcn_image_sample_c_cd_cl:
|
||||
|
||||
// Sample with offsets.
|
||||
case Intrinsic::amdgcn_image_sample_o:
|
||||
case Intrinsic::amdgcn_image_sample_cl_o:
|
||||
case Intrinsic::amdgcn_image_sample_d_o:
|
||||
case Intrinsic::amdgcn_image_sample_d_cl_o:
|
||||
case Intrinsic::amdgcn_image_sample_l_o:
|
||||
case Intrinsic::amdgcn_image_sample_b_o:
|
||||
case Intrinsic::amdgcn_image_sample_b_cl_o:
|
||||
case Intrinsic::amdgcn_image_sample_lz_o:
|
||||
case Intrinsic::amdgcn_image_sample_cd_o:
|
||||
case Intrinsic::amdgcn_image_sample_cd_cl_o:
|
||||
|
||||
// Sample with comparison and offsets.
|
||||
case Intrinsic::amdgcn_image_sample_c_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_cl_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_d_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_d_cl_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_l_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_b_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_b_cl_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_lz_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_cd_o:
|
||||
case Intrinsic::amdgcn_image_sample_c_cd_cl_o:
|
||||
|
||||
// Basic gather4
|
||||
case Intrinsic::amdgcn_image_gather4:
|
||||
case Intrinsic::amdgcn_image_gather4_cl:
|
||||
case Intrinsic::amdgcn_image_gather4_l:
|
||||
case Intrinsic::amdgcn_image_gather4_b:
|
||||
case Intrinsic::amdgcn_image_gather4_b_cl:
|
||||
case Intrinsic::amdgcn_image_gather4_lz:
|
||||
|
||||
// Gather4 with comparison
|
||||
case Intrinsic::amdgcn_image_gather4_c:
|
||||
case Intrinsic::amdgcn_image_gather4_c_cl:
|
||||
case Intrinsic::amdgcn_image_gather4_c_l:
|
||||
case Intrinsic::amdgcn_image_gather4_c_b:
|
||||
case Intrinsic::amdgcn_image_gather4_c_b_cl:
|
||||
case Intrinsic::amdgcn_image_gather4_c_lz:
|
||||
|
||||
// Gather4 with offsets
|
||||
case Intrinsic::amdgcn_image_gather4_o:
|
||||
case Intrinsic::amdgcn_image_gather4_cl_o:
|
||||
case Intrinsic::amdgcn_image_gather4_l_o:
|
||||
case Intrinsic::amdgcn_image_gather4_b_o:
|
||||
case Intrinsic::amdgcn_image_gather4_b_cl_o:
|
||||
case Intrinsic::amdgcn_image_gather4_lz_o:
|
||||
|
||||
// Gather4 with comparison and offsets
|
||||
case Intrinsic::amdgcn_image_gather4_c_o:
|
||||
case Intrinsic::amdgcn_image_gather4_c_cl_o:
|
||||
case Intrinsic::amdgcn_image_gather4_c_l_o:
|
||||
case Intrinsic::amdgcn_image_gather4_c_b_o:
|
||||
case Intrinsic::amdgcn_image_gather4_c_b_cl_o:
|
||||
case Intrinsic::amdgcn_image_gather4_c_lz_o: {
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
Info.opc = ISD::INTRINSIC_W_CHAIN;
|
||||
Info.memVT = MVT::getVT(CI.getType());
|
||||
Info.ptrVal = MFI->getImagePSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(1));
|
||||
Info.align = 0;
|
||||
Info.flags = MachineMemOperand::MOLoad |
|
||||
MachineMemOperand::MODereferenceable;
|
||||
return true;
|
||||
}
|
||||
case Intrinsic::amdgcn_image_store:
|
||||
case Intrinsic::amdgcn_image_store_mip: {
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
Info.opc = ISD::INTRINSIC_VOID;
|
||||
Info.memVT = MVT::getVT(CI.getArgOperand(0)->getType());
|
||||
Info.ptrVal = MFI->getImagePSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(2));
|
||||
Info.flags = MachineMemOperand::MOStore |
|
||||
MachineMemOperand::MODereferenceable;
|
||||
Info.align = 0;
|
||||
return true;
|
||||
}
|
||||
case Intrinsic::amdgcn_image_atomic_swap:
|
||||
case Intrinsic::amdgcn_image_atomic_add:
|
||||
case Intrinsic::amdgcn_image_atomic_sub:
|
||||
case Intrinsic::amdgcn_image_atomic_smin:
|
||||
case Intrinsic::amdgcn_image_atomic_umin:
|
||||
case Intrinsic::amdgcn_image_atomic_smax:
|
||||
case Intrinsic::amdgcn_image_atomic_umax:
|
||||
case Intrinsic::amdgcn_image_atomic_and:
|
||||
case Intrinsic::amdgcn_image_atomic_or:
|
||||
case Intrinsic::amdgcn_image_atomic_xor:
|
||||
case Intrinsic::amdgcn_image_atomic_inc:
|
||||
case Intrinsic::amdgcn_image_atomic_dec: {
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
Info.opc = ISD::INTRINSIC_W_CHAIN;
|
||||
Info.memVT = MVT::getVT(CI.getType());
|
||||
Info.ptrVal = MFI->getImagePSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(2));
|
||||
|
||||
Info.flags = MachineMemOperand::MOLoad |
|
||||
MachineMemOperand::MOStore |
|
||||
MachineMemOperand::MODereferenceable;
|
||||
|
||||
// XXX - Should this be volatile without known ordering?
|
||||
Info.flags |= MachineMemOperand::MOVolatile;
|
||||
return true;
|
||||
}
|
||||
case Intrinsic::amdgcn_image_atomic_cmpswap: {
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
Info.opc = ISD::INTRINSIC_W_CHAIN;
|
||||
Info.memVT = MVT::getVT(CI.getType());
|
||||
Info.ptrVal = MFI->getImagePSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(3));
|
||||
|
||||
Info.flags = MachineMemOperand::MOLoad |
|
||||
MachineMemOperand::MOStore |
|
||||
MachineMemOperand::MODereferenceable;
|
||||
|
||||
// XXX - Should this be volatile without known ordering?
|
||||
Info.flags |= MachineMemOperand::MOVolatile;
|
||||
return true;
|
||||
}
|
||||
case Intrinsic::amdgcn_tbuffer_load:
|
||||
case Intrinsic::amdgcn_buffer_load:
|
||||
case Intrinsic::amdgcn_buffer_load_format: {
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
Info.opc = ISD::INTRINSIC_W_CHAIN;
|
||||
Info.ptrVal = MFI->getBufferPSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(0));
|
||||
Info.memVT = MVT::getVT(CI.getType());
|
||||
Info.flags = MachineMemOperand::MOLoad |
|
||||
MachineMemOperand::MODereferenceable;
|
||||
|
||||
// There is a constant offset component, but there are additional register
|
||||
// offsets which could break AA if we set the offset to anything non-0.
|
||||
return true;
|
||||
}
|
||||
case Intrinsic::amdgcn_tbuffer_store:
|
||||
case Intrinsic::amdgcn_buffer_store:
|
||||
case Intrinsic::amdgcn_buffer_store_format: {
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
Info.opc = ISD::INTRINSIC_VOID;
|
||||
Info.ptrVal = MFI->getBufferPSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(1));
|
||||
Info.memVT = MVT::getVT(CI.getArgOperand(0)->getType());
|
||||
Info.flags = MachineMemOperand::MOStore |
|
||||
MachineMemOperand::MODereferenceable;
|
||||
return true;
|
||||
}
|
||||
case Intrinsic::amdgcn_buffer_atomic_swap:
|
||||
case Intrinsic::amdgcn_buffer_atomic_add:
|
||||
case Intrinsic::amdgcn_buffer_atomic_sub:
|
||||
case Intrinsic::amdgcn_buffer_atomic_smin:
|
||||
case Intrinsic::amdgcn_buffer_atomic_umin:
|
||||
case Intrinsic::amdgcn_buffer_atomic_smax:
|
||||
case Intrinsic::amdgcn_buffer_atomic_umax:
|
||||
case Intrinsic::amdgcn_buffer_atomic_and:
|
||||
case Intrinsic::amdgcn_buffer_atomic_or:
|
||||
case Intrinsic::amdgcn_buffer_atomic_xor: {
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
Info.opc = ISD::INTRINSIC_W_CHAIN;
|
||||
Info.ptrVal = MFI->getBufferPSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(1));
|
||||
Info.memVT = MVT::getVT(CI.getType());
|
||||
Info.flags = MachineMemOperand::MOLoad |
|
||||
MachineMemOperand::MOStore |
|
||||
MachineMemOperand::MODereferenceable |
|
||||
MachineMemOperand::MOVolatile;
|
||||
return true;
|
||||
}
|
||||
case Intrinsic::amdgcn_buffer_atomic_cmpswap: {
|
||||
SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
|
||||
Info.opc = ISD::INTRINSIC_W_CHAIN;
|
||||
Info.ptrVal = MFI->getBufferPSV(
|
||||
*MF.getSubtarget<SISubtarget>().getInstrInfo(),
|
||||
CI.getArgOperand(2));
|
||||
Info.memVT = MVT::getVT(CI.getType());
|
||||
Info.flags = MachineMemOperand::MOLoad |
|
||||
MachineMemOperand::MOStore |
|
||||
MachineMemOperand::MODereferenceable |
|
||||
MachineMemOperand::MOVolatile;
|
||||
return true;
|
||||
}
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue