Skip to content

Commit

Permalink
AMDGPU: Remove old-style image intrinsics
Browse files Browse the repository at this point in the history 
Summary:
This also removes the need for atomic pseudo instructions, since
we select the correct encoding directly in SITargetLowering::lowerImage
for dimension-aware image intrinsics.

Mesa uses dimension-aware image intrinsics since
commit a9a7993441.

Change-Id: I7473d20009476a4ed6d919cae4e6dca9ff42e77a

Reviewers: arsenm, rampitec, mareko, tpr, b-sumner

Subscribers: kzhuravl, wdng, yaxunl, dstuttard, t-tye, llvm-commits

Differential Revision: https://reviews.llvm.org/D48167

llvm-svn: 335231
  • Loading branch information
@nhaehnle
nhaehnle committed Jun 21, 2018
1 parent b29ee70 commit db6911a
Show file tree
Hide file tree
Showing 19 changed files with 2 additions and 3,319 deletions.
169 changes: 0 additions & 169 deletions llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View file
Expand Up @@ -390,175 +390,6 @@ def int_amdgcn_ds_fadd : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_faddf">;
def int_amdgcn_ds_fmin : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fminf">;
def int_amdgcn_ds_fmax : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmaxf">;

defset list<AMDGPURsrcIntrinsic> AMDGPUImageIntrinsics = {

class AMDGPUImageLoad<bit NoMem = 0> : Intrinsic <
[llvm_anyfloat_ty], // vdata(VGPR)
[llvm_anyint_ty, // vaddr(VGPR)
llvm_anyint_ty, // rsrc(SGPR)
llvm_i32_ty, // dmask(imm)
llvm_i1_ty, // glc(imm)
llvm_i1_ty, // slc(imm)
llvm_i1_ty, // lwe(imm)
llvm_i1_ty], // da(imm)
!if(NoMem, [IntrNoMem], [IntrReadMem]), "",
!if(NoMem, [], [SDNPMemOperand])>,
AMDGPURsrcIntrinsic<1, 1>;

def int_amdgcn_image_load : AMDGPUImageLoad;
def int_amdgcn_image_load_mip : AMDGPUImageLoad;
def int_amdgcn_image_getresinfo : AMDGPUImageLoad<1>;

class AMDGPUImageStore : Intrinsic <
[],
[llvm_anyfloat_ty, // vdata(VGPR)
llvm_anyint_ty, // vaddr(VGPR)
llvm_anyint_ty, // rsrc(SGPR)
llvm_i32_ty, // dmask(imm)
llvm_i1_ty, // glc(imm)
llvm_i1_ty, // slc(imm)
llvm_i1_ty, // lwe(imm)
llvm_i1_ty], // da(imm)
[IntrWriteMem], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<2, 1>;

def int_amdgcn_image_store : AMDGPUImageStore;
def int_amdgcn_image_store_mip : AMDGPUImageStore;

class AMDGPUImageSample<bit NoMem = 0> : Intrinsic <
[llvm_anyfloat_ty], // vdata(VGPR)
[llvm_anyfloat_ty, // vaddr(VGPR)
llvm_anyint_ty, // rsrc(SGPR)
llvm_v4i32_ty, // sampler(SGPR)
llvm_i32_ty, // dmask(imm)
llvm_i1_ty, // unorm(imm)
llvm_i1_ty, // glc(imm)
llvm_i1_ty, // slc(imm)
llvm_i1_ty, // lwe(imm)
llvm_i1_ty], // da(imm)
!if(NoMem, [IntrNoMem], [IntrReadMem]), "",
!if(NoMem, [], [SDNPMemOperand])>,
AMDGPURsrcIntrinsic<1, 1>;

// Basic sample
def int_amdgcn_image_sample : AMDGPUImageSample;
def int_amdgcn_image_sample_cl : AMDGPUImageSample;
def int_amdgcn_image_sample_d : AMDGPUImageSample;
def int_amdgcn_image_sample_d_cl : AMDGPUImageSample;
def int_amdgcn_image_sample_l : AMDGPUImageSample;
def int_amdgcn_image_sample_b : AMDGPUImageSample;
def int_amdgcn_image_sample_b_cl : AMDGPUImageSample;
def int_amdgcn_image_sample_lz : AMDGPUImageSample;
def int_amdgcn_image_sample_cd : AMDGPUImageSample;
def int_amdgcn_image_sample_cd_cl : AMDGPUImageSample;

// Sample with comparison
def int_amdgcn_image_sample_c : AMDGPUImageSample;
def int_amdgcn_image_sample_c_cl : AMDGPUImageSample;
def int_amdgcn_image_sample_c_d : AMDGPUImageSample;
def int_amdgcn_image_sample_c_d_cl : AMDGPUImageSample;
def int_amdgcn_image_sample_c_l : AMDGPUImageSample;
def int_amdgcn_image_sample_c_b : AMDGPUImageSample;
def int_amdgcn_image_sample_c_b_cl : AMDGPUImageSample;
def int_amdgcn_image_sample_c_lz : AMDGPUImageSample;
def int_amdgcn_image_sample_c_cd : AMDGPUImageSample;
def int_amdgcn_image_sample_c_cd_cl : AMDGPUImageSample;

// Sample with offsets
def int_amdgcn_image_sample_o : AMDGPUImageSample;
def int_amdgcn_image_sample_cl_o : AMDGPUImageSample;
def int_amdgcn_image_sample_d_o : AMDGPUImageSample;
def int_amdgcn_image_sample_d_cl_o : AMDGPUImageSample;
def int_amdgcn_image_sample_l_o : AMDGPUImageSample;
def int_amdgcn_image_sample_b_o : AMDGPUImageSample;
def int_amdgcn_image_sample_b_cl_o : AMDGPUImageSample;
def int_amdgcn_image_sample_lz_o : AMDGPUImageSample;
def int_amdgcn_image_sample_cd_o : AMDGPUImageSample;
def int_amdgcn_image_sample_cd_cl_o : AMDGPUImageSample;

// Sample with comparison and offsets
def int_amdgcn_image_sample_c_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_cl_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_d_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_d_cl_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_l_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_b_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_b_cl_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_lz_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_cd_o : AMDGPUImageSample;
def int_amdgcn_image_sample_c_cd_cl_o : AMDGPUImageSample;

// Basic gather4
def int_amdgcn_image_gather4 : AMDGPUImageSample;
def int_amdgcn_image_gather4_cl : AMDGPUImageSample;
def int_amdgcn_image_gather4_l : AMDGPUImageSample;
def int_amdgcn_image_gather4_b : AMDGPUImageSample;
def int_amdgcn_image_gather4_b_cl : AMDGPUImageSample;
def int_amdgcn_image_gather4_lz : AMDGPUImageSample;

// Gather4 with comparison
def int_amdgcn_image_gather4_c : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_cl : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_l : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_b : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_b_cl : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_lz : AMDGPUImageSample;

// Gather4 with offsets
def int_amdgcn_image_gather4_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_cl_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_l_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_b_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_b_cl_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_lz_o : AMDGPUImageSample;

// Gather4 with comparison and offsets
def int_amdgcn_image_gather4_c_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_cl_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_l_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_b_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_b_cl_o : AMDGPUImageSample;
def int_amdgcn_image_gather4_c_lz_o : AMDGPUImageSample;

def int_amdgcn_image_getlod : AMDGPUImageSample<1>;

class AMDGPUImageAtomic : Intrinsic <
[llvm_i32_ty],
[llvm_i32_ty, // vdata(VGPR)
llvm_anyint_ty, // vaddr(VGPR)
llvm_v8i32_ty, // rsrc(SGPR)
llvm_i1_ty, // r128(imm)
llvm_i1_ty, // da(imm)
llvm_i1_ty], // slc(imm)
[], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<2, 1>;

def int_amdgcn_image_atomic_swap : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_add : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_sub : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_smin : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_umin : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_smax : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_umax : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_and : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_or : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_xor : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_inc : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_dec : AMDGPUImageAtomic;
def int_amdgcn_image_atomic_cmpswap : Intrinsic <
[llvm_i32_ty],
[llvm_i32_ty, // src(VGPR)
llvm_i32_ty, // cmp(VGPR)
llvm_anyint_ty, // vaddr(VGPR)
llvm_v8i32_ty, // rsrc(SGPR)
llvm_i1_ty, // r128(imm)
llvm_i1_ty, // da(imm)
llvm_i1_ty], // slc(imm)
[], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<3, 1>;

} // defset AMDGPUImageIntrinsics

} // TargetPrefix = "amdgcn"

