[llvm-branch-commits] [llvm] AMDGPU: Mark sendmsg intrinsics as norecurse (PR #125016)

[Matt Arsenault via llvm-branch-commits]

https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/125016

>From 63cb322d5f05ba8bc29cabc841119204302c757f Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 30 Jan 2025 09:34:09 +0700
Subject: [PATCH] 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
---
 llvm/include/llvm/IR/IntrinsicsAMDGPU.td      |  6 +-
 .../FunctionAttrs/sendmsg-norecurse.ll        | 90 +++++++++++++++++++
 2 files changed, 93 insertions(+), 3 deletions(-)
 create mode 100644 llvm/test/Transforms/FunctionAttrs/sendmsg-norecurse.ll

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index eb7bde69994913c..99d68fc12b5002b 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 000000000000000..45cc10f78761509
--- /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: {{.*}}