[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/Intrinsics.td:515
                                    [IntrReadMem], "llvm.read_register">;
-def int_write_register : Intrinsic<[], [llvm_metadata_ty, llvm_anyint_ty],
+def int_write_register : DefaultAttrsIntrinsic<[], [llvm_metadata_ty, llvm_anyint_ty],
                                    [], "llvm.write_register">;
----------------
Probably best to assume this has side effects


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:168
   ClangBuiltin<"__builtin_amdgcn_dispatch_id">,
-  Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>;
+  DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>;
 
----------------
I thought WIllReturn was in the default set?


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:209
   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],
-  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>;
+  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrNoCallback]>;
 def int_amdgcn_s_sendmsghalt : ClangBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
----------------
I'm not sure sendmsg* should be no IntrNoCallback. You can use it for device to host messaging, which I guess could trigger some visible change


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:273
   [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
   [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<2>>, IntrWillReturn]
 >;
----------------
IntrWillReturn default


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D135155