diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 285153695da27..72be557a7af60 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -56,7 +56,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUAddrSpaceMap = { {LangAS::hlsl_input, llvm::AMDGPUAS::PRIVATE_ADDRESS}, {LangAS::hlsl_output, llvm::AMDGPUAS::PRIVATE_ADDRESS}, {LangAS::hlsl_push_constant, llvm::AMDGPUAS::GLOBAL_ADDRESS}, - {LangAS::amdgpu_barrier, llvm::AMDGPUAS::LOCAL_ADDRESS}, + {LangAS::amdgpu_barrier, llvm::AMDGPUAS::BARRIER}, }; } // namespace targets diff --git a/clang/test/CodeGen/target-data.c b/clang/test/CodeGen/target-data.c index e74454ce95f81..3340176bb54c0 100644 --- a/clang/test/CodeGen/target-data.c +++ b/clang/test/CodeGen/target-data.c @@ -160,12 +160,12 @@ // RUN: %clang_cc1 -triple amdgcn-unknown -target-cpu hawaii -o - -emit-llvm %s \ // RUN: | FileCheck %s -check-prefix=R600SI -// R600SI: target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9" +// R600SI: target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-p15:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9" // Test default -target-cpu // RUN: %clang_cc1 -triple amdgcn-unknown -o - -emit-llvm %s \ // RUN: | FileCheck %s -check-prefix=R600SIDefault -// R600SIDefault: target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9" +// R600SIDefault: target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-p15:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9" // RUN: %clang_cc1 -triple arm64-unknown -o - -emit-llvm %s | \ // RUN: FileCheck %s -check-prefix=AARCH64 diff --git a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip index df9e3631c0d1f..cb450044a2f30 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(3) global target("amdgcn.named.barrier", 0) undef, align 4 -// CHECK: @bar_arr = addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] undef, align 4 -// CHECK: @bar_wrapper_str = addrspace(3) global %struct.WrapperStruct undef, align 4 -// CHECK: @bar_wrapperwrapper_str = addrspace(3) global %struct.WrapperWrapperStruct undef, align 4 +// 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 //. __shared__ struct WrapperStruct { __amdgpu_named_workgroup_barrier_t x; @@ -32,10 +32,10 @@ __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); // CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8 // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[TMP0]]) #[[ATTR2:[0-9]+]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @bar to ptr)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @bar_wrapper_str to ptr)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @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(3) @bar_arr to ptr), i64 16)) #[[ATTR2]] +// 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:%.*]] = 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/CodeGenOpenCL/amdgpu-env-amdgcn.cl b/clang/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl index 72ce72644b8ea..fcee9b3b20813 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl @@ -1,5 +1,5 @@ // RUN: %clang_cc1 %s -O0 -triple amdgcn -emit-llvm -o - | FileCheck %s // RUN: %clang_cc1 %s -O0 -triple amdgcn---opencl -emit-llvm -o - | FileCheck %s -// CHECK: target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9" +// CHECK: target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-p15:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9" void foo(void) {} diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl index 332a2fa94ee92..8ea4d4e2b32e2 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl @@ -83,9 +83,9 @@ void test_s_barrier_signal() // CHECK-NEXT: store ptr [[BAR:%.*]], ptr addrspace(5) [[BAR_ADDR]], align 8 // CHECK-NEXT: store i32 [[A:%.*]], ptr addrspace(5) [[A_ADDR]], align 4 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[BAR_ADDR]], align 8 -// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3) +// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(15) // CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[A_ADDR]], align 4 -// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) [[TMP1]], i32 [[TMP2]]) +// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) [[TMP1]], i32 [[TMP2]]) // CHECK-NEXT: ret void // void test_s_barrier_signal_var(void *bar, int a) @@ -132,9 +132,9 @@ void test_s_barrier_signal_isfirst(int* a, int* b, int *c) // CHECK-NEXT: store ptr [[BAR:%.*]], ptr addrspace(5) [[BAR_ADDR]], align 8 // CHECK-NEXT: store i32 [[A:%.*]], ptr addrspace(5) [[A_ADDR]], align 4 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[BAR_ADDR]], align 8 -// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3) +// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(15) // CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[A_ADDR]], align 4 -// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) [[TMP1]], i32 [[TMP2]]) +// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) [[TMP1]], i32 [[TMP2]]) // CHECK-NEXT: ret void // void test_s_barrier_init(void *bar, int a) @@ -147,8 +147,8 @@ void test_s_barrier_init(void *bar, int a) // CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) // CHECK-NEXT: store ptr [[BAR:%.*]], ptr addrspace(5) [[BAR_ADDR]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[BAR_ADDR]], align 8 -// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3) -// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) [[TMP1]]) +// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(15) +// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) [[TMP1]]) // CHECK-NEXT: ret void // void test_s_barrier_join(void *bar) @@ -189,8 +189,8 @@ unsigned test_s_get_barrier_state(int a) // CHECK-NEXT: [[STATE:%.*]] = alloca i32, align 4, addrspace(5) // CHECK-NEXT: store ptr [[BAR:%.*]], ptr addrspace(5) [[BAR_ADDR]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[BAR_ADDR]], align 8 -// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3) -// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) [[TMP1]]) +// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(15) +// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) [[TMP1]]) // CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[STATE]], align 4 // CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[STATE]], align 4 // CHECK-NEXT: ret i32 [[TMP3]] diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl index 8b09216057167..9368c2971a643 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl @@ -1362,8 +1362,8 @@ void test_s_cluster_barrier() // CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) // CHECK-NEXT: store ptr [[BAR:%.*]], ptr addrspace(5) [[BAR_ADDR]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[BAR_ADDR]], align 8 -// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3) -// CHECK-NEXT: call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) [[TMP1]]) +// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(15) +// CHECK-NEXT: call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(15) [[TMP1]]) // CHECK-NEXT: ret void // void test_s_wakeup_barrier(void *bar) diff --git a/lld/test/ELF/lto/amdgcn-oses.ll b/lld/test/ELF/lto/amdgcn-oses.ll index 2ff60bfc66bdd..274efc6929104 100644 --- a/lld/test/ELF/lto/amdgcn-oses.ll +++ b/lld/test/ELF/lto/amdgcn-oses.ll @@ -25,7 +25,7 @@ ;--- amdhsa.ll target triple = "amdgpu7.00-amd-amdhsa" -target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" +target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p15:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" !llvm.module.flags = !{!0} !0 = !{i32 1, !"amdhsa_code_object_version", i32 500} @@ -36,7 +36,7 @@ define void @_start() { ;--- amdpal.ll target triple = "amdgpu7.00-amd-amdpal" -target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" +target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p15:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" define amdgpu_cs void @_start() { ret void @@ -44,7 +44,7 @@ define amdgpu_cs void @_start() { ;--- mesa3d.ll target triple = "amdgpu7.00-amd-mesa3d" -target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" +target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p15:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" define void @_start() { ret void diff --git a/lld/test/ELF/lto/amdgcn.ll b/lld/test/ELF/lto/amdgcn.ll index 186185c44a2c2..1dc2d86c48364 100644 --- a/lld/test/ELF/lto/amdgcn.ll +++ b/lld/test/ELF/lto/amdgcn.ll @@ -5,7 +5,7 @@ ; Make sure the amdgcn triple is handled target triple = "amdgcn-amd-amdhsa" -target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" +target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p15:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" define void @_start() { ret void diff --git a/lld/test/ELF/lto/amdgpu.ll b/lld/test/ELF/lto/amdgpu.ll index c24ce3831df19..7e023046442b0 100644 --- a/lld/test/ELF/lto/amdgpu.ll +++ b/lld/test/ELF/lto/amdgpu.ll @@ -5,7 +5,7 @@ ; Make sure the amdgpu triple is handled target triple = "amdgpu7.00-amd-amdhsa" -target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" +target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p15:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5" define void @_start() { ret void diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst index b3b2a06005377..8e9199bc17bcc 100644 --- a/llvm/docs/AMDGPUUsage.rst +++ b/llvm/docs/AMDGPUUsage.rst @@ -1158,6 +1158,7 @@ supported for the ``amdgcn`` target. *reserved for downstream use (LLPC)* 12 *reserved for future use* 13 *reserved for future use* 14 + Barrier 15 N/A N/A 32 0 *reserved for future use* 16 Streamout Registers 128 N/A GS_REGS ===================================== =============== =========== ================ ======= ============================ @@ -1371,6 +1372,26 @@ supported for the ``amdgcn`` target. a buffer strided pointer, this means that the base pointer is ``align(4)``, that the offset is a multiple of 4 bytes, and that the stride is a multiple of 4. +**Barrier** + This address space represents barrier IDs (introduced in GFX12) as addresses. + It does not map directly to any addressable memory, thus pointers into this address space: + + * Never alias with any other pointers outside this address space. + * Cannot be dereferenced. + * Can only be consumed by intrinsics. + + Pointer are 32 bits and directly correspond to valid barrier IDs. When consumed by an + intrinsic, all barrier pointers must, when interpreted as signed 32 bit integers, + have a value corresponding to a valid barrier ID on the target. + Otherwise, the behavior is undefined. + + The ``NULL`` pointer (as a constant) can be consumed by some intrinsics and + corresponds to the NULL named barrier. + + These pointers do not have a corresponding hardware aperture but safe round-tripping + through the generic address space is still possible. Attempting to dereference a + generic pointer derived from a barrier pointer is undefined behavior. + **Streamout Registers** Dedicated registers used by the GS NGG Streamout Instructions. The register file is modelled as a memory in a distinct address space because it is indexed @@ -1577,10 +1598,8 @@ Named barriers are fixed function hardware barrier objects that are available in gfx12.5+ in addition to the traditional default barriers. In LLVM IR, named barriers are represented by global variables of type -``target("amdgcn.named.barrier", 0)`` in the LDS address space. Named barrier -global variables do not occupy actual LDS memory, but their lifetime and -allocation scope matches that of global variables in LDS. Programs in LLVM IR -refer to named barriers using pointers. +``target("amdgcn.named.barrier", 0)`` in the barrier address space. +Programs in LLVM IR refer to named barriers using pointers. The following named barrier types are supported in global variables, defined recursively: @@ -1591,14 +1610,14 @@ recursively: .. code-block:: llvm - @bar = addrspace(3) global target("amdgcn.named.barrier", 0) undef - @foo = addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] undef - @baz = addrspace(3) global { target("amdgcn.named.barrier", 0) } undef + @bar = addrspace(15) global target("amdgcn.named.barrier", 0) undef + @foo = addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] undef + @baz = addrspace(15) global { target("amdgcn.named.barrier", 0) } undef ... - %foo.i = getelementptr [2 x target("amdgcn.named.barrier", 0)], ptr addrspace(3) @foo, i32 0, i32 %i - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) %foo.i, i32 0) + %foo.i = getelementptr [2 x target("amdgcn.named.barrier", 0)], ptr addrspace(15) @foo, i32 0, i32 %i + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %foo.i, i32 0) Named barrier types may not be used in ``alloca``. diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 565637b36131c..f2a8c4103a1f0 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -13,6 +13,7 @@ def flat_ptr_ty : LLVMQualPointerType<0>; def global_ptr_ty : LLVMQualPointerType<1>; def local_ptr_ty : LLVMQualPointerType<3>; +def barrier_ptr_ty : LLVMQualPointerType<15>; // The amdgpu-no-* attributes (ex amdgpu-no-workitem-id-z) typically inferred // by the backend cause whole-program undefined behavior when violated, such as @@ -295,7 +296,7 @@ def int_amdgcn_s_barrier_signal : ClangBuiltin<"__builtin_amdgcn_s_barrier_signa // If %memberCnt is 0, the member count is retained from the previous // s_barrier_init or s_barrier_signal operation. def int_amdgcn_s_barrier_signal_var : ClangBuiltin<"__builtin_amdgcn_s_barrier_signal_var">, - Intrinsic<[], [local_ptr_ty, llvm_i32_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, + Intrinsic<[], [barrier_ptr_ty, llvm_i32_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>; // bool @llvm.amdgcn.s.barrier.signal.isfirst(i32 %barrierType) @@ -307,20 +308,20 @@ def int_amdgcn_s_barrier_signal_isfirst : ClangBuiltin<"__builtin_amdgcn_s_barri // void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) %barrier, i32 %memberCnt) // The %barrier and %memberCnt argument must be uniform, otherwise behavior is undefined. def int_amdgcn_s_barrier_init : ClangBuiltin<"__builtin_amdgcn_s_barrier_init">, - Intrinsic<[], [local_ptr_ty, llvm_i32_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, + Intrinsic<[], [barrier_ptr_ty, llvm_i32_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>; // void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) %barrier) // The %barrier argument must be uniform, otherwise behavior is undefined. def int_amdgcn_s_barrier_join : ClangBuiltin<"__builtin_amdgcn_s_barrier_join">, - Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, + Intrinsic<[], [barrier_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>; // void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %barrier) // The %barrier argument must be uniform, otherwise behavior is undefined. let TargetFeatures = "s-wakeup-barrier-inst" in def int_amdgcn_s_wakeup_barrier : ClangBuiltin<"__builtin_amdgcn_s_wakeup_barrier">, - Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, + Intrinsic<[], [barrier_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>; // void @llvm.amdgcn.s.barrier.wait(i16 %barrierType) @@ -343,7 +344,7 @@ def int_amdgcn_s_get_barrier_state : ClangBuiltin<"__builtin_amdgcn_s_get_barrie // uint32_t @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) %barrier) // The %barrier argument must be uniform, otherwise behavior is undefined. def int_amdgcn_s_get_named_barrier_state : ClangBuiltin<"__builtin_amdgcn_s_get_named_barrier_state">, - Intrinsic<[llvm_i32_ty], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, + Intrinsic<[llvm_i32_ty], [barrier_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>; def int_amdgcn_wave_barrier : ClangBuiltin<"__builtin_amdgcn_wave_barrier">, diff --git a/llvm/include/llvm/Support/AMDGPUAddrSpace.h b/llvm/include/llvm/Support/AMDGPUAddrSpace.h index 01b1510524d0f..d72ba0a1415c0 100644 --- a/llvm/include/llvm/Support/AMDGPUAddrSpace.h +++ b/llvm/include/llvm/Support/AMDGPUAddrSpace.h @@ -26,8 +26,7 @@ namespace llvm { /// memory locations. namespace AMDGPUAS { enum : unsigned { - // The maximum value for flat, generic, local, private, constant and region. - MAX_AMDGPU_ADDRESS = 9, + MAX_AMDGPU_ADDRESS = 15, FLAT_ADDRESS = 0, ///< Address space for flat memory. GLOBAL_ADDRESS = 1, ///< Address space for global memory (RAT0, VTX0). @@ -47,6 +46,14 @@ enum : unsigned { BUFFER_STRIDED_POINTER = 9, ///< Address space for 192-bit fat buffer ///< pointers with an additional index. + RESERVED_0 = 10, + RESERVED_1 = 11, + RESERVED_2 = 12, + RESERVED_3 = 13, + RESERVED_4 = 14, + + BARRIER = 15, ///< Address space for modeling barrier IDs as addresses. + RESERVED_ADDRESS_SPACE_16 = 16, ///< Reserved for downstream use. /// Internal address spaces. Can be freely renumbered. @@ -84,6 +91,11 @@ enum : unsigned { // Some places use this if the address space can't be determined. UNKNOWN_ADDRESS_SPACE = ~0u, }; + +/// The BARRIER AS does not have an aperture in HW, so when converting +/// BARRIER addresses from/to generic, we represent them as LDS addresses +/// offset by a large amount so they can never alias with real LDS memory. +static constexpr unsigned BarrierAddrLDSOffset = 0x802000u; } // end namespace AMDGPUAS namespace AMDGPU { diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp index 926e18924956f..9e5cfc83c8283 100644 --- a/llvm/lib/IR/AutoUpgrade.cpp +++ b/llvm/lib/IR/AutoUpgrade.cpp @@ -7116,6 +7116,12 @@ std::string llvm::UpgradeDataLayoutString(StringRef DL, StringRef TT) { Res.replace(Res.find(OldP8), OldP8.size(), "-p8:128:128:128:48-"); if (!DL.contains("-p9") && !DL.starts_with("p9")) Res.append("-p9:192:256:256:32"); + + // Add sizing for address space 15, including the reserved address spaces + // in between. + if (!DL.contains("-p15") && !DL.starts_with("p15")) + Res.append( + "-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-p15:32:32"); } // Upgrade the ELF mangling mode. diff --git a/llvm/lib/IR/Type.cpp b/llvm/lib/IR/Type.cpp index 8c108a4d1f275..d01cf4ef53269 100644 --- a/llvm/lib/IR/Type.cpp +++ b/llvm/lib/IR/Type.cpp @@ -1118,6 +1118,8 @@ static TargetTypeInfo getTargetTypeInfo(const TargetExtType *Ty) { TargetExtType::IsTokenLike); // 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); diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h index c72fa69aa1419..2bd0d68955e39 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.h +++ b/llvm/lib/Target/AMDGPU/AMDGPU.h @@ -620,17 +620,23 @@ static inline bool addrspacesMayAlias(unsigned AS1, unsigned AS2) { // clang-format off static const bool ASAliasRules[][AMDGPUAS::MAX_AMDGPU_ADDRESS + 1] = { - /* Flat Global Region Local Constant Private Const32 BufFatPtr BufRsrc BufStrdPtr */ - /* Flat */ {true, true, false, true, true, true, true, true, true, true}, - /* Global */ {true, true, false, false, true, false, true, true, true, true}, - /* Region */ {false, false, true, false, false, false, false, false, false, false}, - /* Local */ {true, false, false, true, false, false, false, false, false, false}, - /* Constant */ {true, true, false, false, false, false, true, true, true, true}, - /* Private */ {true, false, false, false, false, true, false, false, false, false}, - /* Constant 32-bit */ {true, true, false, false, true, false, false, true, true, true}, - /* Buffer Fat Ptr */ {true, true, false, false, true, false, true, true, true, true}, - /* Buffer Resource */ {true, true, false, false, true, false, true, true, true, true}, - /* Buffer Strided Ptr */ {true, true, false, false, true, false, true, true, true, true}, + /* Flat Global Region Local Constant Private Const32 BufFatPtr BufRsrc BufStrdPtr Reserved Reserved Reserved Reserved Reserved Barrier */ + /* Flat */ {true, true, false, true, true, true, true, true, true, true, false, false, false, false, false, false}, + /* Global */ {true, true, false, false, true, false, true, true, true, true, false, false, false, false, false, false}, + /* Region */ {false, false, true, false, false, false, false, false, false, false, false, false, false, false, false, false}, + /* Local */ {true, false, false, true, false, false, false, false, false, false, false, false, false, false, false, false}, + /* Constant */ {true, true, false, false, false, false, true, true, true, true, false, false, false, false, false, false}, + /* Private */ {true, false, false, false, false, true, false, false, false, false, false, false, false, false, false, false}, + /* Constant 32-bit */ {true, true, false, false, true, false, false, true, true, true, false, false, false, false, false, false}, + /* Buffer Fat Ptr */ {true, true, false, false, true, false, true, true, true, true, false, false, false, false, false, false}, + /* Buffer Resource */ {true, true, false, false, true, false, true, true, true, true, false, false, false, false, false, false}, + /* Buffer Strided Ptr */ {true, true, false, false, true, false, true, true, true, true, false, false, false, false, false, false}, + /* Reserved */ {false, false, false, false, false, false, false, false, false, false, false, false, false, false, false, false}, + /* Reserved */ {false, false, false, false, false, false, false, false, false, false, false, false, false, false, false, false}, + /* Reserved */ {false, false, false, false, false, false, false, false, false, false, false, false, false, false, false, false}, + /* Reserved */ {false, false, false, false, false, false, false, false, false, false, false, false, false, false, false, false}, + /* Reserved */ {false, false, false, false, false, false, false, false, false, false, false, false, false, false, false, false}, + /* Barrier */ {false, false, false, false, false, false, false, false, false, false, false, false, false, false, false, true}, }; // clang-format on static_assert(std::size(ASAliasRules) == AMDGPUAS::MAX_AMDGPU_ADDRESS + 1); diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index 517272ac3c1d3..6db35fc8b0522 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -18,6 +18,7 @@ def p3 : PtrValueType; def p4 : PtrValueType; def p5 : PtrValueType; def p6 : PtrValueType; +def p15 : PtrValueType; //===-----------------------------------------------------------------------===// // AMDGPU Subtarget Feature (device properties) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp index 962988ff97e39..ed9256f068ddc 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp @@ -24,6 +24,7 @@ #include "llvm/CodeGen/MachineFrameInfo.h" #include "llvm/IR/DiagnosticInfo.h" #include "llvm/IR/IntrinsicsAMDGPU.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/KnownBits.h" #include "llvm/Target/TargetMachine.h" @@ -1546,29 +1547,33 @@ SDValue AMDGPUTargetLowering::LowerGlobalAddress(AMDGPUMachineFunctionInfo *MFI, GlobalAddressSDNode *G = cast(Op); const GlobalValue *GV = G->getGlobal(); + if (G->getAddressSpace() == AMDGPUAS::BARRIER) { + const GlobalVariable *GVar = cast(GV); + + if (!AMDGPU::isNamedBarrier(*GVar)) { + const Function &Fn = DAG.getMachineFunction().getFunction(); + DAG.getContext()->diagnose(DiagnosticInfoUnsupported( + Fn, "unsupported use of BARRIER address space", + SDLoc(Op).getDebugLoc(), DS_Error)); + return DAG.getPOISON(Op.getValueType()); + } + + unsigned Offset = MFI->allocateBarrierGlobal(DL, *cast(GV)); + return DAG.getConstant(Offset, SDLoc(Op), Op.getValueType()); + } + if (!MFI->isModuleEntryFunction()) { - bool IsNamedBarrier = AMDGPU::isNamedBarrier(*cast(GV)); - std::optional Address = - AMDGPUMachineFunctionInfo::getLDSAbsoluteAddress(*GV); - if (!Address && IsNamedBarrier) - llvm_unreachable("named barrier should have an assigned address"); - if (Address) { - if (IsNamedBarrier) { - unsigned BarCnt = cast(GV)->getGlobalSize(DL) / 16; - MFI->recordNumNamedBarriers(Address.value(), BarCnt); - } - // A constant byte offset (e.g. from a GEP into an array of named - // barriers) folds directly into the fixed LDS address. - return DAG.getConstant(*Address + G->getOffset(), SDLoc(Op), - Op.getValueType()); + if (std::optional Address = + AMDGPUMachineFunctionInfo::get32BitAbsoluteAddress( + *GV, AMDGPUAS::LOCAL_ADDRESS)) { + return DAG.getConstant(*Address, SDLoc(Op), Op.getValueType()); } } if (G->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS || G->getAddressSpace() == AMDGPUAS::REGION_ADDRESS) { if (!MFI->isModuleEntryFunction() && - GV->getName() != "llvm.amdgcn.module.lds" && - !AMDGPU::isNamedBarrier(*cast(GV))) { + GV->getName() != "llvm.amdgcn.module.lds") { SDLoc DL(Op); const Function &Fn = DAG.getMachineFunction().getFunction(); DAG.getContext()->diagnose(DiagnosticInfoUnsupported( diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp index 22b8b10554928..84b176b61b81e 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp @@ -7359,7 +7359,7 @@ bool AMDGPUInstructionSelector::selectNamedBarrierInit( std::optional BarValImm = getIConstantVRegSExtVal(BarOp.getReg(), *MRI); if (BarValImm) { - auto BarID = ((*BarValImm) >> 4) & 0x3F; + uint32_t BarID = *BarValImm & 0x3F; BuildMI(*MBB, &I, DL, TII.get(AMDGPU::S_BARRIER_SIGNAL_IMM)) .addImm(BarID); I.eraseFromParent(); @@ -7368,16 +7368,10 @@ bool AMDGPUInstructionSelector::selectNamedBarrierInit( } } - // BarID = (BarOp >> 4) & 0x3F - Register TmpReg0 = MRI->createVirtualRegister(&AMDGPU::SReg_32RegClass); - BuildMI(*MBB, &I, DL, TII.get(AMDGPU::S_LSHR_B32), TmpReg0) - .add(BarOp) - .addImm(4u) - .setOperandDead(3); // Dead scc - + // BarID = BarOp & 0x3F Register TmpReg1 = MRI->createVirtualRegister(&AMDGPU::SReg_32RegClass); BuildMI(*MBB, &I, DL, TII.get(AMDGPU::S_AND_B32), TmpReg1) - .addReg(TmpReg0) + .add(BarOp) .addImm(0x3F) .setOperandDead(3); // Dead scc @@ -7426,16 +7420,10 @@ bool AMDGPUInstructionSelector::selectNamedBarrierInst( getIConstantVRegSExtVal(BarOp.getReg(), *MRI); if (!BarValImm) { - // BarID = (BarOp >> 4) & 0x3F - Register TmpReg0 = MRI->createVirtualRegister(&AMDGPU::SReg_32RegClass); - BuildMI(*MBB, &I, DL, TII.get(AMDGPU::S_LSHR_B32), TmpReg0) - .addReg(BarOp.getReg()) - .addImm(4u) - .setOperandDead(3); // Dead scc; - + // BarID = BarOp & 0x3F Register TmpReg1 = MRI->createVirtualRegister(&AMDGPU::SReg_32RegClass); BuildMI(*MBB, &I, DL, TII.get(AMDGPU::S_AND_B32), TmpReg1) - .addReg(TmpReg0) + .addReg(BarOp.getReg()) .addImm(0x3F) .setOperandDead(3); // Dead scc; @@ -7458,7 +7446,7 @@ bool AMDGPUInstructionSelector::selectNamedBarrierInst( } if (BarValImm) { - auto BarId = ((*BarValImm) >> 4) & 0x3F; + uint32_t BarId = *BarValImm & 0x3F; MIB.addImm(BarId); } diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp index 3ef5364459350..609e5b2bbeb09 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp @@ -36,6 +36,7 @@ #include "llvm/IR/DiagnosticInfo.h" #include "llvm/IR/IntrinsicsAMDGPU.h" #include "llvm/IR/IntrinsicsR600.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #define DEBUG_TYPE "amdgpu-legalinfo" @@ -2433,15 +2434,15 @@ Register AMDGPULegalizerInfo::getSegmentAperture( const LLT I32 = LLT::integer(32); const LLT I64 = LLT::integer(64); - assert(AS == AMDGPUAS::LOCAL_ADDRESS || AS == AMDGPUAS::PRIVATE_ADDRESS); + bool IsLDS = (AS == AMDGPUAS::LOCAL_ADDRESS || AS == AMDGPUAS::BARRIER); + assert(IsLDS || AS == AMDGPUAS::PRIVATE_ADDRESS); if (ST.hasApertureRegs()) { // Note: this register is somewhat broken. When used as a 32-bit operand, // it only returns zeroes. The real value is in the upper 32 bits. // Thus, we must emit extract the high 32 bits. - const unsigned ApertureRegNo = (AS == AMDGPUAS::LOCAL_ADDRESS) - ? AMDGPU::SRC_SHARED_BASE - : AMDGPU::SRC_PRIVATE_BASE; + const unsigned ApertureRegNo = + IsLDS ? AMDGPU::SRC_SHARED_BASE : AMDGPU::SRC_PRIVATE_BASE; assert((ApertureRegNo != AMDGPU::SRC_PRIVATE_BASE || !ST.hasGloballyAddressableScratch()) && "Cannot use src_private_base with globally addressable scratch!"); @@ -2496,7 +2497,7 @@ Register AMDGPULegalizerInfo::getSegmentAperture( // Offset into amd_queue_t for group_segment_aperture_base_hi / // private_segment_aperture_base_hi. - uint32_t StructOffset = (AS == AMDGPUAS::LOCAL_ADDRESS) ? 0x40 : 0x44; + uint32_t StructOffset = IsLDS ? 0x40 : 0x44; MachineMemOperand *MMO = MF.getMachineMemOperand( PtrInfo, @@ -2565,7 +2566,7 @@ bool AMDGPULegalizerInfo::legalizeAddrSpaceCast( } if (SrcAS == AMDGPUAS::FLAT_ADDRESS && - (DestAS == AMDGPUAS::LOCAL_ADDRESS || + (DestAS == AMDGPUAS::LOCAL_ADDRESS || DestAS == AMDGPUAS::BARRIER || DestAS == AMDGPUAS::PRIVATE_ADDRESS)) { auto castFlatToLocalOrPrivate = [&](const DstOp &Dst) -> Register { if (DestAS == AMDGPUAS::PRIVATE_ADDRESS && @@ -2582,7 +2583,17 @@ bool AMDGPULegalizerInfo::legalizeAddrSpaceCast( return B.buildIntToPtr(Dst, Sub).getReg(0); } - // Extract low 32-bits of the pointer. + if (DestAS == AMDGPUAS::BARRIER) { + // flat -> barrier: extract the low 32 bits, then sub the barrier AS + // offset. + Register LoBits = B.buildExtract(S32, Src, 0).getReg(0); + Register Sub = + B.buildSub(S32, LoBits, + B.buildConstant(S32, AMDGPUAS::BarrierAddrLDSOffset)) + .getReg(0); + return B.buildIntToPtr(Dst, Sub).getReg(0); + } + return B.buildExtract(Dst, Src, 0).getReg(0); }; @@ -2611,7 +2622,7 @@ bool AMDGPULegalizerInfo::legalizeAddrSpaceCast( } if (DestAS == AMDGPUAS::FLAT_ADDRESS && - (SrcAS == AMDGPUAS::LOCAL_ADDRESS || + (SrcAS == AMDGPUAS::LOCAL_ADDRESS || SrcAS == AMDGPUAS::BARRIER || SrcAS == AMDGPUAS::PRIVATE_ADDRESS)) { auto castLocalOrPrivateToFlat = [&](const DstOp &Dst) -> Register { // Coerce the type of the low half of the result so we can use @@ -2653,6 +2664,14 @@ bool AMDGPULegalizerInfo::legalizeAddrSpaceCast( if (!ApertureReg.isValid()) return false; + if (SrcAS == AMDGPUAS::BARRIER) { + // barrier -> flat: add the barrier AS offset + SrcAsInt = + B.buildAdd(S32, SrcAsInt, + B.buildConstant(S32, AMDGPUAS::BarrierAddrLDSOffset)) + .getReg(0); + } + // TODO: Should we allow mismatched types but matching sizes in merges to // avoid the ptrtoint? return B.buildMergeLikeInstr(Dst, {SrcAsInt, ApertureReg}).getReg(0); @@ -3351,10 +3370,27 @@ bool AMDGPULegalizerInfo::legalizeGlobalValue( MachineFunction &MF = B.getMF(); SIMachineFunctionInfo *MFI = MF.getInfo(); + if (AS == AMDGPUAS::BARRIER) { + const GlobalVariable *GVar = cast(GV); + if (!AMDGPU::isNamedBarrier(*GVar)) { + const Function &Fn = MF.getFunction(); + Fn.getContext().diagnose(DiagnosticInfoUnsupported( + Fn, "unsupported use of BARRIER address space", MI.getDebugLoc(), + DS_Error)); + B.buildUndef(DstReg); + MI.eraseFromParent(); + return true; + } + + B.buildConstant(DstReg, + MFI->allocateBarrierGlobal(B.getDataLayout(), *GVar)); + MI.eraseFromParent(); + return true; + } + if (AS == AMDGPUAS::LOCAL_ADDRESS || AS == AMDGPUAS::REGION_ADDRESS) { if (!MFI->isModuleEntryFunction() && - GV->getName() != "llvm.amdgcn.module.lds" && - !AMDGPU::isNamedBarrier(*cast(GV))) { + GV->getName() != "llvm.amdgcn.module.lds") { const Function &Fn = MF.getFunction(); Fn.getContext().diagnose(DiagnosticInfoUnsupported( Fn, "local memory global used by non-kernel function", diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp index 707c3ec975d73..95ba2586efb09 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp @@ -6,12 +6,10 @@ // //===----------------------------------------------------------------------===// // -// Lower LDS global variables with target extension type "amdgpu.named.barrier" +// Lower global variables with target extension type "amdgpu.named.barrier" // that require specialized address assignment. It assigns a unique -// barrier identifier to each named-barrier LDS variable and encodes +// barrier identifier to each named-barrier variable and encodes // this identifier within the !absolute_symbol metadata of that global. -// This encoding ensures that subsequent LDS lowering passes can process these -// barriers correctly without conflicts. // //===----------------------------------------------------------------------===// @@ -37,10 +35,14 @@ using namespace AMDGPU; namespace { +static bool isNamedBarrierToLower(const GlobalVariable &GV) { + return isNamedBarrier(GV) && !GV.isAbsoluteSymbolRef(); +} + // Write the specified address into metadata where it can be retrieved by // the assembler. Format is a half open range, [Address Address+1) -static void recordLDSAbsoluteAddress(Module *M, GlobalVariable *GV, - uint32_t Address) { +static void recordAbsoluteAddress(Module *M, GlobalVariable *GV, + uint32_t Address) { LLVMContext &Ctx = M->getContext(); auto *IntTy = M->getDataLayout().getIntPtrType(Ctx, AMDGPUAS::LOCAL_ADDRESS); auto *MinC = ConstantAsMetadata::get(ConstantInt::get(IntTy, Address)); @@ -137,14 +139,12 @@ static bool lowerExecSyncGlobalVariables(Module &M, GVUsesInfoTy &GVUsesInfo) { LLVM_DEBUG(GV->printAsOperand(dbgs(), false); dbgs() << " was assigned barrier id: " << BarID << " id-count: " << BarCnt << "\n"); - // 4 bits for alignment, 5 bits for the barrier num, - // 3 bits for the barrier scope - Offset = 0x802000u | BarrierScope << 9 | BarID << 4; + Offset = BarID; } else { llvm_unreachable("Unhandled special variable type."); } - recordLDSAbsoluteAddress(&M, GV, Offset); + recordAbsoluteAddress(&M, GV, Offset); } // Also erase those special LDS variables from indirect_access. @@ -214,15 +214,16 @@ static bool runLowerExecSyncGlobals(Module &M) { CallGraph CG = CallGraph(M); bool Changed = false; Changed |= - eliminateGVConstantExprUsesFromAllInstructions(M, isLDSVariableToLower); + eliminateGVConstantExprUsesFromAllInstructions(M, isNamedBarrierToLower); // For each kernel, what variables does it access directly or through // callees - GVUsesInfoTy LDSUsesInfo = getTransitiveUsesOfLDSForLowering(CG, M); + GVUsesInfoTy BarrierUsesInfo = + getTransitiveUsesOfGV(CG, M, isNamedBarrierToLower); - if (hasBarrierToLower(LDSUsesInfo)) { + if (hasBarrierToLower(BarrierUsesInfo)) { // Special LDS variables need special address assignment - Changed |= lowerExecSyncGlobalVariables(M, LDSUsesInfo); + Changed |= lowerExecSyncGlobalVariables(M, BarrierUsesInfo); } return Changed; diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp index ef86d279d193b..12e2478e055b0 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp @@ -937,16 +937,6 @@ class AMDGPULowerModuleLDS { for (auto &[F, Vars] : FunctionLDSUses) AllLDSUses[F].insert(Vars.begin(), Vars.end()); - // Named barriers are handled by AMDGPULowerExecSync; filter them out. - for (auto &[F, Vars] : AllLDSUses) { - SmallVector Barriers; - for (GlobalVariable *V : Vars) - if (AMDGPU::isNamedBarrier(*V)) - Barriers.push_back(V); - for (GlobalVariable *V : Barriers) - Vars.erase(V); - } - // Build reverse map: LDS variable -> functions that use it. DenseMap> VarToFuncs; for (auto &[F, Vars] : AllLDSUses) { diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp b/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp index 3f0fb64be4886..c375f86980a57 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp @@ -31,6 +31,7 @@ #include "llvm/MC/MCInst.h" #include "llvm/MC/MCObjectStreamer.h" #include "llvm/MC/MCStreamer.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/Endian.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/Format.h" @@ -286,7 +287,8 @@ const MCExpr *AMDGPUAsmPrinter::lowerConstant(const Constant *CV, // Intercept LDS variables with known addresses if (const GlobalVariable *GV = dyn_cast(CV)) { if (std::optional Address = - AMDGPUMachineFunctionInfo::getLDSAbsoluteAddress(*GV)) { + AMDGPUMachineFunctionInfo::get32BitAbsoluteAddress( + *GV, AMDGPUAS::LOCAL_ADDRESS)) { auto *IntTy = Type::getInt32Ty(CV->getContext()); return AsmPrinter::lowerConstant(ConstantInt::get(IntTy, *Address), BaseCV, Offset); diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.cpp index 3e8a75a7eb840..d8a47f8f895d7 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.cpp @@ -15,6 +15,7 @@ #include "llvm/IR/ConstantRange.h" #include "llvm/IR/Constants.h" #include "llvm/IR/Metadata.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Target/TargetMachine.h" using namespace llvm; @@ -97,17 +98,8 @@ unsigned AMDGPUMachineFunctionInfo::allocateLDSGlobal(const DataLayout &DL, unsigned Offset; if (GV.getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) { - if (AMDGPU::isNamedBarrier(GV)) { - std::optional BarAddr = getLDSAbsoluteAddress(GV); - if (!BarAddr) - llvm_unreachable("named barrier should have an assigned address"); - Entry.first->second = BarAddr.value(); - unsigned BarCnt = GV.getGlobalSize(DL) / 16; - recordNumNamedBarriers(BarAddr.value(), BarCnt); - return BarAddr.value(); - } - - std::optional MaybeAbs = getLDSAbsoluteAddress(GV); + std::optional MaybeAbs = + get32BitAbsoluteAddress(GV, AMDGPUAS::LOCAL_ADDRESS); if (MaybeAbs) { // Absolute address LDS variables that exist prior to the LDS lowering // pass raise a fatal error in that pass. These failure modes are only @@ -165,6 +157,30 @@ unsigned AMDGPUMachineFunctionInfo::allocateLDSGlobal(const DataLayout &DL, return Offset; } +unsigned +AMDGPUMachineFunctionInfo::allocateBarrierGlobal(const DataLayout &DL, + const GlobalVariable &GV) { + assert(AMDGPU::isNamedBarrier(GV)); + std::optional BarAddr = + get32BitAbsoluteAddress(GV, AMDGPUAS::BARRIER); + if (!BarAddr) { + reportFatalInternalError("named barrier global variable '" + GV.getName() + + "' does not have an address assigned"); + } + + if (*BarAddr == 0) { + // We cannot allow this because some places in CodeGen (rightfully) assume a + // GV address is never null. For example, there are no null checks on + // addrspacecast if the pointer is a GV pointer. + reportFatalInternalError("named barrier global variable '" + GV.getName() + + "' has a NULL address, which is not supported"); + } + + unsigned BarCnt = AMDGPU::getNumNamedBarriersDeclared(DL, GV); + recordNumNamedBarriers(BarAddr.value(), BarCnt); + return BarAddr.value(); +} + std::optional AMDGPUMachineFunctionInfo::getLDSKernelIdMetadata(const Function &F) { // TODO: Would be more consistent with the abs symbols to use a range @@ -182,8 +198,9 @@ AMDGPUMachineFunctionInfo::getLDSKernelIdMetadata(const Function &F) { } std::optional -AMDGPUMachineFunctionInfo::getLDSAbsoluteAddress(const GlobalValue &GV) { - if (GV.getAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) +AMDGPUMachineFunctionInfo::get32BitAbsoluteAddress(const GlobalValue &GV, + unsigned AS) { + if (GV.getAddressSpace() != AS) return {}; std::optional AbsSymRange = GV.getAbsoluteSymbolRange(); @@ -221,7 +238,8 @@ void AMDGPUMachineFunctionInfo::setDynLDSAlign(const Function &F, const GlobalVariable *Dyn = getKernelDynLDSGlobalFromFunction(F); if (Dyn) { unsigned Offset = LDSSize; // return this? - std::optional Expect = getLDSAbsoluteAddress(*Dyn); + std::optional Expect = + get32BitAbsoluteAddress(GV, AMDGPUAS::LOCAL_ADDRESS); if (!Expect || (Offset != *Expect)) { report_fatal_error("Inconsistent metadata on dynamic LDS variable"); } diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.h index 36db6c2dd0d12..c65592bd965ba 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.h +++ b/llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.h @@ -82,7 +82,7 @@ class AMDGPUMachineFunctionInfo : public MachineFunctionInfo { void recordNumNamedBarriers(uint32_t GVAddr, unsigned BarCnt) { NumNamedBarriers = - std::max(NumNamedBarriers, ((GVAddr & 0x1ff) >> 4) + BarCnt - 1); + std::max(NumNamedBarriers, (GVAddr & 0x1ff) + BarCnt - 1); } uint32_t getNumNamedBarriers() const { return NumNamedBarriers; } @@ -109,8 +109,13 @@ class AMDGPUMachineFunctionInfo : public MachineFunctionInfo { unsigned allocateLDSGlobal(const DataLayout &DL, const GlobalVariable &GV, Align Trailing); + unsigned allocateBarrierGlobal(const DataLayout &DL, + const GlobalVariable &GV); + static std::optional getLDSKernelIdMetadata(const Function &F); - static std::optional getLDSAbsoluteAddress(const GlobalValue &GV); + + static std::optional get32BitAbsoluteAddress(const GlobalValue &GV, + unsigned AS); Align getDynLDSAlign() const { return DynLDSAlign; } diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp b/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp index eacce0d5242e1..30e2a380820b3 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp @@ -19,6 +19,7 @@ #include "llvm/IR/IntrinsicsAMDGPU.h" #include "llvm/IR/LLVMContext.h" #include "llvm/IR/ReplaceConstant.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #define DEBUG_TYPE "amdgpu-memory-utils" @@ -76,11 +77,21 @@ static TargetExtType *getTargetExtType(const GlobalVariable &GV) { } TargetExtType *isNamedBarrier(const GlobalVariable &GV) { + if (GV.getAddressSpace() != AMDGPUAS::BARRIER) + return nullptr; if (TargetExtType *Ty = getTargetExtType(GV)) return Ty->getName() == "amdgcn.named.barrier" ? Ty : nullptr; return nullptr; } +unsigned getNumNamedBarriersDeclared(const DataLayout &DL, + const GlobalVariable &GV) { + assert(isNamedBarrier(GV)); + unsigned GVSize = GV.getGlobalSize(DL); + assert(GVSize && (GVSize % NamedBarrierTypeSizeInBytes == 0)); + return GVSize / NamedBarrierTypeSizeInBytes; +} + bool isDynamicLDS(const GlobalVariable &GV) { // external zero size addrspace(3) without initializer is dynlds. const Module *M = GV.getParent(); @@ -292,15 +303,6 @@ GVUsesInfoTy getTransitiveUsesOfLDSForLowering(const CallGraph &CG, Module &M) { if (IsDirectMapDynLDSGV) continue; - // TODO: Remove once barriers are no longer in the LDS AS. - if (isNamedBarrier(*GV)) { - if (IsAbsolute) { - UsesInfo.DirectAccess[Fn].erase(GV); - UsesInfo.IndirectAccess[Fn].erase(GV); - } - continue; - } - if (HasAbsoluteGVs.has_value()) { if (*HasAbsoluteGVs != IsAbsolute) { reportFatalUsageError( diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h b/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h index 4e164d08549dd..93fee16594e69 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h +++ b/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.h @@ -30,6 +30,8 @@ class TargetExtType; namespace AMDGPU { +static constexpr unsigned NamedBarrierTypeSizeInBytes = 16; + using FunctionVariableMap = DenseMap>; using VariableFunctionMap = DenseMap>; @@ -43,6 +45,10 @@ void copyMetadataForWidenedLoad(LoadInst &Dest, const LoadInst &Source); // If GV is a named-barrier return its type. Otherwise return nullptr. TargetExtType *isNamedBarrier(const GlobalVariable &GV); +/// \returns how many named barriers are declared by \p GV. +unsigned getNumNamedBarriersDeclared(const DataLayout &DL, + const GlobalVariable &GV); + bool isDynamicLDS(const GlobalVariable &GV); bool isLDSVariableToLower(const GlobalVariable &GV); diff --git a/llvm/lib/Target/AMDGPU/SIDefines.h b/llvm/lib/Target/AMDGPU/SIDefines.h index a7dd7b5f8dd10..ceffe78ba2676 100644 --- a/llvm/lib/Target/AMDGPU/SIDefines.h +++ b/llvm/lib/Target/AMDGPU/SIDefines.h @@ -1355,10 +1355,6 @@ enum Type { NAMED_BARRIER_LAST = 16, }; -enum { - BARRIER_SCOPE_WORKGROUP = 0, -}; - } // namespace Barrier } // namespace AMDGPU diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index f37e531d39648..b3b8ee8959647 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -44,6 +44,7 @@ #include "llvm/IR/IntrinsicsAMDGPU.h" #include "llvm/IR/IntrinsicsR600.h" #include "llvm/IR/MDBuilder.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/KnownBits.h" #include "llvm/Support/ModRef.h" @@ -8545,9 +8546,11 @@ bool SITargetLowering::shouldUseLDSConstAddress(const GlobalValue *GV) const { // linker can assign their offsets. if (AMDGPUTargetMachine::EnableObjectLinking) { if (const auto *GVar = dyn_cast(GV)) { - if (GVar->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) { - assert(GVar->isDeclaration() && "AS3 GVs should be declaration here " - "when object linking is enabled"); + if (GVar->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS || + GVar->getAddressSpace() == AMDGPUAS::BARRIER) { + assert(GVar->isDeclaration() && + "AS 3 & 13 GVs should be declaration here " + "when object linking is enabled"); return false; } } @@ -9322,10 +9325,11 @@ SDValue SITargetLowering::LowerINLINEASM(SDValue Op, SelectionDAG &DAG) const { SDValue SITargetLowering::getSegmentAperture(unsigned AS, const SDLoc &DL, SelectionDAG &DAG) const { + const bool IsLDS = (AS == AMDGPUAS::LOCAL_ADDRESS || AS == AMDGPUAS::BARRIER); + if (Subtarget->hasApertureRegs()) { - const unsigned ApertureRegNo = (AS == AMDGPUAS::LOCAL_ADDRESS) - ? AMDGPU::SRC_SHARED_BASE - : AMDGPU::SRC_PRIVATE_BASE; + const unsigned ApertureRegNo = + IsLDS ? AMDGPU::SRC_SHARED_BASE : AMDGPU::SRC_PRIVATE_BASE; assert((ApertureRegNo != AMDGPU::SRC_PRIVATE_BASE || !Subtarget->hasGloballyAddressableScratch()) && "Cannot use src_private_base with globally addressable scratch!"); @@ -9347,8 +9351,7 @@ SDValue SITargetLowering::getSegmentAperture(unsigned AS, const SDLoc &DL, // implicit kernargs. const Module *M = DAG.getMachineFunction().getFunction().getParent(); if (AMDGPU::getAMDHSACodeObjectVersion(*M) >= AMDGPU::AMDHSA_COV5) { - ImplicitParameter Param = - (AS == AMDGPUAS::LOCAL_ADDRESS) ? SHARED_BASE : PRIVATE_BASE; + ImplicitParameter Param = IsLDS ? SHARED_BASE : PRIVATE_BASE; return loadImplicitKernelArgument(DAG, MVT::i32, DL, Align(4), Param); } @@ -9366,7 +9369,7 @@ SDValue SITargetLowering::getSegmentAperture(unsigned AS, const SDLoc &DL, // Offset into amd_queue_t for group_segment_aperture_base_hi / // private_segment_aperture_base_hi. - uint32_t StructOffset = (AS == AMDGPUAS::LOCAL_ADDRESS) ? 0x40 : 0x44; + uint32_t StructOffset = IsLDS ? 0x40 : 0x44; SDValue Ptr = DAG.getObjectPtrOffset(DL, QueuePtr, TypeSize::getFixed(StructOffset)); @@ -9422,10 +9425,10 @@ SDValue SITargetLowering::lowerADDRSPACECAST(SDValue Op, SDValue FlatNullPtr = DAG.getConstant(0, SL, MVT::i64); - // flat -> local/private + // flat -> local/private/barrier if (SrcAS == AMDGPUAS::FLAT_ADDRESS) { if (DestAS == AMDGPUAS::LOCAL_ADDRESS || - DestAS == AMDGPUAS::PRIVATE_ADDRESS) { + DestAS == AMDGPUAS::PRIVATE_ADDRESS || DestAS == AMDGPUAS::BARRIER) { SDValue Ptr = DAG.getNode(ISD::TRUNCATE, SL, MVT::i32, Src); if (DestAS == AMDGPUAS::PRIVATE_ADDRESS && @@ -9438,6 +9441,11 @@ SDValue SITargetLowering::lowerADDRSPACECAST(SDValue Op, DAG.getRegister(AMDGPU::SRC_FLAT_SCRATCH_BASE_LO, MVT::i32)), 0); Ptr = DAG.getNode(ISD::SUB, SL, MVT::i32, Ptr, FlatScratchBaseLo); + } else if (DestAS == AMDGPUAS::BARRIER) { + // flat -> barrier: sub the barrier AS offset. + Ptr = DAG.getNode( + ISD::SUB, SL, MVT::i32, Ptr, + DAG.getConstant(AMDGPUAS::BarrierAddrLDSOffset, SL, MVT::i32)); } if (IsNonNull || isKnownNonNull(Op, DAG, TM, SrcAS)) @@ -9452,10 +9460,10 @@ SDValue SITargetLowering::lowerADDRSPACECAST(SDValue Op, } } - // local/private -> flat + // local/private/barrier -> flat if (DestAS == AMDGPUAS::FLAT_ADDRESS) { if (SrcAS == AMDGPUAS::LOCAL_ADDRESS || - SrcAS == AMDGPUAS::PRIVATE_ADDRESS) { + SrcAS == AMDGPUAS::PRIVATE_ADDRESS || SrcAS == AMDGPUAS::BARRIER) { SDValue CvtPtr; if (SrcAS == AMDGPUAS::PRIVATE_ADDRESS && Subtarget->hasGloballyAddressableScratch()) { @@ -9487,7 +9495,19 @@ SDValue SITargetLowering::lowerADDRSPACECAST(SDValue Op, CvtPtr = DAG.getNode(ISD::ADD, SL, MVT::i64, CvtPtr, FlatScratchBase); } else { SDValue Aperture = getSegmentAperture(SrcAS, SL, DAG); - CvtPtr = DAG.getNode(ISD::BUILD_VECTOR, SL, MVT::v2i32, Src, Aperture); + + if (SrcAS == AMDGPUAS::BARRIER) { + // barrier -> flat: add the barrier AS offset. + SDValue SrcOffset = DAG.getNode( + ISD::ADD, SL, MVT::i32, Src, + DAG.getConstant(AMDGPUAS::BarrierAddrLDSOffset, SL, MVT::i32)); + CvtPtr = DAG.getNode(ISD::BUILD_VECTOR, SL, MVT::v2i32, SrcOffset, + Aperture); + } else { + CvtPtr = + DAG.getNode(ISD::BUILD_VECTOR, SL, MVT::v2i32, Src, Aperture); + } + CvtPtr = DAG.getNode(ISD::BITCAST, SL, MVT::i64, CvtPtr); } @@ -10041,12 +10061,11 @@ SDValue SITargetLowering::LowerGlobalAddress(AMDGPUMachineFunctionInfo *MFI, EVT PtrVT = Op.getValueType(); const GlobalValue *GV = GSD->getGlobal(); - if ((GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS && + const unsigned AS = GSD->getAddressSpace(); + if (((AS == AMDGPUAS::LOCAL_ADDRESS || AS == AMDGPUAS::BARRIER) && shouldUseLDSConstAddress(GV)) || - GSD->getAddressSpace() == AMDGPUAS::REGION_ADDRESS || - GSD->getAddressSpace() == AMDGPUAS::PRIVATE_ADDRESS) { - if (GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS && - GV->hasExternalLinkage()) { + AS == AMDGPUAS::REGION_ADDRESS || AS == AMDGPUAS::PRIVATE_ADDRESS) { + if (AS == AMDGPUAS::LOCAL_ADDRESS && GV->hasExternalLinkage()) { const GlobalVariable &GVar = *cast(GV); // HIP uses an unsized array `extern __shared__ T s[]` or similar // zero-sized type in other languages to declare the dynamic shared @@ -10066,7 +10085,13 @@ SDValue SITargetLowering::LowerGlobalAddress(AMDGPUMachineFunctionInfo *MFI, return AMDGPUTargetLowering::LowerGlobalAddress(MFI, Op, DAG); } - if (GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) { + if (AS == AMDGPUAS::BARRIER) { + SDValue GA = DAG.getTargetGlobalAddress(GV, DL, MVT::i32, GSD->getOffset(), + SIInstrInfo::MO_ABS32_LO); + return SDValue(DAG.getMachineNode(AMDGPU::S_MOV_B32, DL, MVT::i32, GA), 0); + } + + if (AS == AMDGPUAS::LOCAL_ADDRESS) { SDValue GA = DAG.getTargetGlobalAddress(GV, DL, MVT::i32, GSD->getOffset(), SIInstrInfo::MO_ABS32_LO); return DAG.getNode(AMDGPUISD::LDS, DL, MVT::i32, GA); @@ -12112,7 +12137,7 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op, if (isa(Op->getOperand(2))) { uint64_t BarID = cast(Op->getOperand(2))->getZExtValue(); if (IntrID == Intrinsic::amdgcn_s_get_named_barrier_state) - BarID = (BarID >> 4) & 0x3F; + BarID = BarID & 0x3F; Opc = AMDGPU::S_GET_BARRIER_STATE_IMM; SDValue K = DAG.getTargetConstant(BarID, DL, MVT::i32); Ops.push_back(K); @@ -12120,11 +12145,8 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op, } else { Opc = AMDGPU::S_GET_BARRIER_STATE_M0; if (IntrID == Intrinsic::amdgcn_s_get_named_barrier_state) { - SDValue M0Val; - M0Val = DAG.getNode(ISD::SRL, DL, MVT::i32, Op->getOperand(2), - DAG.getShiftAmountConstant(4, MVT::i32, DL)); - M0Val = DAG.getNode(ISD::AND, DL, MVT::i32, M0Val, - DAG.getConstant(0x3F, DL, MVT::i32)); + SDValue M0Val = DAG.getNode(ISD::AND, DL, MVT::i32, Op->getOperand(2), + DAG.getConstant(0x3F, DL, MVT::i32)); Ops.push_back(copyToM0(DAG, Chain, DL, M0Val).getValue(0)); } else Ops.push_back(copyToM0(DAG, Chain, DL, Op->getOperand(2)).getValue(0)); @@ -12734,12 +12756,12 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, if (auto *C = dyn_cast(BarOp)) BarVal = C->getZExtValue(); else if (auto *GA = dyn_cast(BarOp)) - if (auto Addr = AMDGPUMachineFunctionInfo::getLDSAbsoluteAddress( - *GA->getGlobal())) + if (auto Addr = AMDGPUMachineFunctionInfo::get32BitAbsoluteAddress( + *GA->getGlobal(), AMDGPUAS::BARRIER)) BarVal = *Addr + GA->getOffset(); if (BarVal) { - unsigned BarID = (*BarVal >> 4) & 0x3F; + unsigned BarID = *BarVal & 0x3F; Ops.push_back(DAG.getTargetConstant(BarID, DL, MVT::i32)); Ops.push_back(Chain); auto *NewMI = DAG.getMachineNode(AMDGPU::S_BARRIER_SIGNAL_IMM, DL, @@ -12759,12 +12781,9 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, unsigned Opc = IntrinsicID == Intrinsic::amdgcn_s_barrier_init ? AMDGPU::S_BARRIER_INIT_M0 : AMDGPU::S_BARRIER_SIGNAL_M0; - // extract the BarrierID from bits 4-9 of BarOp - SDValue BarID; - BarID = DAG.getNode(ISD::SRL, DL, MVT::i32, BarOp, - DAG.getShiftAmountConstant(4, MVT::i32, DL)); - BarID = DAG.getNode(ISD::AND, DL, MVT::i32, BarID, - DAG.getConstant(0x3F, DL, MVT::i32)); + // extract the BarrierID from bits 0-5 of BarOp + SDValue BarID = DAG.getNode(ISD::AND, DL, MVT::i32, BarOp, + DAG.getConstant(0x3F, DL, MVT::i32)); // Member count should be put into M0[ShAmt:+6] // Barrier ID should be put into M0[5:0] SDValue MemberCnt = DAG.getNode(ISD::AND, DL, MVT::i32, CntOp, @@ -12804,8 +12823,8 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, Opc = AMDGPU::S_WAKEUP_BARRIER_IMM; break; } - // extract the BarrierID from bits 4-9 of the immediate - unsigned BarID = (BarVal >> 4) & 0x3F; + // extract the BarrierID from bits 0-5 of the immediate + unsigned BarID = BarVal & 0x3F; SDValue K = DAG.getTargetConstant(BarID, DL, MVT::i32); Ops.push_back(K); Ops.push_back(Chain); @@ -12820,12 +12839,9 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, Opc = AMDGPU::S_WAKEUP_BARRIER_M0; break; } - // extract the BarrierID from bits 4-9 of BarOp, copy to M0[5:0] - SDValue M0Val; - M0Val = DAG.getNode(ISD::SRL, DL, MVT::i32, BarOp, - DAG.getShiftAmountConstant(4, MVT::i32, DL)); - M0Val = DAG.getNode(ISD::AND, DL, MVT::i32, M0Val, - DAG.getConstant(0x3F, DL, MVT::i32)); + // extract the BarrierID from bits 0-5 of BarOp, copy to M0[5:0] + SDValue M0Val = DAG.getNode(ISD::AND, DL, MVT::i32, BarOp, + DAG.getConstant(0x3F, DL, MVT::i32)); Ops.push_back(copyToM0(DAG, Chain, DL, M0Val).getValue(0)); } diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.td b/llvm/lib/Target/AMDGPU/SIRegisterInfo.td index 31f927709e682..95afe4f075dfb 100644 --- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.td +++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.td @@ -596,7 +596,7 @@ class RegisterTypes reg_types> { def Reg16Types : RegisterTypes<[i16, f16, bf16]>; def Reg32DataTypes: RegisterTypes<[i32, f32, v2i16, v2f16, v2bf16]>; -def Reg32PtrTypes: RegisterTypes<[p2, p3, p5, p6]>; +def Reg32PtrTypes: RegisterTypes<[p2, p3, p5, p6, p15]>; def Reg32Types : RegisterTypes; def Reg64DataTypes: RegisterTypes<[i64, f64, v2i32, v2f32, v4i16, v4f16, v4bf16]>; def Reg64PtrTypes: RegisterTypes<[p0, p1, p4]>; diff --git a/llvm/lib/TargetParser/TargetDataLayout.cpp b/llvm/lib/TargetParser/TargetDataLayout.cpp index 8b6f46642e4fa..54b4cd44a1728 100644 --- a/llvm/lib/TargetParser/TargetDataLayout.cpp +++ b/llvm/lib/TargetParser/TargetDataLayout.cpp @@ -274,8 +274,9 @@ static std::string computeAMDDataLayout(const Triple &TT) { // space 8) which cannot be non-trivilally accessed by LLVM memory operations // like getelementptr. return "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32" - "-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-" - "v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-" + "-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-p15:32:32" + "-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:" + "512-" "v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9"; } diff --git a/llvm/test/Analysis/UniformityAnalysis/AMDGPU/always_uniform.ll b/llvm/test/Analysis/UniformityAnalysis/AMDGPU/always_uniform.ll index 6e75d4ead8cdd..d24fe03993482 100644 --- a/llvm/test/Analysis/UniformityAnalysis/AMDGPU/always_uniform.ll +++ b/llvm/test/Analysis/UniformityAnalysis/AMDGPU/always_uniform.ll @@ -160,10 +160,10 @@ define i32 @s_get_barrier_state(i32 %bar) { } ; CHECK-LABEL: for function 's_get_named_barrier_state': -; CHECK: DIVERGENT: ptr addrspace(3) %bar +; CHECK: DIVERGENT: ptr addrspace(15) %bar ; CHECK-NOT: DIVERGENT -define i32 @s_get_named_barrier_state(ptr addrspace(3) %bar) { - %result = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) %bar) +define i32 @s_get_named_barrier_state(ptr addrspace(15) %bar) { + %result = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) %bar) ret i32 %result } diff --git a/llvm/test/CodeGen/AMDGPU/addrspacecast-barrier.ll b/llvm/test/CodeGen/AMDGPU/addrspacecast-barrier.ll new file mode 100644 index 0000000000000..0318d3afde1b8 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/addrspacecast-barrier.ll @@ -0,0 +1,474 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 < %s | FileCheck -check-prefixes=GFX942,GFX942-SDAG %s +; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 < %s | FileCheck -check-prefixes=GFX942,GFX942-GISEL %s + +; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1030 < %s | FileCheck -check-prefixes=GFX1030,GFX1030-SDAG %s +; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1030 < %s | FileCheck -check-prefixes=GFX1030,GFX1030-GISEL %s + +; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1200 < %s | FileCheck -check-prefixes=GFX1200,GFX1200-SDAG %s +; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1200 < %s | FileCheck -check-prefixes=GFX1200,GFX1200-GISEL %s + +; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s | FileCheck -check-prefixes=GFX1250,GFX1250-SDAG %s +; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s | FileCheck -check-prefixes=GFX1250,GFX1250-GISEL %s + +@bar = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison + +define amdgpu_kernel void @barrier_to_generic(ptr addrspace(15) %bar, ptr %out) { +; GFX942-SDAG-LABEL: barrier_to_generic: +; GFX942-SDAG: ; %bb.0: +; GFX942-SDAG-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX942-SDAG-NEXT: s_load_dword s0, s[4:5], 0x0 +; GFX942-SDAG-NEXT: s_load_dwordx2 s[2:3], s[4:5], 0x8 +; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-SDAG-NEXT: s_add_i32 s4, s0, 0x802000 +; GFX942-SDAG-NEXT: s_cmp_lg_u32 s0, 0 +; GFX942-SDAG-NEXT: s_cselect_b32 s0, s4, 0 +; GFX942-SDAG-NEXT: s_cselect_b32 s1, s1, 0 +; GFX942-SDAG-NEXT: v_mov_b32_e32 v2, s0 +; GFX942-SDAG-NEXT: v_mov_b32_e32 v3, s1 +; GFX942-SDAG-NEXT: v_mov_b64_e32 v[0:1], s[2:3] +; GFX942-SDAG-NEXT: flat_store_dwordx2 v[0:1], v[2:3] +; GFX942-SDAG-NEXT: s_endpgm +; +; GFX942-GISEL-LABEL: barrier_to_generic: +; GFX942-GISEL: ; %bb.0: +; GFX942-GISEL-NEXT: s_load_dword s6, s[4:5], 0x0 +; GFX942-GISEL-NEXT: s_load_dwordx2 s[2:3], s[4:5], 0x8 +; GFX942-GISEL-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-GISEL-NEXT: s_add_u32 s0, s6, 0x802000 +; GFX942-GISEL-NEXT: s_cmp_lg_u32 s6, 0 +; GFX942-GISEL-NEXT: s_cselect_b64 s[0:1], s[0:1], 0 +; GFX942-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1] +; GFX942-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[2:3] +; GFX942-GISEL-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX942-GISEL-NEXT: s_endpgm +; +; GFX1030-SDAG-LABEL: barrier_to_generic: +; GFX1030-SDAG: ; %bb.0: +; GFX1030-SDAG-NEXT: s_add_u32 s12, s12, s17 +; GFX1030-SDAG-NEXT: s_addc_u32 s13, s13, 0 +; GFX1030-SDAG-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_LO), s12 +; GFX1030-SDAG-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_HI), s13 +; GFX1030-SDAG-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1030-SDAG-NEXT: s_clause 0x1 +; GFX1030-SDAG-NEXT: s_load_dword s0, s[8:9], 0x0 +; GFX1030-SDAG-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x8 +; GFX1030-SDAG-NEXT: s_waitcnt lgkmcnt(0) +; GFX1030-SDAG-NEXT: s_add_i32 s4, s0, 0x802000 +; GFX1030-SDAG-NEXT: s_cmp_lg_u32 s0, 0 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v2, s2 +; GFX1030-SDAG-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1030-SDAG-NEXT: s_cselect_b32 s1, s1, 0 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v0, s0 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v1, s1 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v3, s3 +; GFX1030-SDAG-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX1030-SDAG-NEXT: s_endpgm +; +; GFX1030-GISEL-LABEL: barrier_to_generic: +; GFX1030-GISEL: ; %bb.0: +; GFX1030-GISEL-NEXT: s_add_u32 s12, s12, s17 +; GFX1030-GISEL-NEXT: s_addc_u32 s13, s13, 0 +; GFX1030-GISEL-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_LO), s12 +; GFX1030-GISEL-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_HI), s13 +; GFX1030-GISEL-NEXT: s_clause 0x1 +; GFX1030-GISEL-NEXT: s_load_dword s4, s[8:9], 0x0 +; GFX1030-GISEL-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x8 +; GFX1030-GISEL-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1030-GISEL-NEXT: s_waitcnt lgkmcnt(0) +; GFX1030-GISEL-NEXT: s_add_u32 s0, s4, 0x802000 +; GFX1030-GISEL-NEXT: s_cmp_lg_u32 s4, 0 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v2, s2 +; GFX1030-GISEL-NEXT: s_cselect_b64 s[0:1], s[0:1], 0 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v3, s3 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v0, s0 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v1, s1 +; GFX1030-GISEL-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX1030-GISEL-NEXT: s_endpgm +; +; GFX1200-SDAG-LABEL: barrier_to_generic: +; GFX1200-SDAG: ; %bb.0: +; GFX1200-SDAG-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1200-SDAG-NEXT: s_clause 0x1 +; GFX1200-SDAG-NEXT: s_load_b32 s0, s[4:5], 0x0 +; GFX1200-SDAG-NEXT: s_load_b64 s[2:3], s[4:5], 0x8 +; GFX1200-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1200-SDAG-NEXT: s_add_co_i32 s4, s0, 0x802000 +; GFX1200-SDAG-NEXT: s_cmp_lg_u32 s0, 0 +; GFX1200-SDAG-NEXT: v_dual_mov_b32 v2, s2 :: v_dual_mov_b32 v3, s3 +; GFX1200-SDAG-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1200-SDAG-NEXT: s_cselect_b32 s1, s1, 0 +; GFX1200-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1200-SDAG-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1 +; GFX1200-SDAG-NEXT: flat_store_b64 v[2:3], v[0:1] +; GFX1200-SDAG-NEXT: s_endpgm +; +; GFX1200-GISEL-LABEL: barrier_to_generic: +; GFX1200-GISEL: ; %bb.0: +; GFX1200-GISEL-NEXT: s_clause 0x1 +; GFX1200-GISEL-NEXT: s_load_b32 s6, s[4:5], 0x0 +; GFX1200-GISEL-NEXT: s_load_b64 s[2:3], s[4:5], 0x8 +; GFX1200-GISEL-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1200-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1200-GISEL-NEXT: s_add_co_u32 s0, s6, 0x802000 +; GFX1200-GISEL-NEXT: s_cmp_lg_u32 s6, 0 +; GFX1200-GISEL-NEXT: v_dual_mov_b32 v2, s2 :: v_dual_mov_b32 v3, s3 +; GFX1200-GISEL-NEXT: s_cselect_b64 s[0:1], s[0:1], 0 +; GFX1200-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1200-GISEL-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1 +; GFX1200-GISEL-NEXT: flat_store_b64 v[2:3], v[0:1] +; GFX1200-GISEL-NEXT: s_endpgm +; +; GFX1250-SDAG-LABEL: barrier_to_generic: +; GFX1250-SDAG: ; %bb.0: +; GFX1250-SDAG-NEXT: global_wb +; GFX1250-SDAG-NEXT: v_nop +; GFX1250-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; GFX1250-SDAG-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1250-SDAG-NEXT: s_clause 0x1 +; GFX1250-SDAG-NEXT: s_load_b32 s0, s[4:5], 0x0 nv +; GFX1250-SDAG-NEXT: s_load_b64 s[2:3], s[4:5], 0x8 nv +; GFX1250-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1250-SDAG-NEXT: s_add_co_i32 s4, s0, 0x802000 +; GFX1250-SDAG-NEXT: s_cmp_lg_u32 s0, 0 +; GFX1250-SDAG-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1250-SDAG-NEXT: s_cselect_b32 s1, s1, 0 +; GFX1250-SDAG-NEXT: v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v0, s0 +; GFX1250-SDAG-NEXT: v_mov_b32_e32 v1, s1 +; GFX1250-SDAG-NEXT: flat_store_b64 v2, v[0:1], s[2:3] +; GFX1250-SDAG-NEXT: s_endpgm +; +; GFX1250-GISEL-LABEL: barrier_to_generic: +; GFX1250-GISEL: ; %bb.0: +; GFX1250-GISEL-NEXT: global_wb +; GFX1250-GISEL-NEXT: v_nop +; GFX1250-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; GFX1250-GISEL-NEXT: s_clause 0x1 +; GFX1250-GISEL-NEXT: s_load_b32 s6, s[4:5], 0x0 nv +; GFX1250-GISEL-NEXT: s_load_b64 s[2:3], s[4:5], 0x8 nv +; GFX1250-GISEL-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1250-GISEL-NEXT: v_mov_b32_e32 v2, 0 +; GFX1250-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1250-GISEL-NEXT: s_add_co_u32 s0, s6, 0x802000 +; GFX1250-GISEL-NEXT: s_cmp_lg_u32 s6, 0 +; GFX1250-GISEL-NEXT: s_cselect_b64 s[0:1], s[0:1], 0 +; GFX1250-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1250-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1] +; GFX1250-GISEL-NEXT: flat_store_b64 v2, v[0:1], s[2:3] +; GFX1250-GISEL-NEXT: s_endpgm + %res = addrspacecast ptr addrspace(15) %bar to ptr + store ptr %res, ptr %out + ret void +} + +define amdgpu_kernel void @barrier_gv_to_generic(ptr %out) { +; GFX942-SDAG-LABEL: barrier_gv_to_generic: +; GFX942-SDAG: ; %bb.0: +; GFX942-SDAG-NEXT: s_load_dwordx2 s[2:3], s[4:5], 0x0 +; GFX942-SDAG-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX942-SDAG-NEXT: v_mov_b32_e32 v0, 0x802001 +; GFX942-SDAG-NEXT: v_mov_b32_e32 v1, s1 +; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-SDAG-NEXT: v_mov_b64_e32 v[2:3], s[2:3] +; GFX942-SDAG-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX942-SDAG-NEXT: s_endpgm +; +; GFX942-GISEL-LABEL: barrier_gv_to_generic: +; GFX942-GISEL: ; %bb.0: +; GFX942-GISEL-NEXT: s_load_dwordx2 s[2:3], s[4:5], 0x0 +; GFX942-GISEL-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX942-GISEL-NEXT: s_mov_b32 s0, 0x802001 +; GFX942-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1] +; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[2:3] +; GFX942-GISEL-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX942-GISEL-NEXT: s_endpgm +; +; GFX1030-SDAG-LABEL: barrier_gv_to_generic: +; GFX1030-SDAG: ; %bb.0: +; GFX1030-SDAG-NEXT: s_add_u32 s12, s12, s17 +; GFX1030-SDAG-NEXT: s_addc_u32 s13, s13, 0 +; GFX1030-SDAG-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_LO), s12 +; GFX1030-SDAG-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_HI), s13 +; GFX1030-SDAG-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x0 +; GFX1030-SDAG-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v0, 0x802001 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v1, s1 +; GFX1030-SDAG-NEXT: s_waitcnt lgkmcnt(0) +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v2, s2 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v3, s3 +; GFX1030-SDAG-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX1030-SDAG-NEXT: s_endpgm +; +; GFX1030-GISEL-LABEL: barrier_gv_to_generic: +; GFX1030-GISEL: ; %bb.0: +; GFX1030-GISEL-NEXT: s_add_u32 s12, s12, s17 +; GFX1030-GISEL-NEXT: s_addc_u32 s13, s13, 0 +; GFX1030-GISEL-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_LO), s12 +; GFX1030-GISEL-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_HI), s13 +; GFX1030-GISEL-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x0 +; GFX1030-GISEL-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1030-GISEL-NEXT: s_mov_b32 s0, 0x802001 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v1, s1 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v0, s0 +; GFX1030-GISEL-NEXT: s_waitcnt lgkmcnt(0) +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v2, s2 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v3, s3 +; GFX1030-GISEL-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX1030-GISEL-NEXT: s_endpgm +; +; GFX1200-SDAG-LABEL: barrier_gv_to_generic: +; GFX1200-SDAG: ; %bb.0: +; GFX1200-SDAG-NEXT: s_load_b64 s[2:3], s[4:5], 0x0 +; GFX1200-SDAG-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1200-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1200-SDAG-NEXT: v_dual_mov_b32 v0, 0x802001 :: v_dual_mov_b32 v1, s1 +; GFX1200-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1200-SDAG-NEXT: v_dual_mov_b32 v2, s2 :: v_dual_mov_b32 v3, s3 +; GFX1200-SDAG-NEXT: flat_store_b64 v[2:3], v[0:1] +; GFX1200-SDAG-NEXT: s_endpgm +; +; GFX1200-GISEL-LABEL: barrier_gv_to_generic: +; GFX1200-GISEL: ; %bb.0: +; GFX1200-GISEL-NEXT: s_load_b64 s[2:3], s[4:5], 0x0 +; GFX1200-GISEL-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1200-GISEL-NEXT: s_mov_b32 s0, 0x802001 +; GFX1200-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1200-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v0, s0 +; GFX1200-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1200-GISEL-NEXT: v_dual_mov_b32 v2, s2 :: v_dual_mov_b32 v3, s3 +; GFX1200-GISEL-NEXT: flat_store_b64 v[2:3], v[0:1] +; GFX1200-GISEL-NEXT: s_endpgm +; +; GFX1250-SDAG-LABEL: barrier_gv_to_generic: +; GFX1250-SDAG: ; %bb.0: +; GFX1250-SDAG-NEXT: global_wb +; GFX1250-SDAG-NEXT: v_nop +; GFX1250-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; GFX1250-SDAG-NEXT: s_load_b64 s[2:3], s[4:5], 0x0 nv +; GFX1250-SDAG-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1250-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1250-SDAG-NEXT: v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v1, s1 +; GFX1250-SDAG-NEXT: v_mov_b32_e32 v0, 0x802001 +; GFX1250-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1250-SDAG-NEXT: flat_store_b64 v2, v[0:1], s[2:3] +; GFX1250-SDAG-NEXT: s_endpgm +; +; GFX1250-GISEL-LABEL: barrier_gv_to_generic: +; GFX1250-GISEL: ; %bb.0: +; GFX1250-GISEL-NEXT: global_wb +; GFX1250-GISEL-NEXT: v_nop +; GFX1250-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; GFX1250-GISEL-NEXT: s_load_b64 s[2:3], s[4:5], 0x0 nv +; GFX1250-GISEL-NEXT: s_mov_b64 s[0:1], src_shared_base +; GFX1250-GISEL-NEXT: s_mov_b32 s0, 0x802001 +; GFX1250-GISEL-NEXT: v_mov_b32_e32 v2, 0 +; GFX1250-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1] +; GFX1250-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1250-GISEL-NEXT: flat_store_b64 v2, v[0:1], s[2:3] +; GFX1250-GISEL-NEXT: s_endpgm + %res = addrspacecast ptr addrspace(15) @bar to ptr + store ptr %res, ptr %out + ret void +} + + +define amdgpu_kernel void @barrier_null_to_generic(ptr %out) { +; GFX942-LABEL: barrier_null_to_generic: +; GFX942: ; %bb.0: +; GFX942-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0 +; GFX942-NEXT: v_mov_b64_e32 v[0:1], 0 +; GFX942-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[0:1] +; GFX942-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX942-NEXT: s_endpgm +; +; GFX1030-SDAG-LABEL: barrier_null_to_generic: +; GFX1030-SDAG: ; %bb.0: +; GFX1030-SDAG-NEXT: s_add_u32 s12, s12, s17 +; GFX1030-SDAG-NEXT: s_addc_u32 s13, s13, 0 +; GFX1030-SDAG-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_LO), s12 +; GFX1030-SDAG-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_HI), s13 +; GFX1030-SDAG-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v0, 0 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v1, v0 +; GFX1030-SDAG-NEXT: s_waitcnt lgkmcnt(0) +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v3, s1 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v2, s0 +; GFX1030-SDAG-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX1030-SDAG-NEXT: s_endpgm +; +; GFX1030-GISEL-LABEL: barrier_null_to_generic: +; GFX1030-GISEL: ; %bb.0: +; GFX1030-GISEL-NEXT: s_add_u32 s12, s12, s17 +; GFX1030-GISEL-NEXT: s_addc_u32 s13, s13, 0 +; GFX1030-GISEL-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_LO), s12 +; GFX1030-GISEL-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_HI), s13 +; GFX1030-GISEL-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v0, 0 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v1, 0 +; GFX1030-GISEL-NEXT: s_waitcnt lgkmcnt(0) +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v3, s1 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v2, s0 +; GFX1030-GISEL-NEXT: flat_store_dwordx2 v[2:3], v[0:1] +; GFX1030-GISEL-NEXT: s_endpgm +; +; GFX1200-SDAG-LABEL: barrier_null_to_generic: +; GFX1200-SDAG: ; %bb.0: +; GFX1200-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x0 +; GFX1200-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1200-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, s1 +; GFX1200-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX1200-SDAG-NEXT: v_dual_mov_b32 v1, v0 :: v_dual_mov_b32 v2, s0 +; GFX1200-SDAG-NEXT: flat_store_b64 v[2:3], v[0:1] +; GFX1200-SDAG-NEXT: s_endpgm +; +; GFX1200-GISEL-LABEL: barrier_null_to_generic: +; GFX1200-GISEL: ; %bb.0: +; GFX1200-GISEL-NEXT: s_load_b64 s[0:1], s[4:5], 0x0 +; GFX1200-GISEL-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, 0 +; GFX1200-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1200-GISEL-NEXT: v_dual_mov_b32 v3, s1 :: v_dual_mov_b32 v2, s0 +; GFX1200-GISEL-NEXT: flat_store_b64 v[2:3], v[0:1] +; GFX1200-GISEL-NEXT: s_endpgm +; +; GFX1250-LABEL: barrier_null_to_generic: +; GFX1250: ; %bb.0: +; GFX1250-NEXT: global_wb +; GFX1250-NEXT: v_nop +; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; GFX1250-NEXT: s_load_b64 s[0:1], s[4:5], 0x0 nv +; GFX1250-NEXT: v_mov_b64_e32 v[0:1], 0 +; GFX1250-NEXT: v_mov_b32_e32 v2, 0 +; GFX1250-NEXT: s_wait_kmcnt 0x0 +; GFX1250-NEXT: flat_store_b64 v2, v[0:1], s[0:1] +; GFX1250-NEXT: s_endpgm + %res = addrspacecast ptr addrspace(15) null to ptr + store ptr %res, ptr %out + ret void +} + +define amdgpu_kernel void @generic_to_barrier(ptr %generic, ptr %out) { +; GFX942-SDAG-LABEL: generic_to_barrier: +; GFX942-SDAG: ; %bb.0: +; GFX942-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0 +; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-SDAG-NEXT: v_mov_b32_e32 v0, s2 +; GFX942-SDAG-NEXT: s_add_i32 s2, s0, 0xff7fe000 +; GFX942-SDAG-NEXT: s_cmp_lg_u64 s[0:1], 0 +; GFX942-SDAG-NEXT: s_cselect_b32 s0, s2, 0 +; GFX942-SDAG-NEXT: v_mov_b32_e32 v1, s3 +; GFX942-SDAG-NEXT: v_mov_b32_e32 v2, s0 +; GFX942-SDAG-NEXT: flat_store_dword v[0:1], v2 +; GFX942-SDAG-NEXT: s_endpgm +; +; GFX942-GISEL-LABEL: generic_to_barrier: +; GFX942-GISEL: ; %bb.0: +; GFX942-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0 +; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-GISEL-NEXT: s_add_u32 s4, s0, 0xff7fe000 +; GFX942-GISEL-NEXT: s_cmp_lg_u64 s[0:1], 0 +; GFX942-GISEL-NEXT: s_cselect_b32 s0, s4, 0 +; GFX942-GISEL-NEXT: v_mov_b32_e32 v2, s0 +; GFX942-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[2:3] +; GFX942-GISEL-NEXT: flat_store_dword v[0:1], v2 +; GFX942-GISEL-NEXT: s_endpgm +; +; GFX1030-SDAG-LABEL: generic_to_barrier: +; GFX1030-SDAG: ; %bb.0: +; GFX1030-SDAG-NEXT: s_add_u32 s12, s12, s17 +; GFX1030-SDAG-NEXT: s_addc_u32 s13, s13, 0 +; GFX1030-SDAG-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_LO), s12 +; GFX1030-SDAG-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_HI), s13 +; GFX1030-SDAG-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0 +; GFX1030-SDAG-NEXT: s_waitcnt lgkmcnt(0) +; GFX1030-SDAG-NEXT: s_add_i32 s4, s0, 0xff7fe000 +; GFX1030-SDAG-NEXT: s_cmp_lg_u64 s[0:1], 0 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v0, s2 +; GFX1030-SDAG-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v1, s3 +; GFX1030-SDAG-NEXT: v_mov_b32_e32 v2, s0 +; GFX1030-SDAG-NEXT: flat_store_dword v[0:1], v2 +; GFX1030-SDAG-NEXT: s_endpgm +; +; GFX1030-GISEL-LABEL: generic_to_barrier: +; GFX1030-GISEL: ; %bb.0: +; GFX1030-GISEL-NEXT: s_add_u32 s12, s12, s17 +; GFX1030-GISEL-NEXT: s_addc_u32 s13, s13, 0 +; GFX1030-GISEL-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_LO), s12 +; GFX1030-GISEL-NEXT: s_setreg_b32 hwreg(HW_REG_FLAT_SCR_HI), s13 +; GFX1030-GISEL-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0 +; GFX1030-GISEL-NEXT: s_waitcnt lgkmcnt(0) +; GFX1030-GISEL-NEXT: s_add_u32 s4, s0, 0xff7fe000 +; GFX1030-GISEL-NEXT: s_cmp_lg_u64 s[0:1], 0 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v0, s2 +; GFX1030-GISEL-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v1, s3 +; GFX1030-GISEL-NEXT: v_mov_b32_e32 v2, s0 +; GFX1030-GISEL-NEXT: flat_store_dword v[0:1], v2 +; GFX1030-GISEL-NEXT: s_endpgm +; +; GFX1200-SDAG-LABEL: generic_to_barrier: +; GFX1200-SDAG: ; %bb.0: +; GFX1200-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x0 +; GFX1200-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1200-SDAG-NEXT: s_add_co_i32 s4, s0, 0xff7fe000 +; GFX1200-SDAG-NEXT: s_cmp_lg_u64 s[0:1], 0 +; GFX1200-SDAG-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3 +; GFX1200-SDAG-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1200-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1200-SDAG-NEXT: v_mov_b32_e32 v2, s0 +; GFX1200-SDAG-NEXT: flat_store_b32 v[0:1], v2 +; GFX1200-SDAG-NEXT: s_endpgm +; +; GFX1200-GISEL-LABEL: generic_to_barrier: +; GFX1200-GISEL: ; %bb.0: +; GFX1200-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x0 +; GFX1200-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1200-GISEL-NEXT: s_add_co_u32 s4, s0, 0xff7fe000 +; GFX1200-GISEL-NEXT: s_cmp_lg_u64 s[0:1], 0 +; GFX1200-GISEL-NEXT: v_mov_b32_e32 v0, s2 +; GFX1200-GISEL-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1200-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1200-GISEL-NEXT: v_dual_mov_b32 v1, s3 :: v_dual_mov_b32 v2, s0 +; GFX1200-GISEL-NEXT: flat_store_b32 v[0:1], v2 +; GFX1200-GISEL-NEXT: s_endpgm +; +; GFX1250-SDAG-LABEL: generic_to_barrier: +; GFX1250-SDAG: ; %bb.0: +; GFX1250-SDAG-NEXT: global_wb +; GFX1250-SDAG-NEXT: v_nop +; GFX1250-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; GFX1250-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x0 nv +; GFX1250-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1250-SDAG-NEXT: s_add_co_i32 s4, s0, 0xff7fe000 +; GFX1250-SDAG-NEXT: s_cmp_lg_u64 s[0:1], 0 +; GFX1250-SDAG-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1250-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1250-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0 +; GFX1250-SDAG-NEXT: flat_store_b32 v0, v1, s[2:3] +; GFX1250-SDAG-NEXT: s_endpgm +; +; GFX1250-GISEL-LABEL: generic_to_barrier: +; GFX1250-GISEL: ; %bb.0: +; GFX1250-GISEL-NEXT: global_wb +; GFX1250-GISEL-NEXT: v_nop +; GFX1250-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; GFX1250-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x0 nv +; GFX1250-GISEL-NEXT: v_mov_b32_e32 v1, 0 +; GFX1250-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1250-GISEL-NEXT: s_add_co_u32 s4, s0, 0xff7fe000 +; GFX1250-GISEL-NEXT: s_cmp_lg_u64 s[0:1], 0 +; GFX1250-GISEL-NEXT: s_cselect_b32 s0, s4, 0 +; GFX1250-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) +; GFX1250-GISEL-NEXT: v_mov_b32_e32 v0, s0 +; GFX1250-GISEL-NEXT: flat_store_b32 v1, v0, s[2:3] +; GFX1250-GISEL-NEXT: s_endpgm + %res = addrspacecast ptr %generic to ptr addrspace(15) + store ptr addrspace(15) %res, ptr %out + ret void +} +;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: +; GFX1030: {{.*}} +; GFX1200: {{.*}} diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-module-lds.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-module-lds.ll index 6d0ac0e8efda9..83fb380b655c2 100644 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-module-lds.ll +++ b/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-module-lds.ll @@ -6,28 +6,28 @@ ; where amdgpu-lower-module-lds pass runs in pipeline after amdgpu-lower-exec-sync pass. %class.ExpAmdWorkgroupWaveBarrier = type { target("amdgcn.named.barrier", 0) } -@bar2 = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison -@bar3 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison -@bar1 = internal addrspace(3) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison +@bar2 = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison +@bar3 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison +@bar1 = internal addrspace(15) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison @lds1 = internal addrspace(3) global [1 x i8] poison, align 4 ;. -; CHECK: @bar2 = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol [[META0:![0-9]+]] -; CHECK: @bar3 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META1:![0-9]+]] -; CHECK: @bar1 = internal addrspace(3) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison, !absolute_symbol [[META2:![0-9]+]] +; CHECK: @bar2 = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol [[META0:![0-9]+]] +; CHECK: @bar3 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META1:![0-9]+]] +; CHECK: @bar1 = internal addrspace(15) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison, !absolute_symbol [[META2:![0-9]+]] ; CHECK: @llvm.amdgcn.module.lds = internal addrspace(3) global %llvm.amdgcn.module.lds.t poison, align 4, !absolute_symbol [[META3:![0-9]+]] ; CHECK: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(3) @llvm.amdgcn.module.lds to ptr)], section "llvm.metadata" ;. define void @func1() #0 { ; CHECK-LABEL: define void @func1( ; CHECK-SAME: ) #[[ATTR0:[0-9]+]] { -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar3) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar3, i32 7) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar3) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar3, i32 7) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar3) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar3, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar3) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar3, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } @@ -35,14 +35,14 @@ define void @func1() #0 { define void @func2() #0 { ; CHECK-LABEL: define void @func2( ; CHECK-SAME: ) #[[ATTR0]] { -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK-NEXT: store i8 7, ptr addrspace(3) @llvm.amdgcn.module.lds, align 4 ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) store i8 7, ptr addrspace(3) @lds1, align 4 ret void @@ -52,20 +52,20 @@ define amdgpu_kernel void @kernel1() #0 { ; CHECK-LABEL: define amdgpu_kernel void @kernel1( ; CHECK-SAME: ) #[[ATTR1:[0-9]+]] { ; CHECK-NEXT: call void @llvm.donothing() [ "ExplicitUse"(ptr addrspace(3) @llvm.amdgcn.module.lds) ] -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 11) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 11) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) -; CHECK-NEXT: [[STATE:%.*]] = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) @bar1) +; CHECK-NEXT: [[STATE:%.*]] = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) @bar1) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier() ; CHECK-NEXT: call void @func1() ; CHECK-NEXT: call void @func2() ; CHECK-NEXT: store i8 9, ptr addrspace(3) @llvm.amdgcn.module.lds, align 4 ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 11) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 11) call void @llvm.amdgcn.s.barrier.wait(i16 1) - %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) @bar1) + %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) @bar1) call void @llvm.amdgcn.s.barrier() call void @func1() call void @func2() @@ -77,15 +77,15 @@ define amdgpu_kernel void @kernel2() #0 { ; CHECK-LABEL: define amdgpu_kernel void @kernel2( ; CHECK-SAME: ) #[[ATTR1]] { ; CHECK-NEXT: call void @llvm.donothing() [ "ExplicitUse"(ptr addrspace(3) @llvm.amdgcn.module.lds) ] -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 9) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 9) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK-NEXT: call void @func2() ; CHECK-NEXT: store i8 10, ptr addrspace(3) @llvm.amdgcn.module.lds, align 4 ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 9) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 9) call void @llvm.amdgcn.s.barrier.wait(i16 1) call void @func2() store i8 10, ptr addrspace(3) @lds1, align 4 @@ -95,13 +95,13 @@ define amdgpu_kernel void @kernel2() #0 { declare void @llvm.amdgcn.s.barrier() #1 declare void @llvm.amdgcn.s.barrier.wait(i16) #1 declare void @llvm.amdgcn.s.barrier.signal(i32) #1 -declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3), i32) #1 +declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15), i32) #1 declare i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32) #1 -declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(3), i32) #1 -declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(3)) #1 +declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(15), i32) #1 +declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(15)) #1 declare void @llvm.amdgcn.s.barrier.leave(i16) #1 -declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3)) #1 -declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3)) #1 +declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(15)) #1 +declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15)) #1 attributes #0 = { nounwind } attributes #1 = { convergent nounwind } @@ -113,8 +113,8 @@ attributes #2 = { nounwind readnone } ; CHECK: attributes #[[ATTR2:[0-9]+]] = { convergent nocallback nofree nounwind willreturn } ; CHECK: attributes #[[ATTR3:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(none) } ;. -; CHECK: [[META0]] = !{i32 8396880, i32 8396881} -; CHECK: [[META1]] = !{i32 8396912, i32 8396913} -; CHECK: [[META2]] = !{i32 8396816, i32 8396817} +; CHECK: [[META0]] = !{i32 5, i32 6} +; CHECK: [[META1]] = !{i32 7, i32 8} +; CHECK: [[META2]] = !{i32 1, i32 2} ; CHECK: [[META3]] = !{i32 0, i32 1} ;. diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-sw-lds.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-sw-lds.ll index 2af602d23af07..a2c8f5233e94b 100644 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-sw-lds.ll +++ b/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-sw-lds.ll @@ -4,24 +4,24 @@ ; Test to ensure that LDS variables like named barriers are lowered correctly in asan scenario, ; where amdgpu-sw-lower-lds pass runs in pipeline after amdgpu-lower-exec-sync pass. %class.ExpAmdWorkgroupWaveBarrier = type { target("amdgcn.named.barrier", 0) } -@bar2 = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison -@bar1 = internal addrspace(3) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison +@bar2 = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison +@bar1 = internal addrspace(15) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison @lds1 = internal addrspace(3) global [1 x i8] poison, align 4 ;. -; CHECK: @bar2 = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol [[META0:![0-9]+]] -; CHECK: @bar1 = internal addrspace(3) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison, !absolute_symbol [[META1:![0-9]+]] +; CHECK: @bar2 = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol [[META0:![0-9]+]] +; CHECK: @bar1 = internal addrspace(15) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison, !absolute_symbol [[META1:![0-9]+]] ; define void @bar() #0 { ; CHECK-LABEL: define void @bar( ; CHECK-SAME: ) #[[ATTR0:[0-9]+]] { -; CHECK: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) -; CHECK: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) +; CHECK: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) +; CHECK: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) ; CHECK: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK: store i8 7, ptr addrspace(1) {{.*}}, align 4 ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) store i8 7, ptr addrspace(3) @lds1, align 4 ret void @@ -32,32 +32,21 @@ define amdgpu_kernel void @barkernel() #0 { ; CHECK-SAME: ) #[[ATTR1:[0-9]+]] !llvm.amdgcn.lds.kernel.id [[META4:![0-9]+]] { ; CHECK: {{.*}} = call i64 @__asan_malloc_impl(i64 {{.*}}, i64 {{.*}}) ; CHECK: call void @llvm.amdgcn.s.barrier() -; CHECK: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) -; CHECK: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 9) +; CHECK: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) +; CHECK: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 9) ; CHECK: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK: call void @bar() ; CHECK: store i8 10, ptr addrspace(1) {{.*}}, align 4 ; CHECK: call void @__asan_free_impl(i64 {{.*}}, i64 {{.*}}) ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 9) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 9) call void @llvm.amdgcn.s.barrier.wait(i16 1) call void @bar() store i8 10, ptr addrspace(3) @lds1, align 4 ret void } -declare void @llvm.amdgcn.s.barrier() #1 -declare void @llvm.amdgcn.s.barrier.wait(i16) #1 -declare void @llvm.amdgcn.s.barrier.signal(i32) #1 -declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3), i32) #1 -declare i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32) #1 -declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(3), i32) #1 -declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(3)) #1 -declare void @llvm.amdgcn.s.barrier.leave(i16) #1 -declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3)) #1 -declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3)) #1 - attributes #0 = { nounwind sanitize_address } attributes #1 = { convergent nounwind } attributes #2 = { nounwind readnone } @@ -68,6 +57,6 @@ attributes #2 = { nounwind readnone } ; CHECK: attributes #[[ATTR0]] = { nounwind sanitize_address } ; CHECK: attributes #[[ATTR1]] = { nounwind sanitize_address "amdgpu-lds-size"="8" } ;. -; CHECK: [[META0]] = !{i32 8396880, i32 8396881} -; CHECK: [[META1]] = !{i32 8396816, i32 8396817} +; CHECK: [[META0]] = !{i32 5, i32 6} +; CHECK: [[META1]] = !{i32 1, i32 2} ;. diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync.ll index f75bf87559955..5b1da24f48aa3 100644 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync.ll +++ b/llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync.ll @@ -4,37 +4,37 @@ %class.ExpAmdWorkgroupWaveBarrier = type { target("amdgcn.named.barrier", 0) } -@bar2 = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison -@bar3 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison -@bar1 = internal addrspace(3) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison +@bar2 = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison +@bar3 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison +@bar1 = internal addrspace(15) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison ;. -; CHECK: @bar2 = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol [[META0:![0-9]+]] -; CHECK: @bar3 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META1:![0-9]+]] -; CHECK: @bar1 = internal addrspace(3) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison, !absolute_symbol [[META2:![0-9]+]] +; CHECK: @bar2 = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol [[META0:![0-9]+]] +; CHECK: @bar3 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META1:![0-9]+]] +; CHECK: @bar1 = internal addrspace(15) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison, !absolute_symbol [[META2:![0-9]+]] ;. define void @func1() { ; CHECK-LABEL: define void @func1() { -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar3) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar3, i32 7) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar3) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar3, i32 7) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar3) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar3, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar3) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar3, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } define void @func2() { ; CHECK-LABEL: define void @func2() { -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } @@ -42,19 +42,19 @@ define void @func2() { define amdgpu_kernel void @kernel1() #0 { ; CHECK-LABEL: define amdgpu_kernel void @kernel1( ; CHECK-SAME: ) #[[ATTR0:[0-9]+]] { -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 11) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 11) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) -; CHECK-NEXT: [[STATE:%.*]] = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) @bar1) +; CHECK-NEXT: [[STATE:%.*]] = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) @bar1) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier() ; CHECK-NEXT: call void @func1() ; CHECK-NEXT: call void @func2() ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 11) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 11) call void @llvm.amdgcn.s.barrier.wait(i16 1) - %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) @bar1) + %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) @bar1) call void @llvm.amdgcn.s.barrier() call void @func1() call void @func2() @@ -64,14 +64,14 @@ define amdgpu_kernel void @kernel1() #0 { define amdgpu_kernel void @kernel2() #0 { ; CHECK-LABEL: define amdgpu_kernel void @kernel2( ; CHECK-SAME: ) #[[ATTR0]] { -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 9) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 9) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK-NEXT: call void @func2() ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 9) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 9) call void @llvm.amdgcn.s.barrier.wait(i16 1) call void @func2() @@ -81,13 +81,13 @@ define amdgpu_kernel void @kernel2() #0 { declare void @llvm.amdgcn.s.barrier() #1 declare void @llvm.amdgcn.s.barrier.wait(i16) #1 declare void @llvm.amdgcn.s.barrier.signal(i32) #1 -declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3), i32) #1 +declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15), i32) #1 declare i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32) #1 -declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(3), i32) #1 -declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(3)) #1 +declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(15), i32) #1 +declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(15)) #1 declare void @llvm.amdgcn.s.barrier.leave(i16) #1 -declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3)) #1 -declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3)) #1 +declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(15)) #1 +declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15)) #1 attributes #0 = { nounwind } attributes #1 = { convergent nounwind } @@ -96,7 +96,7 @@ attributes #2 = { nounwind readnone } ; CHECK: attributes #[[ATTR0]] = { nounwind } ; CHECK: attributes #[[ATTR1:[0-9]+]] = { convergent nocallback nofree nounwind willreturn } ;. -; CHECK: [[META0]] = !{i32 8396880, i32 8396881} -; CHECK: [[META1]] = !{i32 8396912, i32 8396913} -; CHECK: [[META2]] = !{i32 8396816, i32 8396817} +; CHECK: [[META0]] = !{i32 5, i32 6} +; CHECK: [[META1]] = !{i32 7, i32 8} +; CHECK: [[META2]] = !{i32 1, i32 2} ;. diff --git a/llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll b/llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll index f51247ba10964..90d32d0cc48e9 100644 --- a/llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll +++ b/llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll @@ -489,8 +489,8 @@ attributes #1 = { nounwind } ; HSA: attributes #[[ATTR13]] = { nounwind "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-no-wwm" } ; HSA: attributes #[[ATTR14]] = { nounwind "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-no-wwm" } ;. -; HSA: [[META0]] = !{i32 1, i32 3, i32 4, i32 10} -; HSA: [[META1]] = !{i32 1, i32 5, i32 6, i32 10} -; HSA: [[META2]] = !{i32 2, i32 10} -; HSA: [[META3]] = !{i32 1, i32 4, i32 5, i32 10} +; HSA: [[META0]] = !{i32 1, i32 3, i32 4, i32 16} +; HSA: [[META1]] = !{i32 1, i32 5, i32 6, i32 16} +; HSA: [[META2]] = !{i32 2, i32 16} +; HSA: [[META3]] = !{i32 1, i32 4, i32 5, i32 16} ;. diff --git a/llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit-undefined-behavior.ll b/llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit-undefined-behavior.ll index d3e0f581dc03f..f998807c217b1 100644 --- a/llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit-undefined-behavior.ll +++ b/llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit-undefined-behavior.ll @@ -153,7 +153,7 @@ attributes #0 = { "amdgpu-no-flat-scratch-init" } ; GFX10: attributes #[[ATTR0]] = { "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-no-wwm" } ; GFX10: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) } ;. -; GFX9: [[META0]] = !{i32 1, i32 5, i32 6, i32 10} +; GFX9: [[META0]] = !{i32 1, i32 5, i32 6, i32 16} ;. -; GFX10: [[META0]] = !{i32 1, i32 5, i32 6, i32 10} +; GFX10: [[META0]] = !{i32 1, i32 5, i32 6, i32 16} ;. diff --git a/llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit.ll b/llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit.ll index 11fd660813a13..59f644833ea89 100644 --- a/llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit.ll +++ b/llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit.ll @@ -867,15 +867,15 @@ define amdgpu_kernel void @with_inline_asm() { ; GFX10: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) } ; GFX10: attributes #[[ATTR3]] = { "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-no-wwm" } ;. -; GFX9: [[META0]] = !{i32 2, i32 10} -; GFX9: [[META1]] = !{i32 1, i32 2, i32 3, i32 10} -; GFX9: [[META2]] = !{i32 1, i32 3, i32 4, i32 10} -; GFX9: [[META3]] = !{i32 1, i32 4, i32 5, i32 10} -; GFX9: [[META4]] = !{i32 1, i32 5, i32 6, i32 10} +; GFX9: [[META0]] = !{i32 2, i32 16} +; GFX9: [[META1]] = !{i32 1, i32 2, i32 3, i32 16} +; GFX9: [[META2]] = !{i32 1, i32 3, i32 4, i32 16} +; GFX9: [[META3]] = !{i32 1, i32 4, i32 5, i32 16} +; GFX9: [[META4]] = !{i32 1, i32 5, i32 6, i32 16} ;. -; GFX10: [[META0]] = !{i32 2, i32 10} -; GFX10: [[META1]] = !{i32 1, i32 2, i32 3, i32 10} -; GFX10: [[META2]] = !{i32 1, i32 3, i32 4, i32 10} -; GFX10: [[META3]] = !{i32 1, i32 4, i32 5, i32 10} -; GFX10: [[META4]] = !{i32 1, i32 5, i32 6, i32 10} +; GFX10: [[META0]] = !{i32 2, i32 16} +; GFX10: [[META1]] = !{i32 1, i32 2, i32 3, i32 16} +; GFX10: [[META2]] = !{i32 1, i32 3, i32 4, i32 16} +; GFX10: [[META3]] = !{i32 1, i32 4, i32 5, i32 16} +; GFX10: [[META4]] = !{i32 1, i32 5, i32 6, i32 16} ;. diff --git a/llvm/test/CodeGen/AMDGPU/attributor-noalias-addrspace.ll b/llvm/test/CodeGen/AMDGPU/attributor-noalias-addrspace.ll index f9edbd070ae7c..b8458981b1dda 100644 --- a/llvm/test/CodeGen/AMDGPU/attributor-noalias-addrspace.ll +++ b/llvm/test/CodeGen/AMDGPU/attributor-noalias-addrspace.ll @@ -633,7 +633,7 @@ define amdgpu_kernel void @no_alias_addr_space_has_meta(ptr addrspace(3) %sptr, !0 = !{i32 2, i32 3, i32 4, i32 10} ;. -; CHECK: [[META0]] = !{i32 2, i32 3, i32 4, i32 5, i32 6, i32 10} -; CHECK: [[META1]] = !{i32 2, i32 3, i32 5, i32 10} +; CHECK: [[META0]] = !{i32 2, i32 3, i32 4, i32 5, i32 6, i32 16} +; CHECK: [[META1]] = !{i32 2, i32 3, i32 5, i32 16} ; CHECK: [[META2]] = !{i32 2, i32 3, i32 4, i32 10} ;. diff --git a/llvm/test/CodeGen/AMDGPU/barrier-addrspace-dereference.ll b/llvm/test/CodeGen/AMDGPU/barrier-addrspace-dereference.ll new file mode 100644 index 0000000000000..b58bc8551cc82 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/barrier-addrspace-dereference.ll @@ -0,0 +1,16 @@ +; Check we cannot dereference a barrier GV. + +; RUN: not --crash llc -O0 -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s 2>&1 | FileCheck -check-prefixes=DAGISEL %s +; RUN: not llc -O0 -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s 2>&1 | FileCheck -check-prefixes=GISEL %s + +; TODO: It'd be nicer to have a Verifier diagnostic for this. + +; DAGISEL: LLVM ERROR: {{.*}} store<(store (s32) into @bar, addrspace 15)> +; GISEL: LLVM ERROR: {{.*}} G_LOAD %6:sgpr(p15) :: (load (i32) from @bar, addrspace 15) (in function: func1) +@bar = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison + +define amdgpu_kernel void @func1() { + %val = load i32, ptr addrspace(15) @bar + store i32 %val, ptr addrspace(15) @bar + ret void +} diff --git a/llvm/test/CodeGen/AMDGPU/lds-link-time-codegen-named-barrier.ll b/llvm/test/CodeGen/AMDGPU/lds-link-time-codegen-named-barrier.ll index b9313088f41b6..9189ff75531a3 100644 --- a/llvm/test/CodeGen/AMDGPU/lds-link-time-codegen-named-barrier.ll +++ b/llvm/test/CodeGen/AMDGPU/lds-link-time-codegen-named-barrier.ll @@ -7,10 +7,9 @@ ; 3. group_segment_fixed_size = 0 (linker patches it) ; 4. Named barrier is emitted as an SHN_AMDGPU_LDS symbol (.amdgpu_lds) -@bar = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison +@bar = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison ; CHECK-LABEL: kernel: -; CHECK: s_mov_b32 s{{[0-9]+}}, __amdgpu_named_barrier.bar{{[^ @]*}}@abs32@lo ; CHECK: s_barrier_join m0 ; CHECK: s_barrier_signal m0 ; CHECK: s_barrier_wait 1 @@ -26,8 +25,6 @@ ; CHECK: .amdgpu_call helper ; CHECK: .end_amdgpu_info -; CHECK: .amdgpu_lds __amdgpu_named_barrier.bar{{[^ ,]*}}, 32, 4 - ; ELF: Section { ; ELF: Name: .amdgpu.info ; ELF: Type: SHT_PROGBITS @@ -39,16 +36,16 @@ ; ELF-DAG: R_AMDGPU_ABS64 helper define amdgpu_kernel void @kernel() { - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar, i32 3) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar, i32 3) call void @llvm.amdgcn.s.barrier.wait(i16 1) call void @helper() ret void } declare void @helper() -declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(3)) #0 -declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3), i32) #0 +declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(15)) #0 +declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15), i32) #0 declare void @llvm.amdgcn.s.barrier.wait(i16) #0 attributes #0 = { convergent nounwind } diff --git a/llvm/test/CodeGen/AMDGPU/lds-link-time-named-barrier.ll b/llvm/test/CodeGen/AMDGPU/lds-link-time-named-barrier.ll index 75359d575575e..2441d0c7188c9 100644 --- a/llvm/test/CodeGen/AMDGPU/lds-link-time-named-barrier.ll +++ b/llvm/test/CodeGen/AMDGPU/lds-link-time-named-barrier.ll @@ -6,21 +6,21 @@ ; 2. AMDGPULowerModuleLDS does not handle named barriers at all ; 3. amdgpu.lds.uses does NOT contain barrier entries -@bar = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison +@bar = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison @lds = internal addrspace(3) global [4 x i32] poison, align 4 ; Internal named barrier becomes external with a module-unique hash suffix. -; CHECK: @[[BAR:__amdgpu_named_barrier\.bar\.[a-f0-9]+]] = external dso_local addrspace(3) global target("amdgcn.named.barrier", 0) +; CHECK: @[[BAR:__amdgpu_named_barrier\.bar\.[a-f0-9]+]] = external dso_local addrspace(15) global target("amdgcn.named.barrier", 0) ; CHECK-NOT: !absolute_symbol ; Regular LDS is packed into the per-function struct (external, for linker). ; CHECK: @__amdgpu_lds.kernel = external dso_local addrspace(3) global %__amdgpu_lds.kernel.t, align 16 define amdgpu_kernel void @kernel(i32 %idx) { ; CHECK-LABEL: define amdgpu_kernel void @kernel( -; CHECK: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @[[BAR]]) -; CHECK: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @[[BAR]], i32 3) - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar, i32 3) +; CHECK: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @[[BAR]]) +; CHECK: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @[[BAR]], i32 3) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar, i32 3) call void @llvm.amdgcn.s.barrier.wait(i16 1) %gep = getelementptr [4 x i32], ptr addrspace(3) @lds, i32 0, i32 %idx store i32 42, ptr addrspace(3) %gep, align 4 @@ -29,7 +29,7 @@ define amdgpu_kernel void @kernel(i32 %idx) { ; Named barrier metadata: (barrier_sym, func1, ...) -- emitted by ExecSync. ; CHECK-DAG: !amdgpu.named_barrier.uses = !{[[BAR_MD:![0-9]+]]} -; CHECK-DAG: [[BAR_MD]] = !{ptr addrspace(3) @[[BAR]], ptr @kernel} +; CHECK-DAG: [[BAR_MD]] = !{ptr addrspace(15) @[[BAR]], ptr @kernel} ; LDS metadata must have exactly one entry (the LDS struct), no barrier entries. ; CHECK-DAG: !amdgpu.lds.uses = !{[[LDS_MD:![0-9]+]]} ; CHECK-DAG: [[LDS_MD]] = !{ptr @kernel, ptr addrspace(3) @__amdgpu_lds.kernel} diff --git a/llvm/test/CodeGen/AMDGPU/null-named-barrier-gv.ll b/llvm/test/CodeGen/AMDGPU/null-named-barrier-gv.ll new file mode 100644 index 0000000000000..44536a9f3133f --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/null-named-barrier-gv.ll @@ -0,0 +1,31 @@ +; RUN: split-file %s %t + +; RUN: not --crash llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -o - %t/null-named-barrier-kernel.ll 2>&1 | FileCheck %s +; RUN: not --crash llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -o - %t/null-named-barrier-kernel.ll 2>&1 | FileCheck %s + +; RUN: not --crash llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -o - %t/null-named-barrier-func.ll 2>&1 | FileCheck %s +; RUN: not --crash llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -o - %t/null-named-barrier-func.ll 2>&1 | FileCheck %s + +; CHECK: named barrier global variable 'bar' has a NULL address, which is not supported + +;--- null-named-barrier-kernel.ll + +@bar = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol !0 + +define amdgpu_kernel void @func1() { + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar) + ret void +} + +!0 = !{ i32 0, i32 1 } + +;--- null-named-barrier-func.ll + +@bar = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol !0 + +define void @func1() { + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar) + ret void +} + +!0 = !{ i32 0, i32 1 } diff --git a/llvm/test/CodeGen/AMDGPU/s-barrier-id-allocation.ll b/llvm/test/CodeGen/AMDGPU/s-barrier-id-allocation.ll index cbaf6cdc1747d..986987d3e71de 100644 --- a/llvm/test/CodeGen/AMDGPU/s-barrier-id-allocation.ll +++ b/llvm/test/CodeGen/AMDGPU/s-barrier-id-allocation.ll @@ -1,42 +1,42 @@ ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --check-globals all --version 6 ; RUN: opt -S -mtriple=amdgpu-- -passes=amdgpu-lower-exec-sync < %s 2>&1 | FileCheck %s -@bar = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison -@bar2 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison +@bar = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison +@bar2 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison ;. -; CHECK: @bar = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META0:![0-9]+]] -; CHECK: @bar2 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META0]] +; CHECK: @bar = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META0:![0-9]+]] +; CHECK: @bar2 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META0]] ;. define void @func1() { ; CHECK-LABEL: define void @func1() { -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar, i32 7) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar, i32 7) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } define void @func2() { ; CHECK-LABEL: define void @func2() { -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) -; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) +; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) ; CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1) ; CHECK-NEXT: ret void ; - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } -define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) { +define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(15) %in) { ; CHECK-LABEL: define amdgpu_kernel void @kernel1( -; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(3) [[IN:%.*]]) { +; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(15) [[IN:%.*]]) { ; CHECK-NEXT: call void @func1() ; CHECK-NEXT: [[STATE3:%.*]] = call i32 @llvm.amdgcn.s.get.barrier.state(i32 -1) ; CHECK-NEXT: ret void @@ -46,9 +46,9 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ret void } -define amdgpu_kernel void @kernel2(ptr addrspace(1) %out, ptr addrspace(3) %in) { +define amdgpu_kernel void @kernel2(ptr addrspace(1) %out, ptr addrspace(15) %in) { ; CHECK-LABEL: define amdgpu_kernel void @kernel2( -; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(3) [[IN:%.*]]) { +; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(15) [[IN:%.*]]) { ; CHECK-NEXT: call void @func1() ; CHECK-NEXT: ret void ; @@ -56,9 +56,9 @@ define amdgpu_kernel void @kernel2(ptr addrspace(1) %out, ptr addrspace(3) %in) ret void } -define amdgpu_kernel void @kernel3(ptr addrspace(1) %out, ptr addrspace(3) %in) { +define amdgpu_kernel void @kernel3(ptr addrspace(1) %out, ptr addrspace(15) %in) { ; CHECK-LABEL: define amdgpu_kernel void @kernel3( -; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(3) [[IN:%.*]]) { +; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(15) [[IN:%.*]]) { ; CHECK-NEXT: call void @func2() ; CHECK-NEXT: [[STATE3:%.*]] = call i32 @llvm.amdgcn.s.get.barrier.state(i32 -1) ; CHECK-NEXT: ret void @@ -68,9 +68,9 @@ define amdgpu_kernel void @kernel3(ptr addrspace(1) %out, ptr addrspace(3) %in) ret void } -define amdgpu_kernel void @kernel4(ptr addrspace(1) %out, ptr addrspace(3) %in) { +define amdgpu_kernel void @kernel4(ptr addrspace(1) %out, ptr addrspace(15) %in) { ; CHECK-LABEL: define amdgpu_kernel void @kernel4( -; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(3) [[IN:%.*]]) { +; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(15) [[IN:%.*]]) { ; CHECK-NEXT: call void @func2() ; CHECK-NEXT: ret void ; @@ -82,5 +82,5 @@ define amdgpu_kernel void @kernel4(ptr addrspace(1) %out, ptr addrspace(3) %in) ;. ; CHECK: attributes #[[ATTR0:[0-9]+]] = { convergent nocallback nofree nounwind willreturn } ;. -; CHECK: [[META0]] = !{i32 8396816, i32 8396817} +; CHECK: [[META0]] = !{i32 1, i32 2} ;. diff --git a/llvm/test/CodeGen/AMDGPU/s-barrier-lowering-bad-absolute-symbol.ll b/llvm/test/CodeGen/AMDGPU/s-barrier-lowering-bad-absolute-symbol.ll new file mode 100644 index 0000000000000..acab496bb1971 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/s-barrier-lowering-bad-absolute-symbol.ll @@ -0,0 +1,16 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: not --crash llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 < %s 2>&1 | FileCheck %s +; RUN: not --crash llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 < %s 2>&1 | FileCheck %s + +; The absolute_address of the GV can never be null. + +; CHECK: LLVM ERROR: named barrier global variable 'bar' has a NULL address, which is not supported + +@bar = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol !0 + +define void @func() { + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar) + ret void +} + +!0 = !{i32 0, i32 1} diff --git a/llvm/test/CodeGen/AMDGPU/s-barrier-lowering-wrong-gv-signature.ll b/llvm/test/CodeGen/AMDGPU/s-barrier-lowering-wrong-gv-signature.ll new file mode 100644 index 0000000000000..ad3f094132e13 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/s-barrier-lowering-wrong-gv-signature.ll @@ -0,0 +1,27 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 < %s 2>&1 | FileCheck %s +; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 < %s 2>&1 | FileCheck %s + +; Check what happens when the type or the AS of the barrier GV is wrong. +; Using such a GV in a barrier intrinsic would be UB of course, but we should not crash. + +; @addrspacecasted doesn't trip up on this because we are not doing any unsupported +; operation. +; +; CHECK: in function wrong_type void (): unsupported use of BARRIER address space + +@bar = internal global target("amdgcn.named.barrier", 0) poison +@bar2 = internal addrspace(15) global i32 poison + +define void @addrspacecasted() { + %bar.ascast = addrspacecast ptr @bar to ptr addrspace(15) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) %bar.ascast) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %bar.ascast, i32 7) + ret void +} + +define void @wrong_type() { + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) + ret void +} diff --git a/llvm/test/CodeGen/AMDGPU/s-barrier-lowering.ll b/llvm/test/CodeGen/AMDGPU/s-barrier-lowering.ll index 51ff7a1ddb98c..c17d54b2c1c19 100644 --- a/llvm/test/CodeGen/AMDGPU/s-barrier-lowering.ll +++ b/llvm/test/CodeGen/AMDGPU/s-barrier-lowering.ll @@ -3,26 +3,30 @@ %class.ExpAmdWorkgroupWaveBarrier = type { target("amdgcn.named.barrier", 0) } -@bar2 = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison -@bar3 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison -@bar1 = internal addrspace(3) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison +@bar2 = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison +@bar3 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison +@bar1 = internal addrspace(15) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison -; CHECK: @bar2 = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol !0 -; CHECK-NEXT: @bar3 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol !1 -; CHECK-NEXT: @bar1 = internal addrspace(3) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison, !absolute_symbol !2 +; Test using the workgroup barrier with the GV. +@wgbarr = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol !0 + +; CHECK: @bar2 = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol [[META0:![0-9]+]] +; CHECK-NEXT: @bar3 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META1:![0-9]+]] +; CHECK-NEXT: @bar1 = internal addrspace(15) global [4 x %class.ExpAmdWorkgroupWaveBarrier] poison, !absolute_symbol [[META2:![0-9]+]] +; CHECK-NEXT: @wgbarr = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol [[META3:![0-9]+]] ; SOUT: .set .Lfunc1.num_named_barrier, 7 define void @func1() { - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar3) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar3, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar3) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar3, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } ; SOUT: .set .Lfunc2.num_named_barrier, 6 define void @func2() { - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } @@ -30,44 +34,38 @@ define void @func2() { ; SOUT: .amdhsa_named_barrier_count 2 ; SOUT: .set .Lkernel1.num_named_barrier, max(4, .Lfunc1.num_named_barrier, .Lfunc2.num_named_barrier) define amdgpu_kernel void @kernel1() #0 { -; CHECK-DAG: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 11) - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 11) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 11) call void @llvm.amdgcn.s.barrier.wait(i16 1) - %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) @bar1) + %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) @bar1) call void @llvm.amdgcn.s.barrier() call void @func1() call void @func2() ret void } -; SOUT: .amdhsa_named_barrier_count 2 +; SOUT: .amdhsa_kernel kernel2 +; SOUT: .amdhsa_named_barrier_count 2 ; SOUT: .set .Lkernel2.num_named_barrier, max(4, .Lfunc2.num_named_barrier) define amdgpu_kernel void @kernel2() #0 { -; CHECK-DAG: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 9) - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar1) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar1, i32 9) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar1) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar1, i32 9) call void @llvm.amdgcn.s.barrier.wait(i16 1) call void @func2() ret void } -declare void @llvm.amdgcn.s.barrier() #1 -declare void @llvm.amdgcn.s.barrier.wait(i16) #1 -declare void @llvm.amdgcn.s.barrier.signal(i32) #1 -declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3), i32) #1 -declare i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32) #1 -declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(3), i32) #1 -declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(3)) #1 -declare void @llvm.amdgcn.s.barrier.leave(i16) #1 -declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3)) #1 -declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3)) #1 +; SOUT: .amdhsa_kernel wgbarr_as_gv +; SOUT: .amdhsa_named_barrier_count 0 +define amdgpu_kernel void @wgbarr_as_gv() { + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @wgbarr, i32 7) + call void @llvm.amdgcn.s.barrier.wait(i16 -1) + ret void +} attributes #0 = { nounwind } attributes #1 = { convergent nounwind } attributes #2 = { nounwind readnone } -; CHECK: !0 = !{i32 8396880, i32 8396881} -; CHECK-NEXT: !1 = !{i32 8396912, i32 8396913} -; CHECK-NEXT: !2 = !{i32 8396816, i32 8396817} +!0 = !{i32 -1, i32 0} 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 65a6f150ab79a..135b652607788 100644 --- a/llvm/test/CodeGen/AMDGPU/s-barrier-signal-var-gep.ll +++ b/llvm/test/CodeGen/AMDGPU/s-barrier-signal-var-gep.ll @@ -4,7 +4,7 @@ ; RUN: llc -global-isel=0 -amdgpu-enable-object-linking -mtriple=amdgpu12.50-amd-amdhsa < %s | FileCheck -check-prefixes=CHECK-OBJ,CHECK-OBJ-SDAG %s ; RUN: llc -global-isel=1 -amdgpu-enable-object-linking -mtriple=amdgpu12.50-amd-amdhsa < %s | FileCheck -check-prefixes=CHECK-OBJ,CHECK-OBJ-GISEL %s -@bars = internal addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] poison +@bars = internal addrspace(15) global [2 x target("amdgcn.named.barrier", 0)] poison, !absolute_symbol !0 ; A constant-offset GEP into an array of named barriers (&bars[1]) is a ; compile-time-constant barrier address and should select the immediate @@ -18,9 +18,9 @@ define amdgpu_kernel void @signal_var_bar0() { ; CHECK-NEXT: global_wb ; CHECK-NEXT: v_nop ; CHECK-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; CHECK-NEXT: s_mov_b32 m0, 0x100001 +; CHECK-NEXT: s_mov_b32 m0, 0x10003f ; CHECK-NEXT: s_barrier_init m0 -; CHECK-NEXT: s_barrier_signal 1 +; CHECK-NEXT: s_barrier_signal 63 ; CHECK-NEXT: s_barrier_wait 0 ; CHECK-NEXT: s_endpgm ; @@ -29,13 +29,11 @@ define amdgpu_kernel void @signal_var_bar0() { ; CHECK-OBJ-SDAG-NEXT: global_wb ; CHECK-OBJ-SDAG-NEXT: v_nop ; 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_b32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo -; CHECK-OBJ-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; CHECK-OBJ-SDAG-NEXT: s_bfe_u32 s0, s0, 0x60004 +; CHECK-OBJ-SDAG-NEXT: s_and_b32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@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 -; CHECK-OBJ-SDAG-NEXT: s_mov_b32 m0, s0 -; CHECK-OBJ-SDAG-NEXT: s_barrier_signal m0 +; CHECK-OBJ-SDAG-NEXT: s_barrier_signal 63 ; CHECK-OBJ-SDAG-NEXT: s_barrier_wait 0 ; CHECK-OBJ-SDAG-NEXT: s_endpgm ; @@ -44,41 +42,49 @@ define amdgpu_kernel void @signal_var_bar0() { ; CHECK-OBJ-GISEL-NEXT: global_wb ; CHECK-OBJ-GISEL-NEXT: v_nop ; 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_lshr_b32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo, 4 -; 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_or_b32 m0, s0, 0x100000 +; CHECK-OBJ-GISEL-NEXT: s_mov_b32 m0, 0x10003f ; CHECK-OBJ-GISEL-NEXT: s_barrier_init m0 -; CHECK-OBJ-GISEL-NEXT: s_mov_b32 m0, s0 -; CHECK-OBJ-GISEL-NEXT: s_barrier_signal m0 +; CHECK-OBJ-GISEL-NEXT: s_barrier_signal 63 ; CHECK-OBJ-GISEL-NEXT: s_barrier_wait 0 ; CHECK-OBJ-GISEL-NEXT: s_endpgm - call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) @bars, i32 16) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bars, i32 0) + call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) @bars, i32 16) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bars, i32 0) call void @llvm.amdgcn.s.barrier.wait(i16 0) ret void } define amdgpu_kernel void @signal_var_bar1() { -; CHECK-LABEL: signal_var_bar1: -; CHECK: ; %bb.0: -; CHECK-NEXT: global_wb -; CHECK-NEXT: v_nop -; CHECK-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; CHECK-NEXT: s_mov_b32 m0, 0x100002 -; CHECK-NEXT: s_barrier_init m0 -; CHECK-NEXT: s_barrier_signal 2 -; CHECK-NEXT: s_barrier_wait 1 -; CHECK-NEXT: s_endpgm +; CHECK-SDAG-LABEL: signal_var_bar1: +; CHECK-SDAG: ; %bb.0: +; CHECK-SDAG-NEXT: global_wb +; CHECK-SDAG-NEXT: v_nop +; 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_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_bar1: +; CHECK-GISEL: ; %bb.0: +; CHECK-GISEL-NEXT: global_wb +; CHECK-GISEL-NEXT: v_nop +; 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_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: global_wb ; CHECK-OBJ-SDAG-NEXT: v_nop ; 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_b32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo+16 +; 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_bfe_u32 s0, s0, 0x60004 +; 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 @@ -91,20 +97,14 @@ define amdgpu_kernel void @signal_var_bar1() { ; CHECK-OBJ-GISEL-NEXT: global_wb ; CHECK-OBJ-GISEL-NEXT: v_nop ; 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_add_co_u32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo, 16 -; CHECK-OBJ-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; CHECK-OBJ-GISEL-NEXT: s_lshr_b32 s0, s0, 4 -; 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_mov_b32 m0, 0x10000f ; CHECK-OBJ-GISEL-NEXT: s_barrier_init m0 -; CHECK-OBJ-GISEL-NEXT: s_mov_b32 m0, s0 -; CHECK-OBJ-GISEL-NEXT: s_barrier_signal 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(3) @bars, i32 0, i32 1 - call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) %p1, i32 16) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) %p1, i32 0) + %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 } @@ -114,25 +114,37 @@ define amdgpu_kernel void @signal_var_bar1() { ; offset stays within barrier 0 and selects the same barrier as &bars[0]. define amdgpu_kernel void @signal_var_misaligned() { -; CHECK-LABEL: signal_var_misaligned: -; CHECK: ; %bb.0: -; CHECK-NEXT: global_wb -; CHECK-NEXT: v_nop -; CHECK-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 -; CHECK-NEXT: s_mov_b32 m0, 0x100001 -; CHECK-NEXT: s_barrier_init m0 -; CHECK-NEXT: s_barrier_signal 1 -; CHECK-NEXT: s_barrier_wait 1 -; CHECK-NEXT: s_endpgm +; CHECK-SDAG-LABEL: signal_var_misaligned: +; CHECK-SDAG: ; %bb.0: +; CHECK-SDAG-NEXT: global_wb +; CHECK-SDAG-NEXT: v_nop +; 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_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: global_wb +; CHECK-GISEL-NEXT: v_nop +; 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_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: ; %bb.0: ; CHECK-OBJ-SDAG-NEXT: global_wb ; CHECK-OBJ-SDAG-NEXT: v_nop ; 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_b32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo+1 +; CHECK-OBJ-SDAG-NEXT: s_add_co_i32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@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_bfe_u32 s0, s0, 0x60004 +; 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 @@ -145,20 +157,14 @@ define amdgpu_kernel void @signal_var_misaligned() { ; CHECK-OBJ-GISEL-NEXT: global_wb ; CHECK-OBJ-GISEL-NEXT: v_nop ; 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_add_co_u32 s0, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo, 1 -; CHECK-OBJ-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; CHECK-OBJ-GISEL-NEXT: s_lshr_b32 s0, s0, 4 -; 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_mov_b32 m0, 0x100000 ; CHECK-OBJ-GISEL-NEXT: s_barrier_init m0 -; CHECK-OBJ-GISEL-NEXT: s_mov_b32 m0, s0 -; CHECK-OBJ-GISEL-NEXT: s_barrier_signal m0 +; 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(3) @bars, i32 1 - call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) %p1, i32 16) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) %p1, i32 0) + %p1 = getelementptr i8, ptr addrspace(15) @bars, 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 } @@ -175,9 +181,9 @@ define amdgpu_kernel void @signal_var_dynamic(i32 %idx) { ; 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_load_b32 s0, s[4:5], 0x0 nv ; CHECK-SDAG-NEXT: s_wait_kmcnt 0x0 -; CHECK-SDAG-NEXT: s_lshl4_add_u32 s0, s0, 0x802010 +; CHECK-SDAG-NEXT: s_lshl4_add_u32 s0, s0, -1 ; CHECK-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; CHECK-SDAG-NEXT: s_bfe_u32 s0, s0, 0x60004 +; CHECK-SDAG-NEXT: s_and_b32 s0, s0, 63 ; CHECK-SDAG-NEXT: s_or_b32 m0, s0, 0x100000 ; CHECK-SDAG-NEXT: s_barrier_init m0 ; CHECK-SDAG-NEXT: s_mov_b32 m0, s0 @@ -194,10 +200,9 @@ define amdgpu_kernel void @signal_var_dynamic(i32 %idx) { ; 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, 0x802010, s0 -; CHECK-GISEL-NEXT: s_lshr_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_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 @@ -214,7 +219,7 @@ define amdgpu_kernel void @signal_var_dynamic(i32 %idx) { ; 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_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; CHECK-OBJ-SDAG-NEXT: s_bfe_u32 s0, s0, 0x60004 +; 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 @@ -231,21 +236,23 @@ define amdgpu_kernel void @signal_var_dynamic(i32 %idx) { ; 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, __amdgpu_named_barrier.bars.5a19a560517f8a3a4347b4502da34a70@abs32@lo, s0 -; CHECK-OBJ-GISEL-NEXT: s_lshr_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_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 ; CHECK-OBJ-GISEL-NEXT: s_barrier_signal m0 ; CHECK-OBJ-GISEL-NEXT: s_barrier_wait 1 ; CHECK-OBJ-GISEL-NEXT: s_endpgm - %p1 = getelementptr [2 x target("amdgcn.named.barrier", 0)], ptr addrspace(3) @bars, i32 0, i32 %idx - call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) %p1, i32 16) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) %p1, i32 0) + %p1 = getelementptr [2 x target("amdgcn.named.barrier", 0)], ptr addrspace(15) @bars, i32 0, i32 %idx + 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 } + +!0 = !{i32 -1, i32 0} + ;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: ; CHECK-OBJ: {{.*}} diff --git a/llvm/test/CodeGen/AMDGPU/s-barrier.ll b/llvm/test/CodeGen/AMDGPU/s-barrier.ll index ed07b9af50487..40d4e32c5d49a 100644 --- a/llvm/test/CodeGen/AMDGPU/s-barrier.ll +++ b/llvm/test/CodeGen/AMDGPU/s-barrier.ll @@ -2,9 +2,12 @@ ; RUN: llc -global-isel=0 -mtriple=amdgpu12.00 < %s | FileCheck -check-prefixes=GFX12,GFX12-SDAG %s ; RUN: llc -global-isel=1 -mtriple=amdgpu12.00 < %s | FileCheck -check-prefixes=GFX12,GFX12-GISEL %s -@bar = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison -@bar2 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison -@bar3 = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison +@bar = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison +@bar2 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison +@bar3 = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison + +; Test using the workgroup barrier with the GV. +@wgbarr = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison, !absolute_symbol !0 define void @func1() { ; GFX12-SDAG-LABEL: func1: @@ -33,8 +36,8 @@ define void @func1() { ; GFX12-GISEL-NEXT: s_barrier_signal m0 ; GFX12-GISEL-NEXT: s_barrier_wait 1 ; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31] - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar3) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar3, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar3) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar3, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } @@ -66,13 +69,13 @@ define void @func2() { ; GFX12-GISEL-NEXT: s_barrier_signal m0 ; GFX12-GISEL-NEXT: s_barrier_wait 1 ; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31] - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar2) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar2, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar2) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar2, i32 7) call void @llvm.amdgcn.s.barrier.wait(i16 1) ret void } -define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) #0 { +define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(15) %in) #0 { ; GFX12-SDAG-LABEL: kernel1: ; GFX12-SDAG: ; %bb.0: ; GFX12-SDAG-NEXT: s_mov_b64 s[10:11], s[6:7] @@ -85,7 +88,7 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ; GFX12-SDAG-NEXT: s_mov_b64 s[4:5], s[0:1] ; GFX12-SDAG-NEXT: s_mov_b32 s32, 0 ; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0 -; GFX12-SDAG-NEXT: s_bfe_u32 s2, s2, 0x60004 +; GFX12-SDAG-NEXT: s_and_b32 s2, s2, 63 ; GFX12-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; GFX12-SDAG-NEXT: s_or_b32 s3, s2, 0x90000 ; GFX12-SDAG-NEXT: s_cmp_eq_u32 0, 0 @@ -140,9 +143,8 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ; GFX12-GISEL-NEXT: s_mov_b64 s[6:7], s[2:3] ; GFX12-GISEL-NEXT: s_mov_b32 s32, 0 ; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0 -; GFX12-GISEL-NEXT: s_lshr_b32 s0, s0, 4 -; GFX12-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) ; GFX12-GISEL-NEXT: s_and_b32 s0, s0, 63 +; GFX12-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; GFX12-GISEL-NEXT: s_or_b32 s1, s0, 0x90000 ; GFX12-GISEL-NEXT: s_cmp_eq_u32 0, 0 ; GFX12-GISEL-NEXT: s_mov_b32 m0, s1 @@ -187,17 +189,17 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ; GFX12-GISEL-NEXT: s_swappc_b64 s[30:31], s[0:1] ; GFX12-GISEL-NEXT: s_get_barrier_state s0, -1 ; GFX12-GISEL-NEXT: s_endpgm - call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) @bar, i32 12) - call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) %in, i32 9) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar, i32 12) - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) %in, i32 9) + call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) @bar, i32 12) + call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) %in, i32 9) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar, i32 12) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %in, i32 9) call void @llvm.amdgcn.s.barrier.signal(i32 -1) - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) %in) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) %in) %isfirst = call i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32 -1) call void @llvm.amdgcn.s.barrier.wait(i16 1) call void @llvm.amdgcn.s.barrier.leave(i16 1) - %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) @bar) - %state2 = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) %in) + %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) @bar) + %state2 = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) %in) call void @llvm.amdgcn.s.barrier() call void @func1() call void @func2() @@ -205,7 +207,7 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ret void } -define amdgpu_kernel void @kernel2(ptr addrspace(1) %out, ptr addrspace(3) %in) #0 { +define amdgpu_kernel void @kernel2(ptr addrspace(1) %out, ptr addrspace(15) %in) #0 { ; GFX12-SDAG-LABEL: kernel2: ; GFX12-SDAG: ; %bb.0: ; GFX12-SDAG-NEXT: s_mov_b64 s[10:11], s[6:7] @@ -249,8 +251,8 @@ define amdgpu_kernel void @kernel2(ptr addrspace(1) %out, ptr addrspace(3) %in) ; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0 ; GFX12-GISEL-NEXT: s_swappc_b64 s[30:31], s[12:13] ; GFX12-GISEL-NEXT: s_endpgm - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar, i32 7) - call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) @bar) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar, i32 7) + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar) call void @llvm.amdgcn.s.barrier.wait(i16 1) call void @func2() @@ -267,39 +269,26 @@ define void @signal_var_cnt0_const_bar() { ; GFX12-NEXT: s_wait_kmcnt 0x0 ; GFX12-NEXT: s_barrier_signal 1 ; GFX12-NEXT: s_setpc_b64 s[30:31] - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) @bar, i32 0) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar, i32 0) ret void } -define void @signal_var_cnt0_dynamic_bar(ptr addrspace(3) inreg %bar) { -; GFX12-SDAG-LABEL: signal_var_cnt0_dynamic_bar: -; GFX12-SDAG: ; %bb.0: -; GFX12-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 -; GFX12-SDAG-NEXT: s_wait_expcnt 0x0 -; GFX12-SDAG-NEXT: s_wait_samplecnt 0x0 -; GFX12-SDAG-NEXT: s_wait_bvhcnt 0x0 -; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0 -; GFX12-SDAG-NEXT: s_bfe_u32 m0, s0, 0x60004 -; GFX12-SDAG-NEXT: s_barrier_signal m0 -; GFX12-SDAG-NEXT: s_setpc_b64 s[30:31] -; -; GFX12-GISEL-LABEL: signal_var_cnt0_dynamic_bar: -; GFX12-GISEL: ; %bb.0: -; GFX12-GISEL-NEXT: s_wait_loadcnt_dscnt 0x0 -; GFX12-GISEL-NEXT: s_wait_expcnt 0x0 -; GFX12-GISEL-NEXT: s_wait_samplecnt 0x0 -; GFX12-GISEL-NEXT: s_wait_bvhcnt 0x0 -; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0 -; GFX12-GISEL-NEXT: s_lshr_b32 s0, s0, 4 -; GFX12-GISEL-NEXT: s_wait_alu depctr_sa_sdst(0) -; GFX12-GISEL-NEXT: s_and_b32 m0, s0, 63 -; GFX12-GISEL-NEXT: s_barrier_signal m0 -; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31] - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) %bar, i32 0) +define void @signal_var_cnt0_dynamic_bar(ptr addrspace(15) inreg %bar) { +; GFX12-LABEL: signal_var_cnt0_dynamic_bar: +; GFX12: ; %bb.0: +; GFX12-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX12-NEXT: s_wait_expcnt 0x0 +; GFX12-NEXT: s_wait_samplecnt 0x0 +; GFX12-NEXT: s_wait_bvhcnt 0x0 +; GFX12-NEXT: s_wait_kmcnt 0x0 +; GFX12-NEXT: s_and_b32 m0, s0, 63 +; GFX12-NEXT: s_barrier_signal m0 +; GFX12-NEXT: s_setpc_b64 s[30:31] + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %bar, i32 0) ret void } -define void @barrier_init_dynamic_cnt(ptr addrspace(3) inreg %bar, i32 inreg %cnt) { +define void @barrier_init_dynamic_cnt(ptr addrspace(15) inreg %bar, i32 inreg %cnt) { ; GFX12-SDAG-LABEL: barrier_init_dynamic_cnt: ; GFX12-SDAG: ; %bb.0: ; GFX12-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 @@ -308,7 +297,7 @@ define void @barrier_init_dynamic_cnt(ptr addrspace(3) inreg %bar, i32 inreg %cn ; GFX12-SDAG-NEXT: s_wait_bvhcnt 0x0 ; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0 ; GFX12-SDAG-NEXT: s_and_b32 s1, s1, 63 -; GFX12-SDAG-NEXT: s_bfe_u32 s0, s0, 0x60004 +; GFX12-SDAG-NEXT: s_and_b32 s0, s0, 63 ; GFX12-SDAG-NEXT: s_wait_alu depctr_sa_sdst(0) ; GFX12-SDAG-NEXT: s_lshl_b32 s1, s1, 16 ; GFX12-SDAG-NEXT: s_wait_alu depctr_sa_sdst(0) @@ -323,20 +312,19 @@ define void @barrier_init_dynamic_cnt(ptr addrspace(3) inreg %bar, i32 inreg %cn ; GFX12-GISEL-NEXT: s_wait_samplecnt 0x0 ; GFX12-GISEL-NEXT: s_wait_bvhcnt 0x0 ; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0 -; GFX12-GISEL-NEXT: s_lshr_b32 s0, s0, 4 ; GFX12-GISEL-NEXT: s_and_b32 s1, s1, 63 -; GFX12-GISEL-NEXT: s_wait_alu depctr_sa_sdst(0) ; GFX12-GISEL-NEXT: s_and_b32 s0, s0, 63 +; GFX12-GISEL-NEXT: s_wait_alu depctr_sa_sdst(0) ; GFX12-GISEL-NEXT: s_lshl_b32 s1, s1, 16 ; GFX12-GISEL-NEXT: s_wait_alu depctr_sa_sdst(0) ; GFX12-GISEL-NEXT: s_or_b32 m0, s0, s1 ; GFX12-GISEL-NEXT: s_barrier_init m0 ; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31] - call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) %bar, i32 %cnt) + call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) %bar, i32 %cnt) ret void } -define void @signal_var_dynamic_cnt(ptr addrspace(3) inreg %bar, i32 inreg %cnt) { +define void @signal_var_dynamic_cnt(ptr addrspace(15) inreg %bar, i32 inreg %cnt) { ; GFX12-SDAG-LABEL: signal_var_dynamic_cnt: ; GFX12-SDAG: ; %bb.0: ; GFX12-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 @@ -345,7 +333,7 @@ define void @signal_var_dynamic_cnt(ptr addrspace(3) inreg %bar, i32 inreg %cnt) ; GFX12-SDAG-NEXT: s_wait_bvhcnt 0x0 ; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0 ; GFX12-SDAG-NEXT: s_and_b32 s1, s1, 63 -; GFX12-SDAG-NEXT: s_bfe_u32 s0, s0, 0x60004 +; GFX12-SDAG-NEXT: s_and_b32 s0, s0, 63 ; GFX12-SDAG-NEXT: s_wait_alu depctr_sa_sdst(0) ; GFX12-SDAG-NEXT: s_lshl_b32 s1, s1, 16 ; GFX12-SDAG-NEXT: s_wait_alu depctr_sa_sdst(0) @@ -360,16 +348,15 @@ define void @signal_var_dynamic_cnt(ptr addrspace(3) inreg %bar, i32 inreg %cnt) ; GFX12-GISEL-NEXT: s_wait_samplecnt 0x0 ; GFX12-GISEL-NEXT: s_wait_bvhcnt 0x0 ; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0 -; GFX12-GISEL-NEXT: s_lshr_b32 s0, s0, 4 ; GFX12-GISEL-NEXT: s_and_b32 s1, s1, 63 -; GFX12-GISEL-NEXT: s_wait_alu depctr_sa_sdst(0) ; GFX12-GISEL-NEXT: s_and_b32 s0, s0, 63 +; GFX12-GISEL-NEXT: s_wait_alu depctr_sa_sdst(0) ; GFX12-GISEL-NEXT: s_lshl_b32 s1, s1, 16 ; GFX12-GISEL-NEXT: s_wait_alu depctr_sa_sdst(0) ; GFX12-GISEL-NEXT: s_or_b32 m0, s0, s1 ; GFX12-GISEL-NEXT: s_barrier_signal m0 ; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31] - call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) %bar, i32 %cnt) + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %bar, i32 %cnt) ret void } @@ -392,17 +379,41 @@ define amdgpu_ps void @test_barrier_leave_write_to_scc(i32 inreg %val, ptr addrs ret void } + +define amdgpu_kernel void @wgbarr_as_gv() { +; GFX12-LABEL: wgbarr_as_gv: +; GFX12: ; %bb.0: +; GFX12-NEXT: s_mov_b32 m0, 0x7003f +; GFX12-NEXT: s_barrier_signal m0 +; GFX12-NEXT: s_barrier_wait -1 +; GFX12-NEXT: s_endpgm + call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @wgbarr, i32 7) + call void @llvm.amdgcn.s.barrier.wait(i16 -1) + ret void +} + +define amdgpu_kernel void @null_barrier() { +; GFX12-LABEL: null_barrier: +; GFX12: ; %bb.0: +; GFX12-NEXT: s_barrier_join 0 +; GFX12-NEXT: s_endpgm + call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) null) + ret void +} + declare void @llvm.amdgcn.s.barrier() #1 declare void @llvm.amdgcn.s.barrier.wait(i16) #1 declare void @llvm.amdgcn.s.barrier.signal(i32) #1 -declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3), i32) #1 +declare void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15), i32) #1 declare i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32) #1 -declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(3), i32) #1 -declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(3)) #1 +declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(15), i32) #1 +declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(15)) #1 declare void @llvm.amdgcn.s.barrier.leave(i16) #1 declare i32 @llvm.amdgcn.s.get.barrier.state(i32) #1 -declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3)) #1 +declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15)) #1 attributes #0 = { nounwind } attributes #1 = { convergent nounwind } attributes #2 = { nounwind readnone } + +!0 = !{i32 -1, i32 0} diff --git a/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll b/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll index 182e2ac3dd985..0868a2b6a0cbe 100644 --- a/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll +++ b/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll @@ -5,11 +5,11 @@ ; RUN: not llc -global-isel=0 -mtriple=amdgpu12.00 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -mtriple=amdgpu12.00 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s -; ERR: error: :0:0: in function @kernel1 void (ptr addrspace(1), ptr addrspace(3)): llvm.amdgcn.s.wakeup.barrier requires target feature 's-wakeup-barrier-inst' +; ERR: error: :0:0: in function @kernel1 void (ptr addrspace(1), ptr addrspace(15)): llvm.amdgcn.s.wakeup.barrier requires target feature 's-wakeup-barrier-inst' -@bar = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison +@bar = internal addrspace(15) global target("amdgcn.named.barrier", 0) poison -define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) #0 { +define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(15) %in) #0 { ; GFX1250-SDAG-LABEL: kernel1: ; GFX1250-SDAG: ; %bb.0: ; GFX1250-SDAG-NEXT: global_wb @@ -19,7 +19,7 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ; GFX1250-SDAG-NEXT: s_mov_b32 m0, 1 ; GFX1250-SDAG-NEXT: s_wakeup_barrier m0 ; GFX1250-SDAG-NEXT: s_wait_kmcnt 0x0 -; GFX1250-SDAG-NEXT: s_bfe_u32 m0, s0, 0x60004 +; GFX1250-SDAG-NEXT: s_and_b32 m0, s0, 63 ; GFX1250-SDAG-NEXT: s_wakeup_barrier m0 ; GFX1250-SDAG-NEXT: s_endpgm ; @@ -31,18 +31,16 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ; GFX1250-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x2c nv ; GFX1250-GISEL-NEXT: s_wakeup_barrier 1 ; GFX1250-GISEL-NEXT: s_wait_kmcnt 0x0 -; GFX1250-GISEL-NEXT: s_lshr_b32 s0, s0, 4 -; GFX1250-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) ; GFX1250-GISEL-NEXT: s_and_b32 m0, s0, 63 ; GFX1250-GISEL-NEXT: s_wakeup_barrier m0 ; GFX1250-GISEL-NEXT: s_endpgm - call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) @bar) - call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %in) + call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(15) @bar) + call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(15) %in) ret void } -declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3)) #1 +declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(15)) #1 attributes #0 = { nounwind } attributes #1 = { convergent nounwind } diff --git a/llvm/test/CodeGen/AMDGPU/simple-indirect-call.ll b/llvm/test/CodeGen/AMDGPU/simple-indirect-call.ll index c6cd2b15203f7..66f1b950659e9 100644 --- a/llvm/test/CodeGen/AMDGPU/simple-indirect-call.ll +++ b/llvm/test/CodeGen/AMDGPU/simple-indirect-call.ll @@ -59,5 +59,5 @@ define amdgpu_kernel void @test_simple_indirect_call() { ;. ; ATTRIBUTOR_GCN: attributes #[[ATTR0]] = { "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-no-wwm" } ;. -; ATTRIBUTOR_GCN: [[META0]] = !{i32 1, i32 5, i32 6, i32 10} +; ATTRIBUTOR_GCN: [[META0]] = !{i32 1, i32 5, i32 6, i32 16} ;. diff --git a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp index a082adbf6565e..edff0ec1178cd 100644 --- a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp +++ b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp @@ -43,14 +43,17 @@ TEST(DataLayoutUpgradeTest, ValidDataLayoutUpgrade) { // and that ANDGCN adds p7 and p8 as well. EXPECT_EQ(UpgradeDataLayoutString("e-p:64:64", "amdgcn"), "m:e-e-p:64:64-G1-ni:7:8:9-p7:160:256:256:32-p8:128:128:128:48-p9:" - "192:256:256:32"); + "192:256:256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-" + "p15:32:32"); EXPECT_EQ(UpgradeDataLayoutString("e-p:64:64-G1", "amdgcn"), "m:e-e-p:64:64-G1-ni:7:8:9-p7:160:256:256:32-p8:128:128:128:48-p9:" - "192:256:256:32"); + "192:256:256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-" + "p15:32:32"); // Check that the old AMDGCN p8:128:128 definition is upgraded EXPECT_EQ(UpgradeDataLayoutString("e-p:64:64-p8:128:128-G1", "amdgcn"), "m:e-e-p:64:64-p8:128:128:128:48-G1-ni:7:8:9-p7:160:256:256:32-p9:" - "192:256:256:32"); + "192:256:256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-" + "p15:32:32"); // but that r600 does not. EXPECT_EQ(UpgradeDataLayoutString("e-p:32:32-G1", "r600"), "m:e-e-p:32:32-G1"); @@ -66,7 +69,9 @@ TEST(DataLayoutUpgradeTest, ValidDataLayoutUpgrade) { "m:e-e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:" "64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:" "1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9-p7:160:256:256:32-p8:128:128:" - "128:48-p9:192:256:256:32"); + "128:48-p9:192:256:256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:" + "32" + "-p15:32:32"); // Check that SystemZ adds -S64 if needed. EXPECT_EQ(UpgradeDataLayoutString( @@ -158,24 +163,38 @@ TEST(DataLayoutUpgradeTest, NoDataLayoutUpgrade) { EXPECT_EQ(UpgradeDataLayoutString("G2", "r600"), "m:e-G2"); EXPECT_EQ(UpgradeDataLayoutString("e-p:64:64-G2", "amdgcn"), "m:e-e-p:64:64-G2-ni:7:8:9-p7:160:256:256:32-p8:128:128:128:48-p9:" - "192:256:256:32"); + "192:256:256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32" + "-p15:32:32"); EXPECT_EQ(UpgradeDataLayoutString("G2-e-p:64:64", "amdgcn"), "m:e-G2-e-p:64:64-ni:7:8:9-p7:160:256:256:32-p8:128:128:128:48-p9:" - "192:256:256:32"); + "192:256:256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32" + "-p15:32:32"); EXPECT_EQ(UpgradeDataLayoutString("e-p:64:64-G0", "amdgcn"), "m:e-e-p:64:64-G0-ni:7:8:9-p7:160:256:256:32-p8:128:128:128:48-p9:" - "192:256:256:32"); + "192:256:256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32" + "-p15:32:32"); // Check that AMDGCN targets don't add already declared address space 7. EXPECT_EQ( UpgradeDataLayoutString("e-p:64:64-p7:64:64", "amdgcn"), - "m:e-e-p:64:64-p7:64:64-G1-ni:7:8:9-p8:128:128:128:48-p9:192:256:256:32"); + "m:e-e-p:64:64-p7:64:64-G1-ni:7:8:9-p8:128:128:128:48-p9:192:256:" + "256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-p15:32:32"); EXPECT_EQ( UpgradeDataLayoutString("p7:64:64-G2-e-p:64:64", "amdgcn"), - "m:e-p7:64:64-G2-e-p:64:64-ni:7:8:9-p8:128:128:128:48-p9:192:256:256:32"); + "m:e-p7:64:64-G2-e-p:64:64-ni:7:8:9-p8:128:128:128:48-p9:192:256:" + "256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-p15:32:32"); EXPECT_EQ( UpgradeDataLayoutString("e-p:64:64-p7:64:64-G1", "amdgcn"), - "m:e-e-p:64:64-p7:64:64-G1-ni:7:8:9-p8:128:128:128:48-p9:192:256:256:32"); + "m:e-e-p:64:64-p7:64:64-G1-ni:7:8:9-p8:128:128:128:48-p9:192:256:" + "256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-p15:32:32"); + + // Check that AMDGCN targets don't add already declared address space 15. + EXPECT_EQ(UpgradeDataLayoutString("e-p:64:64-p15:32:32", "amdgcn"), + "m:e-e-p:64:64-p15:32:32-G1-ni:7:8:9-p7:160:256:256:32-p8:128:128:" + "128:48-p9:192:256:256:32"); + EXPECT_EQ(UpgradeDataLayoutString("p15:32:32-G2-e-p:64:64", "amdgcn"), + "m:e-p15:32:32-G2-e-p:64:64-ni:7:8:9-p7:160:256:256:32-p8:128:128:" + "128:48-p9:192:256:256:32"); // Check that SPIR & SPIRV targets don't add -G1 if there is already a -G // flag. @@ -218,7 +237,8 @@ TEST(DataLayoutUpgradeTest, EmptyDataLayout) { EXPECT_EQ(UpgradeDataLayoutString("", "r600"), "m:e-G1"); EXPECT_EQ( UpgradeDataLayoutString("", "amdgcn"), - "m:e-G1-ni:7:8:9-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32"); + "m:e-G1-ni:7:8:9-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:" + "256:32-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-p15:32:32"); // Check that SPIR & SPIRV targets add G1 if it's not present. EXPECT_EQ(UpgradeDataLayoutString("", "spir"), "G1"); diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.td index 0807cf8cf04a4..2b4002e98caae 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.td +++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.td @@ -142,6 +142,8 @@ def ROCDL_Dialect : Dialect { static constexpr unsigned kConstantMemoryAddressSpace = 4; /// The address space value that represents private memory. static constexpr unsigned kPrivateMemoryAddressSpace = 5; + /// The address space value that represents barriers. + static constexpr unsigned kBarrierAddressSpace = 15; }]; let discardableAttrs = (ins diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td index ba185dc27dc7e..6e72bda93f5bd 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td @@ -405,16 +405,17 @@ def ROCDL_WaveBarrierOp : ROCDL_ConcreteNonMemIntrOp<"wave.barrier", [], 0> { def ROCDLGlobalBuffer : LLVM_PointerInAddressSpace<1>; def ROCDLBufferLDS : LLVM_PointerInAddressSpace<3>; +def ROCDLBarrier : LLVM_PointerInAddressSpace<15>; def ROCDL_BarrierInitOp : ROCDL_IntrOp<"s.barrier.init", [], [], [], 0, 0, 0, 0, [1], ["memberCnt"]>, - Arguments<(ins Arg:$ptr, I32Attr:$memberCnt)> { + Arguments<(ins Arg:$ptr, I32Attr:$memberCnt)> { let description = [{ Available on gfx1250+. Example: ```mlir // Initialize a named barrier with member count. - rocdl.s.barrier.init %ptr member_cnt = 1 : !llvm.ptr<3> + rocdl.s.barrier.init %ptr member_cnt = 1 : !llvm.ptr<15> ``` }]; let results = (outs); @@ -437,7 +438,7 @@ def ROCDL_BarrierSignalOp : ROCDL_ConcreteNonMemIntrOp<"s.barrier.signal", [], 0 } def ROCDL_BarrierSignalVarOp : ROCDL_IntrOp<"s.barrier.signal.var", [], [], [], 0, 0, 0, 0, [1], ["memberCnt"]>, - Arguments<(ins Arg:$ptr, I32Attr:$memberCnt)> { + Arguments<(ins Arg:$ptr, I32Attr:$memberCnt)> { let description = [{ Available on gfx1250+. @@ -446,7 +447,7 @@ def ROCDL_BarrierSignalVarOp : ROCDL_IntrOp<"s.barrier.signal.var", [], [], [], Example: ```mlir // Signal a named barrier with variable ID. - rocdl.s.barrier.signal.var %ptr member_cnt = 1 : !llvm.ptr<3> + rocdl.s.barrier.signal.var %ptr member_cnt = 1 : !llvm.ptr<15> ``` }]; let results = (outs); @@ -454,14 +455,14 @@ def ROCDL_BarrierSignalVarOp : ROCDL_IntrOp<"s.barrier.signal.var", [], [], [], } def ROCDL_BarrierJoinOp : ROCDL_IntrOp<"s.barrier.join", [], [], [], 0>, - Arguments<(ins Arg:$ptr)> { + Arguments<(ins Arg:$ptr)> { let description = [{ Available on gfx1250+. Example: ```mlir // Join a named barrier. - rocdl.s.barrier.join %ptr : !llvm.ptr<3> + rocdl.s.barrier.join %ptr : !llvm.ptr<15> ``` }]; let results = (outs); @@ -529,14 +530,14 @@ def ROCDL_GetBarrierStateOp : ROCDL_ConcreteNonMemIntrOp<"s.get.barrier.state", } def ROCDL_GetNamedBarrierStateOp : ROCDL_ConcreteNonMemIntrOp<"s.get.named.barrier.state", [], 1, [], []>, - Arguments<(ins Arg:$ptr)> { + Arguments<(ins Arg:$ptr)> { let description = [{ Available on gfx1250+. Example: ```mlir // Query named barrier state by pointer. - %0 = rocdl.s.get.named.barrier.state %ptr : !llvm.ptr<3> -> i32 + %0 = rocdl.s.get.named.barrier.state %ptr : !llvm.ptr<15> -> i32 ``` }]; let results = (outs I32:$res); @@ -544,7 +545,7 @@ def ROCDL_GetNamedBarrierStateOp : ROCDL_ConcreteNonMemIntrOp<"s.get.named.barri } def ROCDL_WakeupBarrierOp : ROCDL_ConcreteNonMemIntrOp<"s.wakeup.barrier", [], 0, [], []>, - Arguments<(ins Arg:$ptr)> { + Arguments<(ins Arg:$ptr)> { let description = [{ Wakes up waves associated with a given named barrier. Note, This op does not release waves waiting at the barrier. It just signal other waves in the same work-group waiting on the indicated named barrier @@ -554,7 +555,7 @@ def ROCDL_WakeupBarrierOp : ROCDL_ConcreteNonMemIntrOp<"s.wakeup.barrier", [], 0 Example: ```mlir // Wake up waves waiting on a named barrier. - rocdl.s.wakeup.barrier %ptr : !llvm.ptr<3> + rocdl.s.wakeup.barrier %ptr : !llvm.ptr<15> ``` }]; let assemblyFormat = "$ptr attr-dict `:` qualified(type($ptr))"; diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp index 90e15b1680446..3b3cc6bd81fec 100644 --- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp +++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp @@ -4553,7 +4553,7 @@ void mlir::amdgpu::populateCommonGPUTypeAndAttributeConversions( }); typeConverter.addConversion([](gpu::NamedBarrierType type) { return LLVM::LLVMPointerType::get( - type.getContext(), ROCDL::ROCDLDialect::kSharedMemoryAddressSpace); + type.getContext(), ROCDL::ROCDLDialect::kBarrierAddressSpace); }); } diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp index a3819df4f8a84..ef1292918a1cf 100644 --- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp +++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp @@ -683,7 +683,8 @@ struct GPUInitializeNamedBarrierOpLowering final auto targetTy = LLVM::LLVMTargetExtType::get( rewriter.getContext(), "amdgcn.named.barrier", {}, {0}); - auto ptrTy = LLVM::LLVMPointerType::get(rewriter.getContext(), 3); + auto ptrTy = LLVM::LLVMPointerType::get( + rewriter.getContext(), ROCDL::ROCDLDialect::kBarrierAddressSpace); // Build the global detached so SymbolTable::insert can both place it and // rename it as needed without creating a transient name conflict in IR. @@ -691,7 +692,8 @@ struct GPUInitializeNamedBarrierOpLowering final auto globalOp = LLVM::GlobalOp::create( detachedBuilder, loc, targetTy, /*isConstant=*/false, LLVM::Linkage::Internal, "__named_barrier", /*value=*/Attribute(), - /*alignment=*/0, /*addrSpace=*/3); + /*alignment=*/0, + /*addrSpace=*/ROCDL::ROCDLDialect::kBarrierAddressSpace); // Initialize with poison. { Region ®ion = globalOp.getInitializerRegion(); diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir index c6a9574ca43c1..402dcae5e9832 100644 --- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir +++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir @@ -5,7 +5,7 @@ gpu.module @test_module { // CHECK-LABEL: func @named_barrier func.func @named_barrier() { %member_count = arith.constant 4 : i32 - // CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @[[NB:__named_barrier[_0-9]*]] : !llvm.ptr<3> + // CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @[[NB:__named_barrier[_0-9]*]] : !llvm.ptr<15> // CHECK: rocdl.s.barrier.init %[[ADDR]] member_cnt = 4 %nb = gpu.initialize_named_barrier %member_count : i32 -> !gpu.named_barrier // CHECK: llvm.fence syncscope("workgroup") release @@ -21,10 +21,10 @@ func.func @named_barrier() { func.func @two_named_barriers() { %c4 = arith.constant 4 : i32 %c8 = arith.constant 8 : i32 - // CHECK: %[[ADDR0:.*]] = llvm.mlir.addressof @[[NB0:__named_barrier[_0-9]*]] : !llvm.ptr<3> + // CHECK: %[[ADDR0:.*]] = llvm.mlir.addressof @[[NB0:__named_barrier[_0-9]*]] : !llvm.ptr<15> // CHECK: rocdl.s.barrier.init %[[ADDR0]] member_cnt = 4 %nb0 = gpu.initialize_named_barrier %c4 : i32 -> !gpu.named_barrier - // CHECK: %[[ADDR1:.*]] = llvm.mlir.addressof @[[NB1:__named_barrier[_0-9]*]] : !llvm.ptr<3> + // CHECK: %[[ADDR1:.*]] = llvm.mlir.addressof @[[NB1:__named_barrier[_0-9]*]] : !llvm.ptr<15> // CHECK: rocdl.s.barrier.init %[[ADDR1]] member_cnt = 8 %nb1 = gpu.initialize_named_barrier %c8 : i32 -> !gpu.named_barrier // CHECK: rocdl.s.barrier.join %[[ADDR0]] @@ -49,6 +49,6 @@ func.func @cluster_scope() { } // One LDS global per gpu.initialize_named_barrier. -// CHECK-COUNT-3: llvm.mlir.global internal @__named_barrier{{[_0-9]*}}() {addr_space = 3 : i32} : !llvm.target<"amdgcn.named.barrier", 0> +// CHECK-COUNT-3: llvm.mlir.global internal @__named_barrier{{[_0-9]*}}() {addr_space = 15 : i32} : !llvm.target<"amdgcn.named.barrier", 0> } diff --git a/mlir/test/Dialect/LLVMIR/rocdl.mlir b/mlir/test/Dialect/LLVMIR/rocdl.mlir index dd0b00faf7f1f..505bff44a6f61 100644 --- a/mlir/test/Dialect/LLVMIR/rocdl.mlir +++ b/mlir/test/Dialect/LLVMIR/rocdl.mlir @@ -1216,10 +1216,10 @@ llvm.func @rocdl.s.barrier() { llvm.return } -llvm.func @rocdl.s.barrier.init(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.barrier.init(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.barrier.init - // CHECK: rocdl.s.barrier.init %{{.*}} member_cnt = 1 : !llvm.ptr<3> - rocdl.s.barrier.init %ptr member_cnt = 1 : !llvm.ptr<3> + // CHECK: rocdl.s.barrier.init %{{.*}} member_cnt = 1 : !llvm.ptr<15> + rocdl.s.barrier.init %ptr member_cnt = 1 : !llvm.ptr<15> llvm.return } @@ -1230,17 +1230,17 @@ llvm.func @rocdl.s.barrier.signal() { llvm.return } -llvm.func @rocdl.s.barrier.signal.var(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.barrier.signal.var(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.barrier.signal.var - // CHECK: rocdl.s.barrier.signal.var %{{.*}} member_cnt = 1 : !llvm.ptr<3> - rocdl.s.barrier.signal.var %ptr member_cnt = 1 : !llvm.ptr<3> + // CHECK: rocdl.s.barrier.signal.var %{{.*}} member_cnt = 1 : !llvm.ptr<15> + rocdl.s.barrier.signal.var %ptr member_cnt = 1 : !llvm.ptr<15> llvm.return } -llvm.func @rocdl.s.barrier.join(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.barrier.join(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.barrier.join - // CHECK: rocdl.s.barrier.join %{{.*}} : !llvm.ptr<3> - rocdl.s.barrier.join %ptr : !llvm.ptr<3> + // CHECK: rocdl.s.barrier.join %{{.*}} : !llvm.ptr<15> + rocdl.s.barrier.join %ptr : !llvm.ptr<15> llvm.return } @@ -1272,17 +1272,17 @@ llvm.func @rocdl.s.get.barrier.state() { llvm.return } -llvm.func @rocdl.s.get.named.barrier.state(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.get.named.barrier.state(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.get.named.barrier.state - // CHECK: rocdl.s.get.named.barrier.state %{{.*}} : !llvm.ptr<3> -> i32 - %0 = rocdl.s.get.named.barrier.state %ptr : !llvm.ptr<3> -> i32 + // CHECK: rocdl.s.get.named.barrier.state %{{.*}} : !llvm.ptr<15> -> i32 + %0 = rocdl.s.get.named.barrier.state %ptr : !llvm.ptr<15> -> i32 llvm.return } -llvm.func @rocdl.s.wakeup.barrier(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.wakeup.barrier(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.wakeup.barrier - // CHECK: rocdl.s.wakeup.barrier %{{.*}} : !llvm.ptr<3> - rocdl.s.wakeup.barrier %ptr : !llvm.ptr<3> + // CHECK: rocdl.s.wakeup.barrier %{{.*}} : !llvm.ptr<15> + rocdl.s.wakeup.barrier %ptr : !llvm.ptr<15> llvm.return } diff --git a/mlir/test/Target/LLVMIR/rocdl.mlir b/mlir/test/Target/LLVMIR/rocdl.mlir index 9b767d8574473..989e9c5919587 100644 --- a/mlir/test/Target/LLVMIR/rocdl.mlir +++ b/mlir/test/Target/LLVMIR/rocdl.mlir @@ -278,10 +278,10 @@ llvm.func @rocdl.wave_barrier() { llvm.return } -llvm.func @rocdl.s.barrier.init(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.barrier.init(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.barrier.init - // CHECK: call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) %{{.*}}, i32 1) - rocdl.s.barrier.init %ptr member_cnt = 1 : !llvm.ptr<3> + // CHECK: call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) %{{.*}}, i32 1) + rocdl.s.barrier.init %ptr member_cnt = 1 : !llvm.ptr<15> llvm.return } @@ -292,17 +292,17 @@ llvm.func @rocdl.s.barrier.signal() { llvm.return } -llvm.func @rocdl.s.barrier.signal.var(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.barrier.signal.var(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.barrier.signal.var - // CHECK: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) %{{.*}}, i32 1) - rocdl.s.barrier.signal.var %ptr member_cnt = 1 : !llvm.ptr<3> + // CHECK: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) %{{.*}}, i32 1) + rocdl.s.barrier.signal.var %ptr member_cnt = 1 : !llvm.ptr<15> llvm.return } -llvm.func @rocdl.s.barrier.join(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.barrier.join(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.barrier.join - // CHECK: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) %{{.*}}) - rocdl.s.barrier.join %ptr : !llvm.ptr<3> + // CHECK: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) %{{.*}}) + rocdl.s.barrier.join %ptr : !llvm.ptr<15> llvm.return } @@ -334,17 +334,17 @@ llvm.func @rocdl.s.get.barrier.state() { llvm.return } -llvm.func @rocdl.s.get.named.barrier.state(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.get.named.barrier.state(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.get.named.barrier.state - // CHECK: %{{.*}} = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) %{{.*}}) - %0 = rocdl.s.get.named.barrier.state %ptr : !llvm.ptr<3> -> i32 + // CHECK: %{{.*}} = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(15) %{{.*}}) + %0 = rocdl.s.get.named.barrier.state %ptr : !llvm.ptr<15> -> i32 llvm.return } -llvm.func @rocdl.s.wakeup.barrier(%ptr : !llvm.ptr<3>) { +llvm.func @rocdl.s.wakeup.barrier(%ptr : !llvm.ptr<15>) { // CHECK-LABEL: rocdl.s.wakeup.barrier - // CHECK: call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %{{.*}}) - rocdl.s.wakeup.barrier %ptr : !llvm.ptr<3> + // CHECK: call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(15) %{{.*}}) + rocdl.s.wakeup.barrier %ptr : !llvm.ptr<15> llvm.return }