AMDGPU: Mark sendmsg intrinsics as nocallback #126782
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@llvm/pr-subscribers-llvm-transforms @llvm/pr-subscribers-llvm-ir Author: Matt Arsenault (arsenm) ChangesThese can be nocallback, but no nosync. These return in the current Fixes #124802 Full diff: https://github.com/llvm/llvm-project/pull/126782.diff 2 Files Affected:
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 }
+;.
|
|
@llvm/pr-subscribers-backend-amdgpu Author: Matt Arsenault (arsenm) ChangesThese can be nocallback, but no nosync. These return in the current Fixes #124802 Full diff: https://github.com/llvm/llvm-project/pull/126782.diff 2 Files Affected:
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 }
+;.
|
This was referenced Feb 11, 2025
shiltian
reviewed
Feb 12, 2025
jayfoad
reviewed
Feb 13, 2025
| call void @sendmsg_rtn_is_norecurse() | ||
| ret void | ||
| } | ||
| ;. |
There was a problem hiding this comment.
Does ;. mean something special? Did update_test_checks insert this?
There was a problem hiding this comment.
Yes, update_test_checks adds this when checking the attribute groups for some reason
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
arsenm
force-pushed
the
users/arsenm/issue124802/mark-sendmsg-nocallback
branch
from
ae1bbdb to
67ada5c
Compare
jayfoad
approved these changes
Feb 13, 2025
flovent
pushed a commit
to flovent/llvm-project
that referenced
this pull request
Feb 13, 2025
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 llvm#124802
joaosaffran
pushed a commit
to joaosaffran/llvm-project
that referenced
this pull request
Feb 14, 2025
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 llvm#124802
sivan-shani
pushed a commit
to sivan-shani/llvm-project
that referenced
this pull request
Feb 24, 2025
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 llvm#124802
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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