[llvm] AMDGPU: Mark sendmsg intrinsics as nocallback (PR #126782)

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-llvm-ir

Author: Matt Arsenault (arsenm)

<details>
<summary>Changes</summary>

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

---
Full diff: https://github.com/llvm/llvm-project/pull/126782.diff


2 Files Affected:

- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+3-3) 
- (added) llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll (+86) 


``````````diff
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index eb7bde6999491..0662d2bfac99a 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/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">,
diff --git a/llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll b/llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll
new file mode 100644
index 0000000000000..da136fb747d9f
--- /dev/null
+++ b/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 }
+;.

``````````

</details>


https://github.com/llvm/llvm-project/pull/126782