[llvm] r328938 - AMDGPU: Make getTgtMemIntrinsic table-driven for resource-based intrinsics
Nicolai Haehnle via llvm-commits
llvm-commits at lists.llvm.org
Sun Apr 1 10:09:07 PDT 2018
Author: nha
Date: Sun Apr 1 10:09:07 2018
New Revision: 328938
URL: http://llvm.org/viewvc/llvm-project?rev=328938&view=rev
Log:
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
Added:
llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td
Modified:
llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/trunk/lib/Target/AMDGPU/AMDGPU.td
llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.cpp
llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.h
llvm/trunk/lib/Target/AMDGPU/CMakeLists.txt
llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp
Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=328938&r1=328937&r2=328938&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Sun Apr 1 10:09:07 2018
@@ -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 : AMDGPULDSF32Int
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> : I
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> :
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 : In
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 : I
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.
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPU.td?rev=328938&r1=328937&r2=328938&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPU.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPU.td Sun Apr 1 10:09:07 2018
@@ -801,3 +801,4 @@ include "AMDGPURegisterInfo.td"
include "AMDGPURegisterBanks.td"
include "AMDGPUInstructions.td"
include "AMDGPUCallingConv.td"
+include "AMDGPUSearchableTables.td"
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.cpp?rev=328938&r1=328937&r2=328938&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.cpp Sun Apr 1 10:09:07 2018
@@ -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() {}
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.h?rev=328938&r1=328937&r2=328938&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.h (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.h Sun Apr 1 10:09:07 2018
@@ -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
Added: llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td?rev=328938&view=auto
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td (added)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td Sun Apr 1 10:09:07 2018
@@ -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)>;
+}
Modified: llvm/trunk/lib/Target/AMDGPU/CMakeLists.txt
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/CMakeLists.txt?rev=328938&r1=328937&r2=328938&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/CMakeLists.txt (original)
+++ llvm/trunk/lib/Target/AMDGPU/CMakeLists.txt Sun Apr 1 10:09:07 2018
@@ -13,6 +13,7 @@ tablegen(LLVM AMDGPUGenAsmMatcher.inc -g
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
Modified: llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp?rev=328938&r1=328937&r2=328938&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp Sun Apr 1 10:09:07 2018
@@ -570,6 +570,49 @@ bool SITargetLowering::getTgtMemIntrinsi
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::getTgtMemIntrinsi
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;
}
More information about the llvm-commits
mailing list