This is an archive of the discontinued LLVM Phabricator instance.

Differential D145072

[AMDGPU] Mark mbcnt as convergent
ClosedPublic

Authored by yaxunl on Mar 1 2023, 8:00 AM.

Details

Reviewers
arsenm
rampitec
ruiling
tra
Commits
rG37114036aa57: [AMDGPU] Mark mbcnt as convergent
Summary

since it depends on CFG.

Otherwise some passes will try to merge them and cause
incorrect results.

Diff Detail

Event Timeline

yaxunl created this revision.Mar 1 2023, 8:00 AM
Herald added a project: Restricted Project. · View Herald TranscriptMar 1 2023, 8:00 AM
Herald added subscribers: kosarev, kerbowa, tpr and 3 others. · View Herald Transcript
yaxunl requested review of this revision.Mar 1 2023, 8:00 AM
Herald added a project: Restricted Project. · View Herald TranscriptMar 1 2023, 8:00 AM
Herald added a subscriber: wdng. · View Herald Transcript
Harbormaster completed remote builds in B216723: Diff 501517.Mar 1 2023, 8:51 AM
tra accepted this revision.Mar 1 2023, 11:03 AM
This revision is now accepted and ready to land.Mar 1 2023, 11:03 AM
This revision was landed with ongoing or failed builds.Mar 2 2023, 8:57 AM
Closed by commit rG37114036aa57: [AMDGPU] Mark mbcnt as convergent (authored by yaxunl). · Explain Why
This revision was automatically updated to reflect the committed changes.
yaxunl added a commit: rG37114036aa57: [AMDGPU] Mark mbcnt as convergent.
Herald added a project: Restricted Project. · View Herald TranscriptMar 2 2023, 8:57 AM
Herald added a subscriber: cfe-commits. · View Herald Transcript
sameerds added a reverting change: D153953: Revert "[AMDGPU] Mark mbcnt as convergent".Jun 28 2023, 1:31 AM
sameerds added a reverting change: rG7a101798b753: Revert "[AMDGPU] Mark mbcnt as convergent".Jun 30 2023, 12:54 AM

Revision Contents

PathSize
clang/
test/
CodeGenOpenCL/
3 lines
llvm/
include/
llvm/
IR/
4 lines

Diff 501886

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

Show First 20 Lines • Show All 670 Lines • ▼ Show 20 Lines
// CHECK-LABEL: @test_gws_sema_p( // CHECK-LABEL: @test_gws_sema_p(
// CHECK: call void @llvm.amdgcn.ds.gws.sema.p(i32 %id) // CHECK: call void @llvm.amdgcn.ds.gws.sema.p(i32 %id)
kernel void test_gws_sema_p(uint id) { kernel void test_gws_sema_p(uint id) {
__builtin_amdgcn_ds_gws_sema_p(id); __builtin_amdgcn_ds_gws_sema_p(id);
} }
// CHECK-LABEL: @test_mbcnt_lo( // CHECK-LABEL: @test_mbcnt_lo(
// CHECK: call i32 @llvm.amdgcn.mbcnt.lo(i32 %src0, i32 %src1) // 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) { kernel void test_mbcnt_lo(global uint* out, uint src0, uint src1) {
*out = __builtin_amdgcn_mbcnt_lo(src0, src1); *out = __builtin_amdgcn_mbcnt_lo(src0, src1);
} }
// CHECK-LABEL: @test_mbcnt_hi( // CHECK-LABEL: @test_mbcnt_hi(
// CHECK: call i32 @llvm.amdgcn.mbcnt.hi(i32 %src0, i32 %src1) // 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) { kernel void test_mbcnt_hi(global uint* out, uint src0, uint src1) {
*out = __builtin_amdgcn_mbcnt_hi(src0, src1); *out = __builtin_amdgcn_mbcnt_hi(src0, src1);
} }
// CHECK-LABEL: @test_alignbit( // CHECK-LABEL: @test_alignbit(
// CHECK: tail call i32 @llvm.fshr.i32(i32 %src0, i32 %src1, i32 %src2) // CHECK: tail call i32 @llvm.fshr.i32(i32 %src0, i32 %src1, i32 %src2)
kernel void test_alignbit(global uint* out, uint src0, uint src1, uint src2) { kernel void test_alignbit(global uint* out, uint src0, uint src1, uint src2) {
*out = __builtin_amdgcn_alignbit(src0, src1, src2); *out = __builtin_amdgcn_alignbit(src0, src1, src2);
▲ Show 20 Lines • Show All 100 Lines • ▼ Show 20 Lines
kernel void test_s_setreg(uint val) { kernel void test_s_setreg(uint val) {
__builtin_amdgcn_s_setreg(8193, val); __builtin_amdgcn_s_setreg(8193, val);
} }
// CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024} // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
// CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025} // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}
// CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { mustprogress nocallback nofree nosync nounwind willreturn memory(read) } // CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { mustprogress nocallback nofree nosync nounwind willreturn memory(read) }
// CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent } // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent }
// CHECK-DAG: attributes #[[$MBCNT_ATTRS]] = {{.* convergent .*}}
// CHECK-DAG: ![[$EXEC]] = !{!"exec"} // CHECK-DAG: ![[$EXEC]] = !{!"exec"}
// CHECK-DAG: ![[$EXEC_LO]] = !{!"exec_lo"} // CHECK-DAG: ![[$EXEC_LO]] = !{!"exec_lo"}
// CHECK-DAG: ![[$EXEC_HI]] = !{!"exec_hi"} // CHECK-DAG: ![[$EXEC_HI]] = !{!"exec_hi"}

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 1,564 Lines • ▼ Show 20 Lines
// Returns true if lane is live (and not a helper lane). // Returns true if lane is live (and not a helper lane).
def int_amdgcn_live_mask : DefaultAttrsIntrinsic <[llvm_i1_ty], def int_amdgcn_live_mask : DefaultAttrsIntrinsic <[llvm_i1_ty],
[], [IntrReadMem, IntrInaccessibleMemOnly] [], [IntrReadMem, IntrInaccessibleMemOnly]
>; >;
def int_amdgcn_mbcnt_lo : def int_amdgcn_mbcnt_lo :
ClangBuiltin<"__builtin_amdgcn_mbcnt_lo">, ClangBuiltin<"__builtin_amdgcn_mbcnt_lo">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem]>; [IntrNoMem, IntrConvergent]>;
def int_amdgcn_mbcnt_hi : def int_amdgcn_mbcnt_hi :
ClangBuiltin<"__builtin_amdgcn_mbcnt_hi">, ClangBuiltin<"__builtin_amdgcn_mbcnt_hi">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem]>; [IntrNoMem, IntrConvergent]>;
// llvm.amdgcn.ds.swizzle src offset // llvm.amdgcn.ds.swizzle src offset
def int_amdgcn_ds_swizzle : def int_amdgcn_ds_swizzle :
ClangBuiltin<"__builtin_amdgcn_ds_swizzle">, ClangBuiltin<"__builtin_amdgcn_ds_swizzle">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree,
ImmArg<ArgIndex<1>>]>; ImmArg<ArgIndex<1>>]>;
▲ Show 20 Lines • Show All 888 Lines • Show Last 20 Lines