[AMDGPU] Add IntrWillReturn to three intrinsics

[sebastian-ne]

Summary

None of these can terminate a wave or lane.
With these, all intrinsic are IntrWillReturn except those that change
exec or can terminate the wave.

Not marking intrinsics as WillReturn may prevent optimizations in the
future: https://lists.llvm.org/pipermail/llvm-dev/2021-January/148047.html

int_r600_kill is marked as IntrWillReturn, is that correct?

[arsenm]

We should probably switch to using the new default attribute list feature, but I guess that's a separate change

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

 
 class AMDGPUBufferAtomicFP : Intrinsic <
   [llvm_anyfloat_ty],
   [LLVMMatchType<0>, // vdata(VGPR)
    llvm_v4i32_ty,    // rsrc(SGPR)
    llvm_i32_ty,      // vindex(VGPR)
    llvm_i32_ty,      // offset(SGPR/VGPR/imm)
    llvm_i1_ty],      // slc(imm)
-  [ImmArg<ArgIndex<4>>], "", [SDNPMemOperand]>,
+  [ImmArg<ArgIndex<4>>, IntrWillReturn], "", [SDNPMemOperand]>,
   AMDGPURsrcIntrinsic<1, 0>;
 
 // Legacy form of the intrinsic. raw and struct forms should be preferred.
 def int_amdgcn_buffer_atomic_fadd : AMDGPUBufferAtomicFP;
 } // defset AMDGPUBufferIntrinsics
 
 // Uses that do not set the done bit should set IntrWriteMem on the
 // call site.

 
 // Note this can be used to set FP environment properties that are
 // unsafe to change in non-strictfp functions. The register properties
 // available (and value required to access them) may differ per
 // subtarget. llvm.amdgcn.s.setreg(hwmode, value)
 def int_amdgcn_s_setreg :
   GCCBuiltin<"__builtin_amdgcn_s_setreg">,
   Intrinsic<[], [llvm_i32_ty, llvm_i32_ty],
-  [IntrNoMem, IntrHasSideEffects, ImmArg<ArgIndex<0>>]
+  [IntrNoMem, IntrHasSideEffects, IntrWillReturn, ImmArg<ArgIndex<0>>]
 >;
 
 // int_amdgcn_s_getpc is provided to allow a specific style of position
 // independent code to determine the high part of its address when it is
 // known (through convention) that the code and any data of interest does
 // not cross a 4Gb address boundary. Use for any other purpose may not
 // produce the desired results as optimizations may cause code movement,
 // especially as we explicitly use IntrNoMem to allow optimizations.

 def int_amdgcn_global_atomic_csub : AMDGPUGlobalAtomicRtn<llvm_i32_ty>;
 
 // uint4 llvm.amdgcn.image.bvh.intersect.ray <node_ptr>, <ray_extent>, <ray_origin>,
 //                                           <ray_dir>, <ray_inv_dir>, <texture_descr>
 def int_amdgcn_image_bvh_intersect_ray :
   Intrinsic<[llvm_v4i32_ty],
             [llvm_anyint_ty, llvm_float_ty, llvm_v4f32_ty, llvm_anyvector_ty,
              LLVMMatchType<1>, llvm_v4i32_ty],
-            [IntrReadMem]>;
+            [IntrReadMem, IntrWillReturn]>;
 
 //===----------------------------------------------------------------------===//
 // Deep learning intrinsics.
 //===----------------------------------------------------------------------===//
 
 // f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)
 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
 def int_amdgcn_fdot2 :