[PATCH] D153953: Revert "[AMDGPU] Mark mbcnt as convergent"
Sameer Sahasrabuddhe via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 28 01:31:49 PDT 2023
sameerds created this revision.
Herald added subscribers: kerbowa, tpr, dstuttard, yaxunl, jvesely, kzhuravl.
Herald added a project: All.
sameerds requested review of this revision.
Herald added subscribers: llvm-commits, cfe-commits, wdng.
Herald added projects: clang, LLVM.
This reverts commit 37114036aa57e53217a57afacd7f47b36114edfb <https://reviews.llvm.org/rG37114036aa57e53217a57afacd7f47b36114edfb>.
The output of mbcnt does not depend on other active lanes, and hence it is not
convergent. The original change was made as a possible fix for
https://github.com/ROCm-Developer-Tools/HIP/issues/3172
But changing mbcnt does not fix that issue.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D153953
Files:
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
Index: llvm/include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -1833,12 +1833,12 @@
def int_amdgcn_mbcnt_lo :
ClangBuiltin<"__builtin_amdgcn_mbcnt_lo">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
- [IntrNoMem, IntrConvergent]>;
+ [IntrNoMem]>;
def int_amdgcn_mbcnt_hi :
ClangBuiltin<"__builtin_amdgcn_mbcnt_hi">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
- [IntrNoMem, IntrConvergent]>;
+ [IntrNoMem]>;
// llvm.amdgcn.ds.swizzle src offset
def int_amdgcn_ds_swizzle :
Index: clang/test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -690,14 +690,12 @@
// 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);
}
@@ -834,7 +832,6 @@
// 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"}
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D153953.535281.patch
Type: text/x-patch
Size: 1993 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230628/6e259376/attachment.bin>
More information about the cfe-commits
mailing list