llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Shilei Tian (shiltian) <details> <summary>Changes</summary> --- Full diff: https://github.com/llvm/llvm-project/pull/222222.diff 3 Files Affected: - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-async-load-store-lds.cl (+6-6) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+15-13) - (modified) llvm/test/Transforms/FunctionAttrs/sendmsg-nocallback.ll (+34-25) ``````````diff 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 0e40ce71ee794..db2dfc305dc13 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -171,12 +171,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>], [], @@ -268,16 +268,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">, @@ -2250,7 +2250,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 : @@ -3182,7 +3182,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>>] >; @@ -3195,7 +3195,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> : @@ -3210,7 +3210,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 = { @@ -3850,13 +3851,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 : @@ -4272,7 +4274,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>>] >; @@ -4290,7 +4292,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: {{.*}} `````````` </details> https://github.com/llvm/llvm-project/pull/222222 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