// New-style image intrinsics
Expand Down
76 changes: 0 additions & 76 deletions llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
View file
Expand Up @@ -4130,82 +4130,6 @@ const char* AMDGPUTargetLowering::getTargetNodeName(unsigned Opcode) const {
NODE_NAME_CASE(BUFFER_ATOMIC_OR)
NODE_NAME_CASE(BUFFER_ATOMIC_XOR)
NODE_NAME_CASE(BUFFER_ATOMIC_CMPSWAP)
NODE_NAME_CASE(IMAGE_LOAD)
NODE_NAME_CASE(IMAGE_LOAD_MIP)
NODE_NAME_CASE(IMAGE_STORE)
NODE_NAME_CASE(IMAGE_STORE_MIP)
// Basic sample.
NODE_NAME_CASE(IMAGE_SAMPLE)
NODE_NAME_CASE(IMAGE_SAMPLE_CL)
NODE_NAME_CASE(IMAGE_SAMPLE_D)
NODE_NAME_CASE(IMAGE_SAMPLE_D_CL)
NODE_NAME_CASE(IMAGE_SAMPLE_L)
NODE_NAME_CASE(IMAGE_SAMPLE_B)
NODE_NAME_CASE(IMAGE_SAMPLE_B_CL)
NODE_NAME_CASE(IMAGE_SAMPLE_LZ)
NODE_NAME_CASE(IMAGE_SAMPLE_CD)
NODE_NAME_CASE(IMAGE_SAMPLE_CD_CL)
// Sample with comparison.
NODE_NAME_CASE(IMAGE_SAMPLE_C)
NODE_NAME_CASE(IMAGE_SAMPLE_C_CL)
NODE_NAME_CASE(IMAGE_SAMPLE_C_D)
NODE_NAME_CASE(IMAGE_SAMPLE_C_D_CL)
NODE_NAME_CASE(IMAGE_SAMPLE_C_L)
NODE_NAME_CASE(IMAGE_SAMPLE_C_B)
NODE_NAME_CASE(IMAGE_SAMPLE_C_B_CL)
NODE_NAME_CASE(IMAGE_SAMPLE_C_LZ)
NODE_NAME_CASE(IMAGE_SAMPLE_C_CD)
NODE_NAME_CASE(IMAGE_SAMPLE_C_CD_CL)
// Sample with offsets.
NODE_NAME_CASE(IMAGE_SAMPLE_O)
NODE_NAME_CASE(IMAGE_SAMPLE_CL_O)
NODE_NAME_CASE(IMAGE_SAMPLE_D_O)
NODE_NAME_CASE(IMAGE_SAMPLE_D_CL_O)
NODE_NAME_CASE(IMAGE_SAMPLE_L_O)
NODE_NAME_CASE(IMAGE_SAMPLE_B_O)
NODE_NAME_CASE(IMAGE_SAMPLE_B_CL_O)
NODE_NAME_CASE(IMAGE_SAMPLE_LZ_O)
NODE_NAME_CASE(IMAGE_SAMPLE_CD_O)
NODE_NAME_CASE(IMAGE_SAMPLE_CD_CL_O)
// Sample with comparison and offsets.
NODE_NAME_CASE(IMAGE_SAMPLE_C_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_CL_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_D_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_D_CL_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_L_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_B_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_B_CL_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_LZ_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_CD_O)
NODE_NAME_CASE(IMAGE_SAMPLE_C_CD_CL_O)
// Basic gather4.
NODE_NAME_CASE(IMAGE_GATHER4)
NODE_NAME_CASE(IMAGE_GATHER4_CL)
NODE_NAME_CASE(IMAGE_GATHER4_L)
NODE_NAME_CASE(IMAGE_GATHER4_B)
NODE_NAME_CASE(IMAGE_GATHER4_B_CL)
NODE_NAME_CASE(IMAGE_GATHER4_LZ)
// Gather4 with comparison.
NODE_NAME_CASE(IMAGE_GATHER4_C)
NODE_NAME_CASE(IMAGE_GATHER4_C_CL)
NODE_NAME_CASE(IMAGE_GATHER4_C_L)
NODE_NAME_CASE(IMAGE_GATHER4_C_B)
NODE_NAME_CASE(IMAGE_GATHER4_C_B_CL)
NODE_NAME_CASE(IMAGE_GATHER4_C_LZ)
// Gather4 with offsets.
NODE_NAME_CASE(IMAGE_GATHER4_O)
NODE_NAME_CASE(IMAGE_GATHER4_CL_O)
NODE_NAME_CASE(IMAGE_GATHER4_L_O)
NODE_NAME_CASE(IMAGE_GATHER4_B_O)
NODE_NAME_CASE(IMAGE_GATHER4_B_CL_O)
NODE_NAME_CASE(IMAGE_GATHER4_LZ_O)
// Gather4 with comparison and offsets.
NODE_NAME_CASE(IMAGE_GATHER4_C_O)
NODE_NAME_CASE(IMAGE_GATHER4_C_CL_O)
NODE_NAME_CASE(IMAGE_GATHER4_C_L_O)
NODE_NAME_CASE(IMAGE_GATHER4_C_B_O)
NODE_NAME_CASE(IMAGE_GATHER4_C_B_CL_O)
NODE_NAME_CASE(IMAGE_GATHER4_C_LZ_O)

case AMDGPUISD::LAST_AMDGPU_ISD_NUMBER: break;
}
Expand Down
84 changes: 0 additions & 84 deletions llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h
View file
Expand Up @@ -486,90 +486,6 @@ enum NodeType : unsigned {
BUFFER_ATOMIC_OR,
BUFFER_ATOMIC_XOR,
BUFFER_ATOMIC_CMPSWAP,
IMAGE_LOAD,
IMAGE_LOAD_MIP,
IMAGE_STORE,
IMAGE_STORE_MIP,

// Basic sample.
IMAGE_SAMPLE,
IMAGE_SAMPLE_CL,
IMAGE_SAMPLE_D,
IMAGE_SAMPLE_D_CL,
IMAGE_SAMPLE_L,
IMAGE_SAMPLE_B,
IMAGE_SAMPLE_B_CL,
IMAGE_SAMPLE_LZ,
IMAGE_SAMPLE_CD,
IMAGE_SAMPLE_CD_CL,

// Sample with comparison.
IMAGE_SAMPLE_C,
IMAGE_SAMPLE_C_CL,
IMAGE_SAMPLE_C_D,
IMAGE_SAMPLE_C_D_CL,
IMAGE_SAMPLE_C_L,
IMAGE_SAMPLE_C_B,
IMAGE_SAMPLE_C_B_CL,
IMAGE_SAMPLE_C_LZ,
IMAGE_SAMPLE_C_CD,
IMAGE_SAMPLE_C_CD_CL,

// Sample with offsets.
IMAGE_SAMPLE_O,
IMAGE_SAMPLE_CL_O,
IMAGE_SAMPLE_D_O,
IMAGE_SAMPLE_D_CL_O,
IMAGE_SAMPLE_L_O,
IMAGE_SAMPLE_B_O,
IMAGE_SAMPLE_B_CL_O,
IMAGE_SAMPLE_LZ_O,
IMAGE_SAMPLE_CD_O,
IMAGE_SAMPLE_CD_CL_O,

// Sample with comparison and offsets.
IMAGE_SAMPLE_C_O,
IMAGE_SAMPLE_C_CL_O,
IMAGE_SAMPLE_C_D_O,
IMAGE_SAMPLE_C_D_CL_O,
IMAGE_SAMPLE_C_L_O,
IMAGE_SAMPLE_C_B_O,
IMAGE_SAMPLE_C_B_CL_O,
IMAGE_SAMPLE_C_LZ_O,
IMAGE_SAMPLE_C_CD_O,
IMAGE_SAMPLE_C_CD_CL_O,

// Basic gather4.
IMAGE_GATHER4,
IMAGE_GATHER4_CL,
IMAGE_GATHER4_L,
IMAGE_GATHER4_B,
IMAGE_GATHER4_B_CL,
IMAGE_GATHER4_LZ,

// Gather4 with comparison.
IMAGE_GATHER4_C,
IMAGE_GATHER4_C_CL,
IMAGE_GATHER4_C_L,
IMAGE_GATHER4_C_B,
IMAGE_GATHER4_C_B_CL,
IMAGE_GATHER4_C_LZ,

// Gather4 with offsets.
IMAGE_GATHER4_O,
IMAGE_GATHER4_CL_O,
IMAGE_GATHER4_L_O,
IMAGE_GATHER4_B_O,
IMAGE_GATHER4_B_CL_O,
IMAGE_GATHER4_LZ_O,

// Gather4 with comparison and offsets.
IMAGE_GATHER4_C_O,
IMAGE_GATHER4_C_CL_O,
IMAGE_GATHER4_C_L_O,
IMAGE_GATHER4_C_B_O,
IMAGE_GATHER4_C_B_CL_O,
IMAGE_GATHER4_C_LZ_O,

LAST_AMDGPU_ISD_NUMBER
};
Expand Down
14 changes: 0 additions & 14 deletions llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td
View file
Expand Up @@ -26,7 +26,6 @@ def RsrcIntrinsics : GenericTable {
}

