[llvm] AMDGPU: Mark sendmsg intrinsics as nocallback (PR #126782)
Matt Arsenault via llvm-commits
llvm-commits at lists.llvm.org
Thu Feb 13 01:46:01 PST 2025
https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/126782
>From 6ace5cde2c40b9b0d5ef1280aff877436b5cc99c 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 1/2] 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
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 6 +-
.../FunctionAttrs/sendmsg-nocallback.ll | 86 +++++++++++++++++++
2 files changed, 89 insertions(+), 3 deletions(-)
create mode 100644 llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll
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 }
+;.
>From 67ada5c3d96dd982ad97c512c53d98c2a291a694 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 13 Feb 2025 16:30:42 +0700
Subject: [PATCH 2/2] fix rtn case
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 2 +-
.../FunctionAttrs/sendmsg-nocallback.ll | 35 +++++++++----------
2 files changed, 18 insertions(+), 19 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 0662d2bfac99a..d5d185ebc12f6 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -246,7 +246,7 @@ def int_amdgcn_s_sendmsghalt : ClangBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
// 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, IntrNoRecurse]>;
+ [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>;
// 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
index da136fb747d9f..4d5db3263f527 100644
--- a/llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll
+++ b/llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll
@@ -2,6 +2,8 @@
; 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
+; Make sure norecurse is inferred on the calling functions
+
define internal void @sendmsg_is_norecurse() {
; FNATTRS: Function Attrs: mustprogress norecurse nounwind willreturn
; FNATTRS-LABEL: define internal void @sendmsg_is_norecurse(
@@ -12,7 +14,7 @@ define internal void @sendmsg_is_norecurse() {
; 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: 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)
@@ -37,10 +39,10 @@ define internal i32 @sendmsg_rtn_is_norecurse() {
; 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: Function Attrs: mustprogress norecurse 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-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)
@@ -54,12 +56,12 @@ define void @user() {
; FNATTRS-NEXT: call void @sendmsg_rtn_is_norecurse()
; FNATTRS-NEXT: ret void
;
-; ATTRIBUTOR: Function Attrs: nounwind
+; ATTRIBUTOR: Function Attrs: norecurse 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-SAME: ) #[[ATTR1]] {
+; ATTRIBUTOR-NEXT: call void @sendmsg_is_norecurse() #[[ATTR5:[0-9]+]]
+; ATTRIBUTOR-NEXT: call void @sendmsghalt_is_norecurse() #[[ATTR6:[0-9]+]]
+; ATTRIBUTOR-NEXT: call void @sendmsg_rtn_is_norecurse() #[[ATTR6]]
; ATTRIBUTOR-NEXT: ret void
;
call void @sendmsg_is_norecurse()
@@ -71,16 +73,13 @@ define void @user() {
; 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 }
+; FNATTRS: attributes #[[ATTR3:[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 }
+; ATTRIBUTOR: attributes #[[ATTR2:[0-9]+]] = { nocallback nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR3:[0-9]+]] = { nocallback nounwind }
+; ATTRIBUTOR: attributes #[[ATTR4]] = { willreturn }
+; ATTRIBUTOR: attributes #[[ATTR5]] = { nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR6]] = { nounwind }
;.
More information about the llvm-commits
mailing list