AMDGPU: Mark sendmsg intrinsics as nocallback

These can be nocallback, but no nosync. These return in the current
thread immediately, but may send a signal to the host, which could trigger
the async execution of code in the module.

Fixes #124802

Author: arsenm

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

@@ -237,16 +237,16 @@ def int_amdgcn_reloc_constant : DefaultAttrsIntrinsic<
 // the second one is copied to m0
 def int_amdgcn_s_sendmsg : ClangBuiltin<"__builtin_amdgcn_s_sendmsg">,
   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],
-  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;
+  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>;
 def int_amdgcn_s_sendmsghalt : ClangBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],
-  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>;
+  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrNoCallback]>;
 
 
 // gfx11 intrinsic
 // The first parameter is s_sendmsg immediate (i16). Return type is i32 or i64.
 def int_amdgcn_s_sendmsg_rtn : Intrinsic <[llvm_anyint_ty], [llvm_i32_ty],
-  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;
+  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoRecurse]>;
 
 // Vanilla workgroup sync-barrier
 def int_amdgcn_s_barrier : ClangBuiltin<"__builtin_amdgcn_s_barrier">,

llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll

@@ -0,0 +1,86 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --check-attributes --check-globals all --version 5
+; RUN: opt -S -passes=function-attrs < %s | FileCheck --check-prefixes=COMMON,FNATTRS %s
+; RUN: opt -S -passes=attributor-light < %s | FileCheck --check-prefixes=COMMON,ATTRIBUTOR %s
+
+define internal void @sendmsg_is_norecurse() {
+; FNATTRS: Function Attrs: mustprogress norecurse nounwind willreturn
+; FNATTRS-LABEL: define internal void @sendmsg_is_norecurse(
+; FNATTRS-SAME: ) #[[ATTR0:[0-9]+]] {
+; FNATTRS-NEXT:    call void @llvm.amdgcn.s.sendmsg(i32 1, i32 0)
+; FNATTRS-NEXT:    ret void
+;
+; ATTRIBUTOR: Function Attrs: mustprogress norecurse nounwind willreturn
+; ATTRIBUTOR-LABEL: define internal void @sendmsg_is_norecurse(
+; ATTRIBUTOR-SAME: ) #[[ATTR0:[0-9]+]] {
+; ATTRIBUTOR-NEXT:    call void @llvm.amdgcn.s.sendmsg(i32 1, i32 0) #[[ATTR7:[0-9]+]]
+; ATTRIBUTOR-NEXT:    ret void
+;
+  call void @llvm.amdgcn.s.sendmsg(i32 1, i32 0)
+  ret void
+}
+
+define internal void @sendmsghalt_is_norecurse() {
+; COMMON: Function Attrs: norecurse nounwind
+; COMMON-LABEL: define internal void @sendmsghalt_is_norecurse(
+; COMMON-SAME: ) #[[ATTR1:[0-9]+]] {
+; COMMON-NEXT:    call void @llvm.amdgcn.s.sendmsghalt(i32 1, i32 0)
+; COMMON-NEXT:    ret void
+;
+  call void @llvm.amdgcn.s.sendmsghalt(i32 1, i32 0)
+  ret void
+}
+
+define internal i32 @sendmsg_rtn_is_norecurse() {
+; FNATTRS: Function Attrs: mustprogress norecurse nounwind willreturn
+; FNATTRS-LABEL: define internal i32 @sendmsg_rtn_is_norecurse(
+; FNATTRS-SAME: ) #[[ATTR0]] {
+; FNATTRS-NEXT:    [[RES:%.*]] = call i32 @llvm.amdgcn.s.sendmsg.rtn.i32(i32 1)
+; FNATTRS-NEXT:    ret i32 [[RES]]
+;
+; ATTRIBUTOR: Function Attrs: mustprogress nounwind willreturn
+; ATTRIBUTOR-LABEL: define internal i32 @sendmsg_rtn_is_norecurse(
+; ATTRIBUTOR-SAME: ) #[[ATTR2:[0-9]+]] {
+; ATTRIBUTOR-NEXT:    [[RES:%.*]] = call i32 @llvm.amdgcn.s.sendmsg.rtn.i32(i32 1) #[[ATTR7]]
+; ATTRIBUTOR-NEXT:    ret i32 [[RES]]
+;
+  %res = call i32 @llvm.amdgcn.s.sendmsg.rtn(i32 1)
+  ret i32 %res
+}
+
+define void @user() {
+; FNATTRS-LABEL: define void @user() {
+; FNATTRS-NEXT:    call void @sendmsg_is_norecurse()
+; FNATTRS-NEXT:    call void @sendmsghalt_is_norecurse()
+; FNATTRS-NEXT:    call void @sendmsg_rtn_is_norecurse()
+; FNATTRS-NEXT:    ret void
+;
+; ATTRIBUTOR: Function Attrs: nounwind
+; ATTRIBUTOR-LABEL: define void @user(
+; ATTRIBUTOR-SAME: ) #[[ATTR3:[0-9]+]] {
+; ATTRIBUTOR-NEXT:    call void @sendmsg_is_norecurse() #[[ATTR8:[0-9]+]]
+; ATTRIBUTOR-NEXT:    call void @sendmsghalt_is_norecurse() #[[ATTR3]]
+; ATTRIBUTOR-NEXT:    call void @sendmsg_rtn_is_norecurse() #[[ATTR3]]
+; ATTRIBUTOR-NEXT:    ret void
+;
+  call void @sendmsg_is_norecurse()
+  call void @sendmsghalt_is_norecurse()
+  call void @sendmsg_rtn_is_norecurse()
+  ret void
+}
+;.
+; FNATTRS: attributes #[[ATTR0]] = { mustprogress norecurse nounwind willreturn }
+; FNATTRS: attributes #[[ATTR1]] = { norecurse nounwind }
+; FNATTRS: attributes #[[ATTR2:[0-9]+]] = { nocallback nounwind willreturn }
+; FNATTRS: attributes #[[ATTR3:[0-9]+]] = { norecurse nounwind willreturn }
+; FNATTRS: attributes #[[ATTR4:[0-9]+]] = { nocallback nounwind }
+;.
+; ATTRIBUTOR: attributes #[[ATTR0]] = { mustprogress norecurse nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR1]] = { norecurse nounwind }
+; ATTRIBUTOR: attributes #[[ATTR2]] = { mustprogress nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR3]] = { nounwind }
+; ATTRIBUTOR: attributes #[[ATTR4:[0-9]+]] = { nocallback nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR5:[0-9]+]] = { norecurse nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR6:[0-9]+]] = { nocallback nounwind }
+; ATTRIBUTOR: attributes #[[ATTR7]] = { willreturn }
+; ATTRIBUTOR: attributes #[[ATTR8]] = { nounwind willreturn }
+;.