AMDGPU: Mark exp/exp.compr as inaccessiblememonly

[arsenm]

Summary

Should also be marked writeonly, but I think that would require splitting the version with done set to a separate intrinsic

Test change is only from renumbering the attribute group numbers,
which for some reason the generated check lines consider.

[nhaehnle]

LGTM.

I don't quite understand why you think this can't be writeonly? (The flip side is, I don't see how having this writeonly would help codegen)

[arsenm]

In D63406#1545891, @nhaehnle wrote:

LGTM.

I don't quite understand why you think this can't be writeonly? (The flip side is, I don't see how having this writeonly would help codegen)

I'm worried about reordering exp with done set before an exp without it

[arsenm]

r363560

include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_exp : Intrinsic <[], [
   llvm_i32_ty,       // en
   llvm_any_ty,       // src0 (f32 or i32)
   LLVMMatchType<0>,  // src1
   LLVMMatchType<0>,  // src2
   LLVMMatchType<0>,  // src3
   llvm_i1_ty,        // done
   llvm_i1_ty         // vm
   ],
-  [ImmArg<0>, ImmArg<1>, ImmArg<6>, ImmArg<7>]
+  [ImmArg<0>, ImmArg<1>, ImmArg<6>, ImmArg<7>, IntrInaccessibleMemOnly]
 >;
 
 // exp with compr bit set.
 def int_amdgcn_exp_compr : Intrinsic <[], [
   llvm_i32_ty,       // tgt,
   llvm_i32_ty,       // en
   llvm_anyvector_ty, // src0 (v2f16 or v2i16)
   LLVMMatchType<0>,  // src1
   llvm_i1_ty,        // done
   llvm_i1_ty],       // vm
-  [ImmArg<0>, ImmArg<1>, ImmArg<4>, ImmArg<5>]
+  [ImmArg<0>, ImmArg<1>, ImmArg<4>, ImmArg<5>, IntrInaccessibleMemOnly]
 >;
 
 def int_amdgcn_buffer_wbinvl1_sc :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_buffer_wbinvl1 :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,

test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll

 ; CHECK-NEXT:    ret double 2.000000e+00
 ;
   %val = call double @llvm.amdgcn.rcp.f64(double 0.5) nounwind readnone
   ret double %val
 }
 
 define float @test_constant_fold_rcp_f32_43() nounwind {
 ; CHECK-LABEL: @test_constant_fold_rcp_f32_43(
-; CHECK-NEXT:    [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #4
+; CHECK-NEXT:    [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #5
 ; CHECK-NEXT:    ret float [[VAL]]
 ;
   %val = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) nounwind readnone
   ret float %val
 }
 
 define double @test_constant_fold_rcp_f64_43() nounwind {
 ; CHECK-LABEL: @test_constant_fold_rcp_f64_43(
-; CHECK-NEXT:    [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) #4
+; CHECK-NEXT:    [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) #5
 ; CHECK-NEXT:    ret double [[VAL]]
 ;
   %val = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) nounwind readnone
   ret double %val
 }
 
 ; --------------------------------------------------------------------
 ; llvm.amdgcn.rsq

 ; CHECK-NEXT:    ret i64 0
 ;
   %result = call i64 @llvm.amdgcn.icmp.i64.i32(i32 9, i32 8, i32 32)
   ret i64 %result
 }
 
 define i64 @icmp_constant_inputs_true() {
 ; CHECK-LABEL: @icmp_constant_inputs_true(
-; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.read_register.i64(metadata !0) #5
+; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.read_register.i64(metadata !0) #6
 ; CHECK-NEXT:    ret i64 [[RESULT]]
 ;
   %result = call i64 @llvm.amdgcn.icmp.i64.i32(i32 9, i32 8, i32 34)
   ret i64 %result
 }
 
 define i64 @icmp_constant_to_rhs_slt(i32 %x) {
 ; CHECK-LABEL: @icmp_constant_to_rhs_slt(

 ; CHECK-NEXT:    ret i64 0
 ;
   %result = call i64 @llvm.amdgcn.fcmp.i64.f32(float 2.0, float 4.0, i32 1)
   ret i64 %result
 }
 
 define i64 @fcmp_constant_inputs_true() {
 ; CHECK-LABEL: @fcmp_constant_inputs_true(
-; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.read_register.i64(metadata !0) #5
+; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.read_register.i64(metadata !0) #6
 ; CHECK-NEXT:    ret i64 [[RESULT]]
 ;
   %result = call i64 @llvm.amdgcn.fcmp.i64.f32(float 2.0, float 4.0, i32 4)
   ret i64 %result
 }
 
 define i64 @fcmp_constant_to_rhs_olt(float %x) {
 ; CHECK-LABEL: @fcmp_constant_to_rhs_olt(

 ; --------------------------------------------------------------------
 
 declare i32 @llvm.amdgcn.readfirstlane(i32)
 
 @gv = constant i32 0
 
 define amdgpu_kernel void @readfirstlane_constant(i32 %arg) {
 ; CHECK-LABEL: @readfirstlane_constant(
-; CHECK-NEXT: %var = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
-; CHECK-NEXT: store volatile i32 %var, i32* undef, align 4
-; CHECK-NEXT: store volatile i32 0, i32* undef, align 4
-; CHECK-NEXT: store volatile i32 123, i32* undef, align 4
-; CHECK-NEXT: store volatile i32 ptrtoint (i32* @gv to i32), i32* undef, align 4
-; CHECK-NEXT: store volatile i32 undef, i32* undef, align 4
+; CHECK-NEXT:    [[VAR:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[ARG:%.*]])
+; CHECK-NEXT:    store volatile i32 [[VAR]], i32* undef, align 4
+; CHECK-NEXT:    store volatile i32 0, i32* undef, align 4
+; CHECK-NEXT:    store volatile i32 123, i32* undef, align 4
+; CHECK-NEXT:    store volatile i32 ptrtoint (i32* @gv to i32), i32* undef, align 4
+; CHECK-NEXT:    store volatile i32 undef, i32* undef, align 4
+; CHECK-NEXT:    ret void
+;
   %var = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
   %zero = call i32 @llvm.amdgcn.readfirstlane(i32 0)
   %imm = call i32 @llvm.amdgcn.readfirstlane(i32 123)
   %constexpr = call i32 @llvm.amdgcn.readfirstlane(i32 ptrtoint (i32* @gv to i32))
   %undef = call i32 @llvm.amdgcn.readfirstlane(i32 undef)
   store volatile i32 %var, i32* undef
   store volatile i32 %zero, i32* undef
   store volatile i32 %imm, i32* undef
   store volatile i32 %constexpr, i32* undef
   store volatile i32 %undef, i32* undef
   ret void
 }
 
 ; --------------------------------------------------------------------
 ; llvm.amdgcn.readlane
 ; --------------------------------------------------------------------
 
 declare i32 @llvm.amdgcn.readlane(i32, i32)
 
 define amdgpu_kernel void @readlane_constant(i32 %arg, i32 %lane) {
 ; CHECK-LABEL: @readlane_constant(
-; CHECK-NEXT: %var = call i32 @llvm.amdgcn.readlane(i32 %arg, i32 7)
-; CHECK-NEXT: store volatile i32 %var, i32* undef, align 4
-; CHECK-NEXT: store volatile i32 0, i32* undef, align 4
-; CHECK-NEXT: store volatile i32 123, i32* undef, align 4
-; CHECK-NEXT: store volatile i32 ptrtoint (i32* @gv to i32), i32* undef, align 4
-; CHECK-NEXT: store volatile i32 undef, i32* undef, align 4
+; CHECK-NEXT:    [[VAR:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[ARG:%.*]], i32 7)
+; CHECK-NEXT:    store volatile i32 [[VAR]], i32* undef, align 4
+; CHECK-NEXT:    store volatile i32 0, i32* undef, align 4
+; CHECK-NEXT:    store volatile i32 123, i32* undef, align 4
+; CHECK-NEXT:    store volatile i32 ptrtoint (i32* @gv to i32), i32* undef, align 4
+; CHECK-NEXT:    store volatile i32 undef, i32* undef, align 4
+; CHECK-NEXT:    ret void
+;
   %var = call i32 @llvm.amdgcn.readlane(i32 %arg, i32 7)
   %zero = call i32 @llvm.amdgcn.readlane(i32 0, i32 %lane)
   %imm = call i32 @llvm.amdgcn.readlane(i32 123, i32 %lane)
   %constexpr = call i32 @llvm.amdgcn.readlane(i32 ptrtoint (i32* @gv to i32), i32 %lane)
   %undef = call i32 @llvm.amdgcn.readlane(i32 undef, i32 %lane)
   store volatile i32 %var, i32* undef
   store volatile i32 %zero, i32* undef
   store volatile i32 %imm, i32* undef