AMDGPU: Mark sendmsg intrinsics as norecurse

We cannot mark these as nocallback or nosync. These send a message
to a host thread which could trigger arbitrary code. However this
signaled code can't directly cause this thread to start recursing.

FunctionAttrs manages to propagate norecurse to the caller, but
attributor does not.

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, IntrNoRecurse]>;
 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, IntrNoRecurse]>;
 
 
 // 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-norecurse.ll

@@ -0,0 +1,90 @@
+; 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
+
+; FIXME: attributor misses inferring norecurse
+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 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) #[[ATTR4:[0-9]+]]
+; ATTRIBUTOR-NEXT:    ret void
+;
+  call void @llvm.amdgcn.s.sendmsg(i32 1, i32 0)
+  ret void
+}
+
+define internal void @sendmsghalt_is_norecurse() {
+; FNATTRS: Function Attrs: norecurse nounwind
+; FNATTRS-LABEL: define internal void @sendmsghalt_is_norecurse(
+; FNATTRS-SAME: ) #[[ATTR1:[0-9]+]] {
+; FNATTRS-NEXT:    call void @llvm.amdgcn.s.sendmsghalt(i32 1, i32 0)
+; FNATTRS-NEXT:    ret void
+;
+; ATTRIBUTOR: Function Attrs: nounwind
+; ATTRIBUTOR-LABEL: define internal void @sendmsghalt_is_norecurse(
+; ATTRIBUTOR-SAME: ) #[[ATTR1:[0-9]+]] {
+; ATTRIBUTOR-NEXT:    call void @llvm.amdgcn.s.sendmsghalt(i32 1, i32 0)
+; ATTRIBUTOR-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: ) #[[ATTR0]] {
+; ATTRIBUTOR-NEXT:    [[RES:%.*]] = call i32 @llvm.amdgcn.s.sendmsg.rtn.i32(i32 1) #[[ATTR4]]
+; 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: ) #[[ATTR1]] {
+; ATTRIBUTOR-NEXT:    call void @sendmsg_is_norecurse() #[[ATTR5:[0-9]+]]
+; ATTRIBUTOR-NEXT:    call void @sendmsghalt_is_norecurse() #[[ATTR1]]
+; ATTRIBUTOR-NEXT:    call void @sendmsg_rtn_is_norecurse() #[[ATTR1]]
+; 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]+]] = { norecurse nounwind willreturn }
+;.
+; ATTRIBUTOR: attributes #[[ATTR0]] = { mustprogress nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR1]] = { nounwind }
+; ATTRIBUTOR: attributes #[[ATTR2:[0-9]+]] = { norecurse nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR3:[0-9]+]] = { norecurse nounwind }
+; ATTRIBUTOR: attributes #[[ATTR4]] = { willreturn }
+; ATTRIBUTOR: attributes #[[ATTR5]] = { nounwind willreturn }
+;.
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; COMMON: {{.*}}