https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/230516
>From edb27a66044053bca4e2cfd95519ebe66277a8e7 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Fri, 9 Oct 2026 09:28:58 -0400 Subject: [PATCH] [AMDGPU] Make named barrier type 1 byte to fix barrier IDs of array elements Since #209746, a pointer in the barrier address space (15) is the barrier ID, and the backend reads the ID as `ptr & 0x3F`. But `target("amdgcn.named.barrier", 0)` is still 16 bytes, so GEP to element `i` of a barrier array adds `16 * i` to the barrier ID. For example, if `@bars` gets barrier ID 1, `&bars[2]` selects barrier 33 instead of 3. This PR changes the type to 1 byte, so one array element is one barrier ID. Fixes LCOMPILER-2898. --- clang/include/clang/Basic/AMDGPUTypes.def | 2 +- clang/test/CodeGenHIP/amdgpu-barrier-type.hip | 10 +-- .../Generic/amdgpu-barrier-type-debug-info.c | 2 +- clang/test/SemaCXX/amdgpu-barrier.cpp | 4 +- clang/test/SemaHIP/amdgpu-barrier.hip | 4 +- llvm/lib/IR/Type.cpp | 6 +- .../lib/Target/AMDGPU/AMDGPULowerExecSync.cpp | 2 +- llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h | 2 +- .../CodeGen/AMDGPU/s-barrier-array-index.ll | 36 +++----- .../AMDGPU/s-barrier-signal-var-gep.ll | 88 +++---------------- 10 files changed, 39 insertions(+), 117 deletions(-) diff --git a/clang/include/clang/Basic/AMDGPUTypes.def b/clang/include/clang/Basic/AMDGPUTypes.def index 8065393413d5ed..45caa1607f94be 100644 --- a/clang/include/clang/Basic/AMDGPUTypes.def +++ b/clang/include/clang/Basic/AMDGPUTypes.def @@ -28,7 +28,7 @@ AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_buffer_rsrc_t", AMDGPUBufferRsrc, AMDGPUBufferRsrcTy, 128, 128, 8) AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_texture_t", AMDGPUTexture, AMDGPUTextureTy, 256, 256, 0) -AMDGPU_NAMED_BARRIER_TYPE("__amdgpu_named_workgroup_barrier_t", AMDGPUNamedWorkgroupBarrier, AMDGPUNamedWorkgroupBarrierTy, 128, 32, 0) +AMDGPU_NAMED_BARRIER_TYPE("__amdgpu_named_workgroup_barrier_t", AMDGPUNamedWorkgroupBarrier, AMDGPUNamedWorkgroupBarrierTy, 8, 8, 0) AMDGPU_FEATURE_PREDICATE_TYPE("__amdgpu_feature_predicate_t", AMDGPUFeaturePredicate, AMDGPUFeaturePredicateTy, 1, 1) diff --git a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip index f19c91a52a70ce..57af6e53db06d2 100644 --- a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip +++ b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip @@ -8,10 +8,10 @@ __shared__ __amdgpu_named_workgroup_barrier_t bar; __shared__ __amdgpu_named_workgroup_barrier_t bar_arr[2]; //. -// CHECK: @bar = addrspace(15) global target("amdgcn.named.barrier", 0) undef, align 4 -// CHECK: @bar_arr = addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] undef, align 4 -// CHECK: @bar_wrapper_str = addrspace(15) global %struct.WrapperStruct undef, align 4 -// CHECK: @bar_wrapperwrapper_str = addrspace(15) global %struct.WrapperWrapperStruct undef, align 4 +// CHECK: @bar = addrspace(15) global target("amdgcn.named.barrier", 0) undef, align 1 +// CHECK: @bar_arr = addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] undef, align 1 +// CHECK: @bar_wrapper_str = addrspace(15) global %struct.WrapperStruct undef, align 1 +// CHECK: @bar_wrapperwrapper_str = addrspace(15) global %struct.WrapperWrapperStruct undef, align 1 //. __shared__ struct WrapperStruct { __amdgpu_named_workgroup_barrier_t x; @@ -35,7 +35,7 @@ __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(15) @bar to ptr)) #[[ATTR2]] // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(15) @bar_wrapper_str to ptr)) #[[ATTR2]] // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(15) @bar_wrapperwrapper_str to ptr)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(15) @bar_arr to ptr), i64 16)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(15) @bar_arr to ptr), i64 1)) #[[ATTR2]] // CHECK-NEXT: [[CALL:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]] // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[CALL]]) #[[ATTR2]] // CHECK-NEXT: [[CALL1:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]] diff --git a/clang/test/DebugInfo/Generic/amdgpu-barrier-type-debug-info.c b/clang/test/DebugInfo/Generic/amdgpu-barrier-type-debug-info.c index ad5e961c5e6375..c70f11667336f6 100644 --- a/clang/test/DebugInfo/Generic/amdgpu-barrier-type-debug-info.c +++ b/clang/test/DebugInfo/Generic/amdgpu-barrier-type-debug-info.c @@ -2,7 +2,7 @@ // RUN: %clang_cc1 -triple amdgpu -emit-llvm -o - %s -debug-info-kind=limited 2>&1 | FileCheck %s // CHECK: name: "__amdgpu_named_workgroup_barrier_t",{{.*}}baseType: ![[BT:[0-9]+]] -// CHECK: [[BT]] = !DIBasicType(name: "__amdgpu_named_workgroup_barrier_t", size: 128, encoding: DW_ATE_unsigned) +// CHECK: [[BT]] = !DIBasicType(name: "__amdgpu_named_workgroup_barrier_t", size: 8, encoding: DW_ATE_unsigned) void test_locals(void) { __amdgpu_named_workgroup_barrier_t *k0; } diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index 3d43dcd1bba6bb..09bc29cabcf077 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -84,5 +84,5 @@ class TemplatedWrapperImplNoStandardLayout : public CRTPWrapperBase<TemplatedWra int k = 0; }; -static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size"); -static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 4, "wrong alignment"); +static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 1, "wrong size"); +static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 1, "wrong alignment"); diff --git a/clang/test/SemaHIP/amdgpu-barrier.hip b/clang/test/SemaHIP/amdgpu-barrier.hip index a14ca874a40980..1a9afeff7c8e84 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -88,5 +88,5 @@ class TemplatedWrapperImplNoStandardLayout : public CRTPWrapperBase<TemplatedWra }; -static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size"); -static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 4, "wrong alignment"); +static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 1, "wrong size"); +static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 1, "wrong alignment"); diff --git a/llvm/lib/IR/Type.cpp b/llvm/lib/IR/Type.cpp index cd87ab5e20e538..7af638499cd868 100644 --- a/llvm/lib/IR/Type.cpp +++ b/llvm/lib/IR/Type.cpp @@ -1089,10 +1089,8 @@ static TargetTypeInfo getTargetTypeInfo(const TargetExtType *Ty) { // Opaque types in the AMDGPU name space. // NOTE: If the size of the type is changed, it must be also updated in // AMDGPUMemoryUtils.h ! - if (Name == "amdgcn.named.barrier") { - return TargetTypeInfo(FixedVectorType::get(Type::getInt32Ty(C), 4), - TargetExtType::CanBeGlobal); - } + if (Name == "amdgcn.named.barrier") + return TargetTypeInfo(Type::getInt8Ty(C), TargetExtType::CanBeGlobal); if (Name == "amdgpu.stridemark") return TargetTypeInfo(Type::getVoidTy(C), TargetExtType::IsTokenLike); diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp index 1fad3489cf1a7f..48395bd0b43672 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp @@ -127,7 +127,7 @@ static bool lowerExecSyncGlobalVariables(Module &M, GVUsesInfoTy &GVUsesInfo) { unsigned Offset; if (TargetExtType *ExtTy = isNamedBarrier(*GV)) { unsigned BarrierScope = ExtTy->getIntParameter(0); - unsigned BarCnt = GV->getGlobalSize(DL) / 16; + unsigned BarCnt = getNumNamedBarriersDeclared(DL, *GV); unsigned BarID = allocateExecSyncID(KernelBarrierIDs, Kernels, BarrierScope, NumBarScopes, BarCnt); diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h b/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h index 7755011bd60c06..95f8c91a4dd939 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h +++ b/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h @@ -31,7 +31,7 @@ class TargetExtType; namespace AMDGPU { -static constexpr unsigned NamedBarrierTypeSizeInBytes = 16; +static constexpr unsigned NamedBarrierTypeSizeInBytes = 1; using FunctionVariableMap = MapVector<Function *, DenseSet<GlobalVariable *>>; using VariableFunctionMap = DenseMap<GlobalVariable *, DenseSet<Function *>>; diff --git a/llvm/test/CodeGen/AMDGPU/s-barrier-array-index.ll b/llvm/test/CodeGen/AMDGPU/s-barrier-array-index.ll index d48d31c1c9b2c8..65775e7ca39c10 100644 --- a/llvm/test/CodeGen/AMDGPU/s-barrier-array-index.ll +++ b/llvm/test/CodeGen/AMDGPU/s-barrier-array-index.ll @@ -13,7 +13,7 @@ define amdgpu_kernel void @signal_var_const_idx() { ; SDAG-NEXT: s_mov_b64 s[64:65], 0 ; SDAG-NEXT: v_nop ; SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; SDAG-NEXT: s_mov_b32 m0, 33 +; SDAG-NEXT: s_mov_b32 m0, 3 ; SDAG-NEXT: s_barrier_signal m0 ; SDAG-NEXT: s_endpgm ; @@ -23,7 +23,7 @@ define amdgpu_kernel void @signal_var_const_idx() { ; GISEL-NEXT: s_mov_b64 s[64:65], 0 ; GISEL-NEXT: v_nop ; GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; GISEL-NEXT: s_barrier_signal 33 +; GISEL-NEXT: s_barrier_signal 3 ; GISEL-NEXT: s_endpgm %bar = getelementptr [4 x target("amdgcn.named.barrier", 0)], ptr addrspace(15) @bars, i32 0, i32 2 call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %bar, i32 0) @@ -39,11 +39,10 @@ define amdgpu_kernel void @init_dynamic_idx(i32 %idx, i32 %cnt) { ; SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x0 nv ; SDAG-NEXT: s_wait_kmcnt 0x0 -; SDAG-NEXT: s_lshl_b32 s0, s0, 4 +; SDAG-NEXT: s_add_co_i32 s0, s0, 1 ; SDAG-NEXT: s_and_b32 s1, s1, 63 -; SDAG-NEXT: s_or_b32 s0, s0, 1 +; SDAG-NEXT: s_and_b32 s0, s0, 63 ; SDAG-NEXT: s_lshl_b32 s1, s1, 16 -; SDAG-NEXT: s_and_b32 s0, s0, 49 ; SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; SDAG-NEXT: s_or_b32 m0, s1, s0 ; SDAG-NEXT: s_barrier_init m0 @@ -57,11 +56,10 @@ define amdgpu_kernel void @init_dynamic_idx(i32 %idx, i32 %cnt) { ; GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; GISEL-NEXT: s_load_b64 s[0:1], s[4:5], 0x0 nv ; GISEL-NEXT: s_wait_kmcnt 0x0 -; GISEL-NEXT: s_lshl_b32 s0, s0, 4 -; GISEL-NEXT: s_and_b32 s1, s1, 63 ; GISEL-NEXT: s_add_co_u32 s0, 1, s0 -; GISEL-NEXT: s_lshl_b32 s1, s1, 16 +; GISEL-NEXT: s_and_b32 s1, s1, 63 ; GISEL-NEXT: s_and_b32 s0, s0, 63 +; GISEL-NEXT: s_lshl_b32 s1, s1, 16 ; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; GISEL-NEXT: s_or_b32 m0, s0, s1 ; GISEL-NEXT: s_barrier_init m0 @@ -80,10 +78,9 @@ define amdgpu_kernel void @join_dynamic_idx(i32 %idx) { ; SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; SDAG-NEXT: s_load_b32 s0, s[4:5], 0x0 nv ; SDAG-NEXT: s_wait_kmcnt 0x0 -; SDAG-NEXT: s_lshl_b32 s0, s0, 4 -; SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; SDAG-NEXT: s_or_b32 s0, s0, 1 -; SDAG-NEXT: s_and_b32 m0, s0, 49 +; SDAG-NEXT: s_add_co_i32 s0, s0, 1 +; SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; SDAG-NEXT: s_and_b32 m0, s0, 63 ; SDAG-NEXT: s_barrier_join m0 ; SDAG-NEXT: s_endpgm ; @@ -95,9 +92,8 @@ define amdgpu_kernel void @join_dynamic_idx(i32 %idx) { ; GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; GISEL-NEXT: s_load_b32 s0, s[4:5], 0x0 nv ; GISEL-NEXT: s_wait_kmcnt 0x0 -; GISEL-NEXT: s_lshl_b32 s0, s0, 4 -; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; GISEL-NEXT: s_add_co_u32 s0, 1, s0 +; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; GISEL-NEXT: s_and_b32 m0, s0, 63 ; GISEL-NEXT: s_barrier_join m0 ; GISEL-NEXT: s_endpgm @@ -115,13 +111,11 @@ define amdgpu_kernel void @get_state_dynamic_idx(ptr addrspace(1) %out, i32 %idx ; SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x0 nv ; SDAG-NEXT: s_wait_kmcnt 0x0 -; SDAG-NEXT: s_lshl_b32 s2, s2, 4 -; SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; SDAG-NEXT: s_or_b32 s2, s2, 1 -; SDAG-NEXT: s_and_b32 m0, s2, 49 +; SDAG-NEXT: s_add_co_i32 s2, s2, 1 +; SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_2) +; SDAG-NEXT: s_and_b32 m0, s2, 63 ; SDAG-NEXT: s_get_barrier_state s2, m0 ; SDAG-NEXT: s_wait_kmcnt 0x0 -; SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_2) ; SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2 ; SDAG-NEXT: global_store_b32 v0, v1, s[0:1] ; SDAG-NEXT: s_endpgm @@ -135,13 +129,11 @@ define amdgpu_kernel void @get_state_dynamic_idx(ptr addrspace(1) %out, i32 %idx ; GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x0 nv ; GISEL-NEXT: v_mov_b32_e32 v1, 0 ; GISEL-NEXT: s_wait_kmcnt 0x0 -; GISEL-NEXT: s_lshl_b32 s2, s2, 4 -; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; GISEL-NEXT: s_add_co_u32 s2, 1, s2 +; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_2) ; GISEL-NEXT: s_and_b32 m0, s2, 63 ; GISEL-NEXT: s_get_barrier_state s2, m0 ; GISEL-NEXT: s_wait_kmcnt 0x0 -; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_2) ; GISEL-NEXT: v_mov_b32_e32 v0, s2 ; GISEL-NEXT: global_store_b32 v1, v0, s[0:1] ; GISEL-NEXT: s_endpgm diff --git a/llvm/test/CodeGen/AMDGPU/s-barrier-signal-var-gep.ll b/llvm/test/CodeGen/AMDGPU/s-barrier-signal-var-gep.ll index 337b6c053e931f..f11ccaba5c9882 100644 --- a/llvm/test/CodeGen/AMDGPU/s-barrier-signal-var-gep.ll +++ b/llvm/test/CodeGen/AMDGPU/s-barrier-signal-var-gep.ll @@ -10,7 +10,7 @@ ; compile-time-constant barrier address and should select the immediate ; s_barrier_signal form, like the bare &bars[0]. With object linking the ; barrier is externalized and the offset folds into the relocation addend -; (...@abs32@lo+16) instead of a runtime add. +; (...@abs32@lo+1) instead of a runtime add. define amdgpu_kernel void @signal_var_bar0() { ; CHECK-LABEL: signal_var_bar0: @@ -31,7 +31,7 @@ define amdgpu_kernel void @signal_var_bar0() { ; CHECK-OBJ-SDAG-NEXT: s_mov_b64 s[64:65], 0 ; CHECK-OBJ-SDAG-NEXT: v_nop ; CHECK-OBJ-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; CHECK-OBJ-SDAG-NEXT: s_and_b32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo, 63 +; CHECK-OBJ-SDAG-NEXT: s_and_b32 s0, __amdgpu_named_barrier.bars.7672b1ccf82b48a897c14e26b2081787@abs32@lo, 63 ; CHECK-OBJ-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; CHECK-OBJ-SDAG-NEXT: s_or_b32 m0, s0, 0x100000 ; CHECK-OBJ-SDAG-NEXT: s_barrier_init m0 @@ -76,83 +76,19 @@ define amdgpu_kernel void @signal_var_bar1() { ; CHECK-GISEL-NEXT: s_mov_b64 s[64:65], 0 ; CHECK-GISEL-NEXT: v_nop ; CHECK-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; CHECK-GISEL-NEXT: s_mov_b32 m0, 0x10000f -; CHECK-GISEL-NEXT: s_barrier_init m0 -; CHECK-GISEL-NEXT: s_barrier_signal 15 -; CHECK-GISEL-NEXT: s_barrier_wait 1 -; CHECK-GISEL-NEXT: s_endpgm -; -; CHECK-OBJ-SDAG-LABEL: signal_var_bar1: -; CHECK-OBJ-SDAG: ; %bb.0: -; CHECK-OBJ-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; CHECK-OBJ-SDAG-NEXT: s_mov_b64 s[64:65], 0 -; CHECK-OBJ-SDAG-NEXT: v_nop -; CHECK-OBJ-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; CHECK-OBJ-SDAG-NEXT: s_add_co_i32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo, 16 -; CHECK-OBJ-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; CHECK-OBJ-SDAG-NEXT: s_and_b32 s0, s0, 63 -; CHECK-OBJ-SDAG-NEXT: s_or_b32 m0, s0, 0x100000 -; CHECK-OBJ-SDAG-NEXT: s_barrier_init m0 -; CHECK-OBJ-SDAG-NEXT: s_mov_b32 m0, s0 -; CHECK-OBJ-SDAG-NEXT: s_barrier_signal m0 -; CHECK-OBJ-SDAG-NEXT: s_barrier_wait 1 -; CHECK-OBJ-SDAG-NEXT: s_endpgm -; -; CHECK-OBJ-GISEL-LABEL: signal_var_bar1: -; CHECK-OBJ-GISEL: ; %bb.0: -; CHECK-OBJ-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; CHECK-OBJ-GISEL-NEXT: s_mov_b64 s[64:65], 0 -; CHECK-OBJ-GISEL-NEXT: v_nop -; CHECK-OBJ-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; CHECK-OBJ-GISEL-NEXT: s_mov_b32 m0, 0x10000f -; CHECK-OBJ-GISEL-NEXT: s_barrier_init m0 -; CHECK-OBJ-GISEL-NEXT: s_barrier_signal 15 -; CHECK-OBJ-GISEL-NEXT: s_barrier_wait 1 -; CHECK-OBJ-GISEL-NEXT: s_endpgm - %p1 = getelementptr inbounds [2 x target("amdgcn.named.barrier", 0)], ptr addrspace(15) @bars, i32 0, i32 1 - call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) %p1, i32 16) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %p1, i32 0) - call void @llvm.amdgcn.s.barrier.wait(i16 1) - ret void -} - -; A byte offset that does not land on a named-barrier boundary is valid IR and -; must not crash the compiler. The offset folds into the address; a sub-object -; offset stays within barrier 0 and selects the same barrier as &bars[0]. - -define amdgpu_kernel void @signal_var_misaligned() { -; CHECK-SDAG-LABEL: signal_var_misaligned: -; CHECK-SDAG: ; %bb.0: -; CHECK-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; CHECK-SDAG-NEXT: s_mov_b64 s[64:65], 0 -; CHECK-SDAG-NEXT: v_nop -; CHECK-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; CHECK-SDAG-NEXT: s_mov_b32 m0, 0x100000 -; CHECK-SDAG-NEXT: s_barrier_init m0 -; CHECK-SDAG-NEXT: s_mov_b32 m0, 0 -; CHECK-SDAG-NEXT: s_barrier_signal m0 -; CHECK-SDAG-NEXT: s_barrier_wait 1 -; CHECK-SDAG-NEXT: s_endpgm -; -; CHECK-GISEL-LABEL: signal_var_misaligned: -; CHECK-GISEL: ; %bb.0: -; CHECK-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; CHECK-GISEL-NEXT: s_mov_b64 s[64:65], 0 -; CHECK-GISEL-NEXT: v_nop -; CHECK-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; CHECK-GISEL-NEXT: s_mov_b32 m0, 0x100000 ; CHECK-GISEL-NEXT: s_barrier_init m0 ; CHECK-GISEL-NEXT: s_barrier_signal 0 ; CHECK-GISEL-NEXT: s_barrier_wait 1 ; CHECK-GISEL-NEXT: s_endpgm ; -; CHECK-OBJ-SDAG-LABEL: signal_var_misaligned: +; CHECK-OBJ-SDAG-LABEL: signal_var_bar1: ; CHECK-OBJ-SDAG: ; %bb.0: ; CHECK-OBJ-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; CHECK-OBJ-SDAG-NEXT: s_mov_b64 s[64:65], 0 ; CHECK-OBJ-SDAG-NEXT: v_nop ; CHECK-OBJ-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE -; CHECK-OBJ-SDAG-NEXT: s_add_co_i32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo, 1 +; CHECK-OBJ-SDAG-NEXT: s_add_co_i32 s0, __amdgpu_named_barrier.bars.7672b1ccf82b48a897c14e26b2081787@abs32@lo, 1 ; CHECK-OBJ-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; CHECK-OBJ-SDAG-NEXT: s_and_b32 s0, s0, 63 ; CHECK-OBJ-SDAG-NEXT: s_or_b32 m0, s0, 0x100000 @@ -162,7 +98,7 @@ define amdgpu_kernel void @signal_var_misaligned() { ; CHECK-OBJ-SDAG-NEXT: s_barrier_wait 1 ; CHECK-OBJ-SDAG-NEXT: s_endpgm ; -; CHECK-OBJ-GISEL-LABEL: signal_var_misaligned: +; CHECK-OBJ-GISEL-LABEL: signal_var_bar1: ; CHECK-OBJ-GISEL: ; %bb.0: ; CHECK-OBJ-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; CHECK-OBJ-GISEL-NEXT: s_mov_b64 s[64:65], 0 @@ -173,7 +109,7 @@ define amdgpu_kernel void @signal_var_misaligned() { ; CHECK-OBJ-GISEL-NEXT: s_barrier_signal 0 ; CHECK-OBJ-GISEL-NEXT: s_barrier_wait 1 ; CHECK-OBJ-GISEL-NEXT: s_endpgm - %p1 = getelementptr i8, ptr addrspace(15) @bars, i32 1 + %p1 = getelementptr inbounds [2 x target("amdgcn.named.barrier", 0)], ptr addrspace(15) @bars, i32 0, i32 1 call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) %p1, i32 16) call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %p1, i32 0) call void @llvm.amdgcn.s.barrier.wait(i16 1) @@ -193,7 +129,7 @@ define amdgpu_kernel void @signal_var_dynamic(i32 %idx) { ; CHECK-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; CHECK-SDAG-NEXT: s_load_b32 s0, s[4:5], 0x0 nv ; CHECK-SDAG-NEXT: s_wait_kmcnt 0x0 -; CHECK-SDAG-NEXT: s_lshl4_add_u32 s0, s0, -1 +; CHECK-SDAG-NEXT: s_add_co_i32 s0, s0, -1 ; CHECK-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; CHECK-SDAG-NEXT: s_and_b32 s0, s0, 63 ; CHECK-SDAG-NEXT: s_or_b32 m0, s0, 0x100000 @@ -211,11 +147,9 @@ define amdgpu_kernel void @signal_var_dynamic(i32 %idx) { ; CHECK-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; CHECK-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x0 nv ; CHECK-GISEL-NEXT: s_wait_kmcnt 0x0 -; CHECK-GISEL-NEXT: s_lshl_b32 s0, s0, 4 -; CHECK-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; CHECK-GISEL-NEXT: s_add_co_u32 s0, -1, s0 +; CHECK-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; CHECK-GISEL-NEXT: s_and_b32 s0, s0, 63 -; CHECK-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; CHECK-GISEL-NEXT: s_or_b32 m0, s0, 0x100000 ; CHECK-GISEL-NEXT: s_barrier_init m0 ; CHECK-GISEL-NEXT: s_mov_b32 m0, s0 @@ -231,7 +165,7 @@ define amdgpu_kernel void @signal_var_dynamic(i32 %idx) { ; CHECK-OBJ-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; CHECK-OBJ-SDAG-NEXT: s_load_b32 s0, s[4:5], 0x0 nv ; CHECK-OBJ-SDAG-NEXT: s_wait_kmcnt 0x0 -; CHECK-OBJ-SDAG-NEXT: s_lshl4_add_u32 s0, s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo +; CHECK-OBJ-SDAG-NEXT: s_add_co_i32 s0, s0, __amdgpu_named_barrier.bars.7672b1ccf82b48a897c14e26b2081787@abs32@lo ; CHECK-OBJ-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; CHECK-OBJ-SDAG-NEXT: s_and_b32 s0, s0, 63 ; CHECK-OBJ-SDAG-NEXT: s_or_b32 m0, s0, 0x100000 @@ -249,11 +183,9 @@ define amdgpu_kernel void @signal_var_dynamic(i32 %idx) { ; CHECK-OBJ-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE ; CHECK-OBJ-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x0 nv ; CHECK-OBJ-GISEL-NEXT: s_wait_kmcnt 0x0 -; CHECK-OBJ-GISEL-NEXT: s_lshl_b32 s0, s0, 4 -; CHECK-OBJ-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; CHECK-OBJ-GISEL-NEXT: s_add_co_u32 s0, -1, s0 +; CHECK-OBJ-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; CHECK-OBJ-GISEL-NEXT: s_and_b32 s0, s0, 63 -; CHECK-OBJ-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; CHECK-OBJ-GISEL-NEXT: s_or_b32 m0, s0, 0x100000 ; CHECK-OBJ-GISEL-NEXT: s_barrier_init m0 ; CHECK-OBJ-GISEL-NEXT: s_mov_b32 m0, s0 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
