AMDGPU: Mark sendmsg intrinsics as norecurse

[arsenm]

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

[arsenm]

Warning

This pull request is not mergeable via GitHub because a downstack PR is open. Once all requirements are satisfied, merge this PR as a stack on Graphite.
Learn more

• AMDGPU: Mark sendmsg intrinsics as norecurse #125016 👈 (View in Graphite)
• TableGen: Add intrinsic property for norecurse #125015
• AMDGPU: Add willreturn to sendmsg intrinsics #125014
• main

This stack of pull requests is managed by Graphite. Learn more about stacking.

[llvmbot]

@llvm/pr-subscribers-llvm-ir
@llvm/pr-subscribers-llvm-transforms

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

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

---

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

2 Files Affected:

• (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+3-3)
• (added) llvm/test/Transforms/FunctionAttrs/sendmsg-norecurse.ll (+90)

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index eb7bde69994913..99d68fc12b5002 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, 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">,
diff --git a/llvm/test/Transforms/FunctionAttrs/sendmsg-norecurse.ll b/llvm/test/Transforms/FunctionAttrs/sendmsg-norecurse.ll
new file mode 100644
index 00000000000000..45cc10f7876150
--- /dev/null
+++ b/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: {{.*}}

[jayfoad]

We cannot mark these as nocallback or nosync.

I think we can and should mark these as nocallback. I don't know how well it is documented, but I think in practice nocallback means that the intrinsic does not call back into user code synchronously, in the current thread, before the intrinsic returns.

[arsenm]

Replaced by #126782