foreach intr = !listconcat(AMDGPUBufferIntrinsics,
AMDGPUImageIntrinsics,
AMDGPUImageDimIntrinsics,
AMDGPUImageDimAtomicIntrinsics) in {
def : RsrcIntrinsic<!cast<AMDGPURsrcIntrinsic>(intr)>;
Expand Down Expand Up @@ -60,19 +59,6 @@ def : SourceOfDivergence<int_amdgcn_atomic_dec>;
def : SourceOfDivergence<int_amdgcn_ds_fadd>;
def : SourceOfDivergence<int_amdgcn_ds_fmin>;
def : SourceOfDivergence<int_amdgcn_ds_fmax>;
def : SourceOfDivergence<int_amdgcn_image_atomic_swap>;
def : SourceOfDivergence<int_amdgcn_image_atomic_add>;
def : SourceOfDivergence<int_amdgcn_image_atomic_sub>;
def : SourceOfDivergence<int_amdgcn_image_atomic_smin>;
def : SourceOfDivergence<int_amdgcn_image_atomic_umin>;
def : SourceOfDivergence<int_amdgcn_image_atomic_smax>;
def : SourceOfDivergence<int_amdgcn_image_atomic_umax>;
def : SourceOfDivergence<int_amdgcn_image_atomic_and>;
def : SourceOfDivergence<int_amdgcn_image_atomic_or>;
def : SourceOfDivergence<int_amdgcn_image_atomic_xor>;
def : SourceOfDivergence<int_amdgcn_image_atomic_inc>;
def : SourceOfDivergence<int_amdgcn_image_atomic_dec>;
def : SourceOfDivergence<int_amdgcn_image_atomic_cmpswap>;
def : SourceOfDivergence<int_amdgcn_buffer_atomic_swap>;
def : SourceOfDivergence<int_amdgcn_buffer_atomic_add>;
def : SourceOfDivergence<int_amdgcn_buffer_atomic_sub>;
Expand Down

0 comments on commit db6911a

Please sign in to comment.