Index: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td =================================================================== --- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td +++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td @@ -17,6 +17,13 @@ class AMDGPUReadPreloadRegisterIntrinsicNamed : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, GCCBuiltin; +// Used to tag image and resource intrinsics with information used to generate +// mem operands. +class AMDGPURsrcIntrinsic { + int RsrcArg = rsrcarg; + bit IsImage = isimage; +} + let TargetPrefix = "r600" in { multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz { @@ -330,6 +337,8 @@ def int_amdgcn_ds_fmin : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmin">; def int_amdgcn_ds_fmax : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmax">; +defset list AMDGPUImageIntrinsics = { + class AMDGPUImageLoad : Intrinsic < [llvm_anyfloat_ty], // vdata(VGPR) [llvm_anyint_ty, // vaddr(VGPR) @@ -340,7 +349,8 @@ 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 @@ 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 @@ 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 @@ 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 @@ llvm_i1_ty, // r128(imm) llvm_i1_ty, // da(imm) llvm_i1_ty], // slc(imm) - [], "", [SDNPMemOperand]>; + [], "", [SDNPMemOperand]>, + AMDGPURsrcIntrinsic<3, 1>; + +} // defset AMDGPUImageIntrinsics + +defset list AMDGPUBufferIntrinsics = { class AMDGPUBufferLoad : Intrinsic < [llvm_anyfloat_ty], @@ -497,7 +515,8 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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. Index: llvm/trunk/lib/Target/AMDGPU/AMDGPU.td =================================================================== --- llvm/trunk/lib/Target/AMDGPU/AMDGPU.td +++ llvm/trunk/lib/Target/AMDGPU/AMDGPU.td @@ -801,3 +801,4 @@ include "AMDGPURegisterBanks.td" include "AMDGPUInstructions.td" include "AMDGPUCallingConv.td" +include "AMDGPUSearchableTables.td" Index: llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.h =================================================================== --- llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.h +++ llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.h @@ -53,6 +53,17 @@ 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 Index: llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.cpp =================================================================== --- llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.cpp +++ llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.cpp @@ -25,6 +25,13 @@ #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() {} Index: llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td =================================================================== --- llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td +++ llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td @@ -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 : SearchableTable { + let SearchableFields = ["Intr"]; + let EnumNameField = ?; + + Intrinsic Intr = !cast(intr); + bits<8> RsrcArg = intr.RsrcArg; + bit IsImage = intr.IsImage; +} + +foreach intr = !listconcat(AMDGPUBufferIntrinsics, + AMDGPUImageIntrinsics) in { + def : RsrcIntrinsic(intr)>; +} Index: llvm/trunk/lib/Target/AMDGPU/CMakeLists.txt =================================================================== --- llvm/trunk/lib/Target/AMDGPU/CMakeLists.txt +++ llvm/trunk/lib/Target/AMDGPU/CMakeLists.txt @@ -13,6 +13,7 @@ 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 Index: llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp =================================================================== --- llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp +++ llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp @@ -570,6 +570,49 @@ 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(); + + if (RsrcIntr->IsImage) { + Info.ptrVal = MFI->getImagePSV( + *MF.getSubtarget().getInstrInfo(), + CI.getArgOperand(RsrcIntr->RsrcArg)); + Info.align = 0; + } else { + Info.ptrVal = MFI->getBufferPSV( + *MF.getSubtarget().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 @@ 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(); - Info.opc = ISD::INTRINSIC_W_CHAIN; - Info.memVT = MVT::getVT(CI.getType()); - Info.ptrVal = MFI->getImagePSV( - *MF.getSubtarget().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(); - Info.opc = ISD::INTRINSIC_VOID; - Info.memVT = MVT::getVT(CI.getArgOperand(0)->getType()); - Info.ptrVal = MFI->getImagePSV( - *MF.getSubtarget().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(); - Info.opc = ISD::INTRINSIC_W_CHAIN; - Info.memVT = MVT::getVT(CI.getType()); - Info.ptrVal = MFI->getImagePSV( - *MF.getSubtarget().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(); - Info.opc = ISD::INTRINSIC_W_CHAIN; - Info.memVT = MVT::getVT(CI.getType()); - Info.ptrVal = MFI->getImagePSV( - *MF.getSubtarget().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(); - Info.opc = ISD::INTRINSIC_W_CHAIN; - Info.ptrVal = MFI->getBufferPSV( - *MF.getSubtarget().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(); - Info.opc = ISD::INTRINSIC_VOID; - Info.ptrVal = MFI->getBufferPSV( - *MF.getSubtarget().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(); - Info.opc = ISD::INTRINSIC_W_CHAIN; - Info.ptrVal = MFI->getBufferPSV( - *MF.getSubtarget().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(); - Info.opc = ISD::INTRINSIC_W_CHAIN; - Info.ptrVal = MFI->getBufferPSV( - *MF.getSubtarget().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; }