llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-mlir @llvm/pr-subscribers-mlir-gpu Author: Pierre van Houtryve (Pierre-vh) <details> <summary>Changes</summary> (Recreated from #<!-- -->195613 due to a rebase issue. Apologies for the inconvenience) Add a new BARRIER address space that is used for global variables that are used to represent the barrier IDs in GFX12.5. These barrier addresses just have values corresponding 1-1 to barrier IDs. They are still implemented on top of LDS, but the offsetting happens during an addrspacecast to generic, not whenever the barrier GV is used. The motivation for this is to make the relation between LDS and barrier GVs explicit in the compiler. It does add a bit more complexity, but that complexity was already there, just hidden by pretending barrier GVs were actual LDS. --- Patch is 178.10 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/209746.diff 54 Files Affected: - (modified) clang/lib/Basic/Targets/AMDGPU.cpp (+1-1) - (modified) clang/test/CodeGen/target-data.c (+2-2) - (modified) clang/test/CodeGenHIP/amdgpu-barrier-type.hip (+8-8) - (modified) clang/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl (+1-1) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl (+8-8) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+2-2) - (modified) lld/test/ELF/lto/amdgcn-oses.ll (+3-3) - (modified) lld/test/ELF/lto/amdgcn.ll (+1-1) - (modified) llvm/docs/AMDGPUUsage.rst (+25-9) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+6-5) - (modified) llvm/include/llvm/Support/AMDGPUAddrSpace.h (+14-2) - (modified) llvm/lib/Target/AMDGPU/AMDGPU.h (+17-11) - (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+1) - (modified) llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp (+31-21) - (modified) llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp (+6-18) - (modified) llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp (+52-14) - (modified) llvm/lib/Target/AMDGPU/AMDGPULowerExecSync.cpp (+15-14) - (modified) llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp (-10) - (modified) llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp (+3-1) - (modified) llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.cpp (+28-14) - (modified) llvm/lib/Target/AMDGPU/AMDGPUMachineFunctionInfo.h (+7-2) - (modified) llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp (+3-9) - (modified) llvm/lib/Target/AMDGPU/SIDefines.h (-4) - (modified) llvm/lib/Target/AMDGPU/SIISelLowering.cpp (+59-43) - (modified) llvm/lib/Target/AMDGPU/SIRegisterInfo.td (+1-1) - (modified) llvm/lib/TargetParser/TargetDataLayout.cpp (+3-2) - (modified) llvm/test/Analysis/UniformityAnalysis/AMDGPU/always_uniform.ll (+3-3) - (added) llvm/test/CodeGen/AMDGPU/addrspacecast-barrier.ll (+474) - (modified) llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-module-lds.ll (+32-32) - (modified) llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync-and-sw-lds.ll (+14-25) - (modified) llvm/test/CodeGen/AMDGPU/amdgpu-lower-exec-sync.ll (+32-32) - (modified) llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll (+4-4) - (modified) llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit-undefined-behavior.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/attributor-flatscratchinit.ll (+10-10) - (modified) llvm/test/CodeGen/AMDGPU/attributor-noalias-addrspace.ll (+2-2) - (added) llvm/test/CodeGen/AMDGPU/barrier-addrspace-dereference.ll (+16) - (modified) llvm/test/CodeGen/AMDGPU/lds-link-time-codegen-named-barrier.ll (+5-8) - (modified) llvm/test/CodeGen/AMDGPU/lds-link-time-named-barrier.ll (+7-7) - (added) llvm/test/CodeGen/AMDGPU/null-named-barrier-gv.ll (+31) - (modified) llvm/test/CodeGen/AMDGPU/s-barrier-id-allocation.ll (+21-21) - (added) llvm/test/CodeGen/AMDGPU/s-barrier-lowering-bad-absolute-symbol.ll (+16) - (added) llvm/test/CodeGen/AMDGPU/s-barrier-lowering-wrong-gv-signature.ll (+27) - (modified) llvm/test/CodeGen/AMDGPU/s-barrier-lowering.ll (+29-31) - (modified) llvm/test/CodeGen/AMDGPU/s-barrier-signal-var-gep.ll (+81-74) - (modified) llvm/test/CodeGen/AMDGPU/s-barrier.ll (+72-61) - (modified) llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll (+6-8) - (modified) llvm/test/CodeGen/AMDGPU/simple-indirect-call.ll (+1-1) - (modified) mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.td (+2) - (modified) mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td (+11-10) - (modified) mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp (+1-1) - (modified) mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp (+4-2) - (modified) mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir (+4-4) - (modified) mlir/test/Dialect/LLVMIR/rocdl.mlir (+15-15) - (modified) mlir/test/Target/LLVMIR/rocdl.mlir (+15-15) ``````````diff diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 4665dac9cb787..a2330dd1d4c8b 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -57,7 +57,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUAddrSpaceMap = { llvm::AMDGPUAS::PRIVATE_ADDRESS, // hlsl_output llvm::AMDGPUAS::GLOBAL_ADDRESS, // hlsl_push_constant llvm::AMDGPUAS::FLAT_ADDRESS, // wasm_funcref - llvm::AMDGPUAS::LOCAL_ADDRESS, // amdgpu_barrier + llvm::AMDGPUAS::BARRIER, // amdgpu_barrier }; } // namespace targets diff --git a/clang/test/CodeGen/target-data.c b/clang/test/CodeGen/target-data.c index a5e0b814c7042..1a0e8a0a7fc3f 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 b3caf0f0de3b9..7c2101266a85c 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 = "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" !llvm.module.flags = !{!0} !0 = !{i32 1, !"amdhsa_code_object_version", i32 500} @@ -36,7 +36,7 @@ define void @_start() { ;--- amdpal.ll target triple = "amdgcn-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 = "amdgcn-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/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst index cbe9d15074457..42345ad1ebe2f 100644 --- a/llvm/docs/AMDGPUUsage.rst +++ b/llvm/docs/AMDGPUUsage.rst @@ -1120,6 +1120,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 ===================================== =============== =========== ================ ======= ============================ @@ -1333,6 +1334,23 @@ 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 + + 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 @@ -1539,10 +1557,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: @@ -1553,14 +1569,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 21882e247c027..da8cf32b8ead4 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,19 +308,19 @@ 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. 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) @@ -342,7 +343,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_st... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/209746 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
