Author: Yaxun (Sam) Liu Date: 2023-03-02T11:56:32-05:00 New Revision: 37114036aa57e53217a57afacd7f47b36114edfb
URL: https://github.com/llvm/llvm-project/commit/37114036aa57e53217a57afacd7f47b36114edfb DIFF: https://github.com/llvm/llvm-project/commit/37114036aa57e53217a57afacd7f47b36114edfb.diff LOG: [AMDGPU] Mark mbcnt as convergent since it depends on CFG. Otherwise some passes will try to merge them and cause incorrect results. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D145072 Added: Modified: clang/test/CodeGenOpenCL/builtins-amdgcn.cl llvm/include/llvm/IR/IntrinsicsAMDGPU.td Removed: ################################################################################ diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 094851b218898..ff13357855413 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -676,12 +676,14 @@ kernel void test_gws_sema_p(uint id) { // CHECK-LABEL: @test_mbcnt_lo( // CHECK: call i32 @llvm.amdgcn.mbcnt.lo(i32 %src0, i32 %src1) +// CHECK: declare i32 @llvm.amdgcn.mbcnt.lo(i32, i32) #[[$MBCNT_ATTRS:[0-9]+]] kernel void test_mbcnt_lo(global uint* out, uint src0, uint src1) { *out = __builtin_amdgcn_mbcnt_lo(src0, src1); } // CHECK-LABEL: @test_mbcnt_hi( // CHECK: call i32 @llvm.amdgcn.mbcnt.hi(i32 %src0, i32 %src1) +// CHECK: declare i32 @llvm.amdgcn.mbcnt.hi(i32, i32) #[[$MBCNT_ATTRS]] kernel void test_mbcnt_hi(global uint* out, uint src0, uint src1) { *out = __builtin_amdgcn_mbcnt_hi(src0, src1); } @@ -798,6 +800,7 @@ kernel void test_s_setreg(uint val) { // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025} // CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { mustprogress nocallback nofree nosync nounwind willreturn memory(read) } // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent } +// CHECK-DAG: attributes #[[$MBCNT_ATTRS]] = {{.* convergent .*}} // CHECK-DAG: ![[$EXEC]] = !{!"exec"} // CHECK-DAG: ![[$EXEC_LO]] = !{!"exec_lo"} // CHECK-DAG: ![[$EXEC_HI]] = !{!"exec_hi"} diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index f82682aff7c4d..ba01f383f0ce6 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -1570,12 +1570,12 @@ def int_amdgcn_live_mask : DefaultAttrsIntrinsic <[llvm_i1_ty], def int_amdgcn_mbcnt_lo : ClangBuiltin<"__builtin_amdgcn_mbcnt_lo">, DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem]>; + [IntrNoMem, IntrConvergent]>; def int_amdgcn_mbcnt_hi : ClangBuiltin<"__builtin_amdgcn_mbcnt_hi">, DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem]>; + [IntrNoMem, IntrConvergent]>; // llvm.amdgcn.ds.swizzle src offset def int_amdgcn_ds_swizzle : _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
