[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:1581
   ClangBuiltin<"__builtin_amdgcn_ds_swizzle">,
-  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
-            [IntrNoMem, IntrConvergent, IntrWillReturn,
+  DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
+            [IntrNoMem, IntrConvergent,
----------------
Shouldn't be default, nosync is wrong


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1649-1686
 def int_amdgcn_icmp :
-  Intrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>, llvm_i32_ty],
-            [IntrNoMem, IntrConvergent, IntrWillReturn,
+  DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>, llvm_i32_ty],
+            [IntrNoMem, IntrConvergent,
              ImmArg<ArgIndex<2>>]>;
 
 def int_amdgcn_fcmp :
+  DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty],
----------------
Not default, probably should not get nosync


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1990
 def int_amdgcn_permlane64 :
-  Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
-            [IntrNoMem, IntrConvergent, IntrWillReturn]>;
+  DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty],
+            [IntrNoMem, IntrConvergent]>;
----------------
Shouldn't get nosync


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1993-2012
 def int_amdgcn_ds_add_gs_reg_rtn :
   ClangBuiltin<"__builtin_amdgcn_ds_add_gs_reg_rtn">,
-  Intrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty],
-            [ImmArg<ArgIndex<1>>, IntrHasSideEffects, IntrWillReturn]>;
+  DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty],
+            [ImmArg<ArgIndex<1>>, IntrHasSideEffects]>;
 
 def int_amdgcn_ds_sub_gs_reg_rtn :
   ClangBuiltin<"__builtin_amdgcn_ds_sub_gs_reg_rtn">,
----------------
Not sure about these ones, but maybe shouldn't get nosync


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:2020-2044
 class AMDGPUWmmaIntrinsic<LLVMType AB, LLVMType CD> :
-  Intrinsic<
+  DefaultAttrsIntrinsic<
     [CD],               // %D
     [
       AB,               // %A
       AB,               // %B
       LLVMMatchType<0>, // %C
----------------
Also not sure about nosync for these


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:2457-2465
 def int_amdgcn_fdiv_fast : Intrinsic<
   [llvm_float_ty], [llvm_float_ty, llvm_float_ty],
-  [IntrNoMem, IntrSpeculatable, IntrWillReturn]
+  [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback, IntrNoFree]
 >;
 
 // Represent a relocation constant.
 def int_amdgcn_reloc_constant : Intrinsic<
----------------
Defaults


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D135155