maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[AMDGPU] Add IntrWillReturn to three intrinsics 

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

Differential Revision: https://reviews.llvm.org/D95436
This commit is contained in:
Sebastian Neubauer 2021-01-26 12:44:02 +01:00
parent 608ac62540
commit b36370d153
1 changed files with 3 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -1204,7 +1204,7 @@ class AMDGPUBufferAtomicFP : Intrinsic <
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.
@ -1289,7 +1289,7 @@ def int_amdgcn_s_getreg :
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
@ -1725,7 +1725,7 @@ 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.