[PATCH] D135155: [AMDGPU] Annotate the intrinsics to be default and nocallback

[Matt Arsenault via Phabricator via llvm-commits]

arsenm added inline comments.


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:467
     [IntrArgMemOnly, IntrWillReturn, NoCapture<ArgIndex<0>>,
-     ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>]
+     ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, IntrNoCallback]
 >;
----------------
NoFree


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:495
    [IntrConvergent, IntrWillReturn, IntrArgMemOnly,
-    NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<1>>],
+    NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<1>>, IntrNoCallback],
    "",
----------------
NoFree


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:993
 
 class AMDGPUBufferLoad<LLVMType data_ty = llvm_any_ty> : Intrinsic <
   [data_ty],
----------------
This can use Default 


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1081
 
 class AMDGPUStructBufferStore<LLVMType data_ty = llvm_any_ty> : Intrinsic <
   [],
----------------
Can be default


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1104
    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
-  [ImmArg<ArgIndex<4>>, IntrWillReturn], "", [SDNPMemOperand]>,
+  [ImmArg<ArgIndex<4>>, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand]>,
   AMDGPURsrcIntrinsic<1, 0>;
----------------
NoFree


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1177
 // Obsolescent tbuffer intrinsics.
 def int_amdgcn_tbuffer_load : Intrinsic <
     [llvm_any_ty],    // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
----------------
Can be default


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1243
 
 def int_amdgcn_struct_tbuffer_load : Intrinsic <
     [llvm_any_ty],    // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
----------------
All these buffer loads can be default 


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1356
 // call site.
 def int_amdgcn_exp : Intrinsic <[], [
   llvm_i32_ty,       // tgt,
----------------
exp_* can be default 


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1398-1408
 def int_amdgcn_buffer_wbinvl1_sc :
   ClangBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
-  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;
+  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>;
 
 def int_amdgcn_buffer_wbinvl1 :
   ClangBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
+  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>;
----------------
Default 


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1692-1715
 def int_amdgcn_alignbyte : ClangBuiltin<"__builtin_amdgcn_alignbyte">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
-  [IntrNoMem, IntrSpeculatable, IntrWillReturn]
+  [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback]
 >;
 
 def int_amdgcn_mul_i24 : Intrinsic<[llvm_i32_ty],
   [llvm_i32_ty, llvm_i32_ty],
----------------
Defaults


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1855-1861
 def int_amdgcn_s_dcache_inv_vol :
   ClangBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
-  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;
+  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>;
 
 def int_amdgcn_buffer_wbinvl1_vol :
   ClangBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
+  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>;
----------------
Defaults


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D135155/new/

https://reviews.llvm.org/D135155