[clang] 5ad3dd8 - [AMDGPU] Add more attributes to some intrinsics to help with attribute propagation (#222222)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 9 11:57:24 PDT 2026
Author: Shilei Tian
Date: 2026-09-09T14:57:17-04:00
New Revision: 5ad3dd8a19d4ef59a13ac0052abe5fbca7b2823c
URL: https://github.com/llvm/llvm-project/commit/5ad3dd8a19d4ef59a13ac0052abe5fbca7b2823c
DIFF: https://github.com/llvm/llvm-project/commit/5ad3dd8a19d4ef59a13ac0052abe5fbca7b2823c.diff
LOG: [AMDGPU] Add more attributes to some intrinsics to help with attribute propagation (#222222)
Added:
Modified:
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-async-load-store-lds.cl
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll
Removed:
################################################################################
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-async-load-store-lds.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-async-load-store-lds.cl
index 34c17a5e204fe..c8e372c8731df 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-async-load-store-lds.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-async-load-store-lds.cl
@@ -149,10 +149,10 @@ void test_amdgcn_ds_atomic_async_barrier_arrive_b64(local long* addr)
}
// CHECK-GFX1250-LABEL: define dso_local void @test_amdgcn_ds_atomic_barrier_arrive_rtn_b64(
-// CHECK-GFX1250-SAME: ptr addrspace(3) noundef captures(none) [[ADDR:%.*]], i64 noundef [[DATA:%.*]], ptr nofree noundef writeonly captures(none) initializes((0, 8)) [[OUT:%.*]]) local_unnamed_addr #[[ATTR4:[0-9]+]] {
+// CHECK-GFX1250-SAME: ptr addrspace(3) nofree noundef captures(none) [[ADDR:%.*]], i64 noundef [[DATA:%.*]], ptr nofree noundef writeonly captures(none) initializes((0, 8)) [[OUT:%.*]]) local_unnamed_addr #[[ATTR4:[0-9]+]] {
// CHECK-GFX1250-NEXT: [[ENTRY:.*:]]
// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i64 @llvm.amdgcn.ds.atomic.barrier.arrive.rtn.b64(ptr addrspace(3) [[ADDR]], i64 [[DATA]])
-// CHECK-GFX1250-NEXT: store i64 [[TMP0]], ptr [[OUT]], align 8, !tbaa [[LONG_TBAA7:![0-9]+]]
+// CHECK-GFX1250-NEXT: store i64 [[TMP0]], ptr [[OUT]], align 8, !tbaa [[LONG_TBAA8:![0-9]+]]
// CHECK-GFX1250-NEXT: ret void
//
void test_amdgcn_ds_atomic_barrier_arrive_rtn_b64(local long* addr, long data, long *out)
@@ -160,8 +160,8 @@ void test_amdgcn_ds_atomic_barrier_arrive_rtn_b64(local long* addr, long data, l
*out = __builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64(addr, data);
}
//.
-// CHECK-GFX1250: [[META5:![0-9]+]] = !{!"omnipotent char", [[META6:![0-9]+]], i64 0}
-// CHECK-GFX1250: [[META6]] = !{!"Simple C/C++ TBAA"}
-// CHECK-GFX1250: [[LONG_TBAA7]] = !{[[META8:![0-9]+]], [[META8]], i64 0}
-// CHECK-GFX1250: [[META8]] = !{!"long", [[META5]], i64 0}
+// CHECK-GFX1250: [[META6:![0-9]+]] = !{!"omnipotent char", [[META7:![0-9]+]], i64 0}
+// CHECK-GFX1250: [[META7]] = !{!"Simple C/C++ TBAA"}
+// CHECK-GFX1250: [[LONG_TBAA8]] = !{[[META9:![0-9]+]], [[META9]], i64 0}
+// CHECK-GFX1250: [[META9]] = !{!"long", [[META6]], i64 0}
//.
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 6b15075f75b9c..39c843c5c5c87 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -172,12 +172,12 @@ defm int_amdgcn_cluster_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_na
<"__builtin_amdgcn_cluster_workgroup_id">;
def int_amdgcn_cluster_workgroup_flat_id:
ClangBuiltin<"__builtin_amdgcn_cluster_workgroup_flat_id">,
- Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
+ DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
defm int_amdgcn_cluster_workgroup_max_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
<"__builtin_amdgcn_cluster_workgroup_max_id">;
def int_amdgcn_cluster_workgroup_max_flat_id:
ClangBuiltin<"__builtin_amdgcn_cluster_workgroup_max_flat_id">,
- Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
+ DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_dispatch_ptr :
DefaultAttrsIntrinsic<[LLVMQualPointerType<4>], [],
@@ -269,16 +269,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, IntrNoCallback]>;
+ [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
def int_amdgcn_s_sendmsghalt : ClangBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],
- [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrNoCallback]>;
+ [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrNoCallback, IntrNoFree]>;
// 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, IntrNoCallback]>;
+ [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
// Vanilla workgroup sync-barrier
def int_amdgcn_s_barrier : ClangBuiltin<"__builtin_amdgcn_s_barrier">,
@@ -2251,7 +2251,7 @@ def int_amdgcn_s_sleep :
def int_amdgcn_s_sleep_var
: ClangBuiltin<"__builtin_amdgcn_s_sleep_var">,
Intrinsic<[], [llvm_i32_ty],
- [IntrNoMem, IntrHasSideEffects, IntrWillReturn]> {
+ [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]> {
}
def int_amdgcn_s_nop :
@@ -3183,7 +3183,7 @@ class AMDGPUSWmmacIntrinsicIdxReuse<LLVMType A, LLVMType B, LLVMType CD, LLVMTyp
llvm_i1_ty, // matrix_a_reuse
llvm_i1_ty, // matrix_b_reuse
],
- [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCreateUndefOrPoison,
+ [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, IntrNoCreateUndefOrPoison,
ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]
>;
@@ -3196,7 +3196,7 @@ class AMDGPUSWmmacIntrinsicIdx<LLVMType A, LLVMType B, LLVMType CD, LLVMType Ind
LLVMMatchType<0>, // %C
Index // %Sparsity index for A
],
- [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCreateUndefOrPoison]
+ [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, IntrNoCreateUndefOrPoison]
>;
class AMDGPUSWmmacIntrinsicIUIdx<LLVMType A, LLVMType B, LLVMType CD, LLVMType Index> :
@@ -3211,7 +3211,8 @@ class AMDGPUSWmmacIntrinsicIUIdx<LLVMType A, LLVMType B, LLVMType CD, LLVMType I
Index, // %Sparsity index for A
llvm_i1_ty, // %clamp
],
- [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCreateUndefOrPoison, ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<6>>]
+ [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, IntrNoCreateUndefOrPoison,
+ ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<6>>]
>;
defset list<Intrinsic> AMDGPUWMMAIntrinsicsGFX12 = {
@@ -3851,13 +3852,14 @@ def int_amdgcn_ds_atomic_async_barrier_arrive_b64 :
ClangBuiltin<"__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64">,
Intrinsic<[], [local_ptr_ty],
// Atomically updates LDS and also ASYNC_CNT which is modeled as InaccessibleMem.
- [IntrConvergent, IntrWillReturn, IntrInaccessibleMemOrArgMemOnly],
+ [IntrConvergent, IntrWillReturn, IntrInaccessibleMemOrArgMemOnly, IntrNoCallback, IntrNoFree],
"", [SDNPMemOperand]>;
def int_amdgcn_ds_atomic_barrier_arrive_rtn_b64 :
ClangBuiltin<"__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64">,
Intrinsic<[llvm_i64_ty], [local_ptr_ty, llvm_i64_ty],
- [IntrConvergent, IntrWillReturn, IntrArgMemOnly, NoCapture<ArgIndex<0>>],
+ [IntrConvergent, IntrWillReturn, IntrArgMemOnly, IntrNoCallback, IntrNoFree,
+ NoCapture<ArgIndex<0>>],
"", [SDNPMemOperand]>;
def int_amdgcn_s_monitor_sleep :
@@ -4273,7 +4275,7 @@ class AMDGPUSWmmacIntrinsicABIdx<LLVMType A, LLVMType B, LLVMType CD, LLVMType I
llvm_i1_ty, // matrix_a_reuse
llvm_i1_ty, // matrix_b_reuse
],
- [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCreateUndefOrPoison,
+ [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, IntrNoCreateUndefOrPoison,
ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<6>>, ImmArg<ArgIndex<7>>]
>;
@@ -4291,7 +4293,7 @@ class AMDGPUSWmmacIntrinsicABIdxClamp<LLVMType A, LLVMType B, LLVMType CD, LLVMT
llvm_i1_ty, // matrix_b_reuse
llvm_i1_ty, // %clamp
],
- [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCreateUndefOrPoison,
+ [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, IntrNoCreateUndefOrPoison,
ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<6>>, ImmArg<ArgIndex<7>>, ImmArg<ArgIndex<8>>]
>;
diff --git a/llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll b/llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll
index 04575e4cc940b..784f5b950ecca 100644
--- a/llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll
+++ b/llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll
@@ -5,13 +5,13 @@
; Make sure norecurse is inferred on the calling functions
define internal void @sendmsg_is_norecurse() {
-; FNATTRS: Function Attrs: mustprogress norecurse nounwind willreturn
+; FNATTRS: Function Attrs: mustprogress nofree 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: Function Attrs: mustprogress nofree 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) #[[ATTR4:[0-9]+]]
@@ -22,24 +22,30 @@ define internal void @sendmsg_is_norecurse() {
}
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
+; FNATTRS: Function Attrs: nofree 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: nofree norecurse 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) #[[ATTR5:[0-9]+]]
+; 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: Function Attrs: mustprogress nofree 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 norecurse nounwind willreturn
+; ATTRIBUTOR: Function Attrs: mustprogress nofree norecurse 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]]
@@ -50,7 +56,7 @@ define internal i32 @sendmsg_rtn_is_norecurse() {
}
define void @user() {
-; FNATTRS: Function Attrs: norecurse nounwind
+; FNATTRS: Function Attrs: nofree norecurse nounwind
; FNATTRS-LABEL: define void @user(
; FNATTRS-SAME: ) #[[ATTR1]] {
; FNATTRS-NEXT: call void @sendmsg_is_norecurse()
@@ -58,12 +64,12 @@ define void @user() {
; FNATTRS-NEXT: [[TMP1:%.*]] = call i32 @sendmsg_rtn_is_norecurse()
; FNATTRS-NEXT: ret void
;
-; ATTRIBUTOR: Function Attrs: norecurse nounwind
+; ATTRIBUTOR: Function Attrs: nofree norecurse 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() #[[ATTR6:[0-9]+]]
-; ATTRIBUTOR-NEXT: [[TMP1:%.*]] = call i32 @sendmsg_rtn_is_norecurse() #[[ATTR6]]
+; ATTRIBUTOR-NEXT: call void @sendmsg_is_norecurse() #[[ATTR6:[0-9]+]]
+; ATTRIBUTOR-NEXT: call void @sendmsghalt_is_norecurse() #[[ATTR7:[0-9]+]]
+; ATTRIBUTOR-NEXT: [[TMP1:%.*]] = call i32 @sendmsg_rtn_is_norecurse() #[[ATTR7]]
; ATTRIBUTOR-NEXT: ret void
;
call void @sendmsg_is_norecurse()
@@ -72,16 +78,19 @@ define void @user() {
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]+]] = { nocallback nounwind }
+; FNATTRS: attributes #[[ATTR0]] = { mustprogress nofree norecurse nounwind willreturn }
+; FNATTRS: attributes #[[ATTR1]] = { nofree norecurse nounwind }
+; FNATTRS: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nounwind willreturn }
+; FNATTRS: attributes #[[ATTR3:[0-9]+]] = { nocallback nofree nounwind }
;.
-; ATTRIBUTOR: attributes #[[ATTR0]] = { mustprogress norecurse nounwind willreturn }
-; ATTRIBUTOR: attributes #[[ATTR1]] = { norecurse nounwind }
-; 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 }
+; ATTRIBUTOR: attributes #[[ATTR0]] = { mustprogress nofree norecurse nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR1]] = { nofree norecurse nounwind }
+; ATTRIBUTOR: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR3:[0-9]+]] = { nocallback nofree nounwind }
+; ATTRIBUTOR: attributes #[[ATTR4]] = { nofree willreturn }
+; ATTRIBUTOR: attributes #[[ATTR5]] = { nofree }
+; ATTRIBUTOR: attributes #[[ATTR6]] = { nofree nounwind willreturn }
+; ATTRIBUTOR: attributes #[[ATTR7]] = { nofree nounwind }
;.
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; COMMON: {{.*}}
More information about the cfe-commits
mailing list