diff --git a/clang/test/CodeGen/target-data.c b/clang/test/CodeGen/target-data.c index f2a09a40ee685..3bc58aa46ce69 100644 --- a/clang/test/CodeGen/target-data.c +++ b/clang/test/CodeGen/target-data.c @@ -160,12 +160,12 @@ // RUN: %clang_cc1 -triple amdgpu7.01-unknown -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-i64:64-i128:128-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 amdgpu-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-i64:64-i128:128-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/printf_nonhostcall.cpp b/clang/test/CodeGenHIP/printf_nonhostcall.cpp index 7e8a468a33527..fd16f685ad0af 100644 --- a/clang/test/CodeGenHIP/printf_nonhostcall.cpp +++ b/clang/test/CodeGenHIP/printf_nonhostcall.cpp @@ -300,7 +300,7 @@ __device__ _BitInt(128) Int128 = 45637; // CHECK-NEXT: [[TMP19:%.*]] = zext i44 [[LOADEDV2]] to i64 // CHECK-NEXT: store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8 // CHECK-NEXT: [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 8 -// CHECK-NEXT: store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], align 8 +// CHECK-NEXT: store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], align 16 // CHECK-NEXT: [[PRINTBUFFNEXTPTR12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], i32 16 // CHECK-NEXT: br label [[END_BLOCK]] // @@ -359,7 +359,7 @@ __device__ _BitInt(128) Int128 = 45637; // CHECK_CONSTRAINED-NEXT: [[TMP19:%.*]] = zext i44 [[LOADEDV2]] to i64 // CHECK_CONSTRAINED-NEXT: store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8 // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 8 -// CHECK_CONSTRAINED-NEXT: store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], align 8 +// CHECK_CONSTRAINED-NEXT: store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], align 16 // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], i32 16 // CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]] // diff --git a/clang/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl b/clang/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl index 858a7db574f38..c39d22c16606d 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 amdgpu -emit-llvm -o - | FileCheck %s // RUN: %clang_cc1 %s -O0 -triple amdgpu---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-i64:64-i128:128-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/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp index 4502759417c5a..d0cb983fe8ed9 100644 --- a/llvm/lib/IR/AutoUpgrade.cpp +++ b/llvm/lib/IR/AutoUpgrade.cpp @@ -7163,6 +7163,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 the alignment of i128, which used to be inherited from the i64 + // entry. Must come after the address space upgrades above, which rely on + // matching against the tail of the string. + if (!DL.contains("-i128") && !DL.starts_with("i128")) + Res.append("-i128:128"); } // Upgrade the ELF mangling mode. diff --git a/llvm/lib/TargetParser/TargetDataLayout.cpp b/llvm/lib/TargetParser/TargetDataLayout.cpp index 8b6f46642e4fa..1e74aa68cce16 100644 --- a/llvm/lib/TargetParser/TargetDataLayout.cpp +++ b/llvm/lib/TargetParser/TargetDataLayout.cpp @@ -273,10 +273,16 @@ static std::string computeAMDDataLayout(const Triple &TT) { // (address space 7), and 128-bit non-integral buffer resourcees (address // space 8) which cannot be non-trivilally accessed by LLVM memory operations // like getelementptr. + // + // i128 is aligned to 16 bytes to match the ABI implemented by Clang, whose + // AMDGPUTargetInfo leaves Int128Align at its 128-bit default. Without an + // explicit entry the alignment would be inherited from i64:64, and any + // frontend that lays out aggregates itself would disagree with the layout + // LLVM computes for the corresponding LLVM struct type. 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-" - "v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9"; + "i128:128-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"; } static std::string computeRISCVDataLayout(const Triple &TT, StringRef ABIName) { diff --git a/llvm/test/Analysis/CostModel/AMDGPU/load-to-trunc.ll b/llvm/test/Analysis/CostModel/AMDGPU/load-to-trunc.ll index 92412d706d224..6ed945682b9a2 100644 --- a/llvm/test/Analysis/CostModel/AMDGPU/load-to-trunc.ll +++ b/llvm/test/Analysis/CostModel/AMDGPU/load-to-trunc.ll @@ -8,7 +8,7 @@ ; Check that cost is 1 for unusual load to register sized load. define i32 @loadUnusualIntegerWithTrunc(ptr %ptr) { ; CHECK-LABEL: 'loadUnusualIntegerWithTrunc' -; CHECK-NEXT: Cost Model: Found an estimated cost of 1 for instruction: %out = load i128, ptr %ptr, align 8 +; CHECK-NEXT: Cost Model: Found an estimated cost of 1 for instruction: %out = load i128, ptr %ptr, align 16 ; CHECK-NEXT: Cost Model: Found an estimated cost of 0 for instruction: %trunc = trunc i128 %out to i32 ; CHECK-NEXT: Cost Model: Found an estimated cost of 1 for instruction: ret i32 %trunc ; @@ -19,7 +19,7 @@ define i32 @loadUnusualIntegerWithTrunc(ptr %ptr) { define i128 @loadUnusualInteger(ptr %ptr) { ; CHECK-LABEL: 'loadUnusualInteger' -; CHECK-NEXT: Cost Model: Found an estimated cost of 2 for instruction: %out = load i128, ptr %ptr, align 8 +; CHECK-NEXT: Cost Model: Found an estimated cost of 2 for instruction: %out = load i128, ptr %ptr, align 16 ; CHECK-NEXT: Cost Model: Found an estimated cost of 1 for instruction: ret i128 %out ; %out = load i128, ptr %ptr diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/function-returns.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/function-returns.ll index 898e987f4a826..706b6580b1ab0 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/function-returns.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/function-returns.ll @@ -297,7 +297,7 @@ define i65 @i65_func_void() #0 { ; CHECK-LABEL: name: i65_func_void ; CHECK: bb.1 (%ir-block.0): ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i65) = G_LOAD [[DEF]](p1) :: (load (i65) from `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i65) = G_LOAD [[DEF]](p1) :: (load (i65) from `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: [[ANYEXT:%[0-9]+]]:_(i96) = G_ANYEXT [[LOAD]](i65) ; CHECK-NEXT: [[UV:%[0-9]+]]:_(i32), [[UV1:%[0-9]+]]:_(i32), [[UV2:%[0-9]+]]:_(i32) = G_UNMERGE_VALUES [[ANYEXT]](i96) ; CHECK-NEXT: $vgpr0 = COPY [[UV]](i32) @@ -312,7 +312,7 @@ define signext i65 @i65_signext_func_void() #0 { ; CHECK-LABEL: name: i65_signext_func_void ; CHECK: bb.1 (%ir-block.0): ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i65) = G_LOAD [[DEF]](p1) :: (load (i65) from `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i65) = G_LOAD [[DEF]](p1) :: (load (i65) from `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: [[SEXT:%[0-9]+]]:_(i96) = G_SEXT [[LOAD]](i65) ; CHECK-NEXT: [[UV:%[0-9]+]]:_(i32), [[UV1:%[0-9]+]]:_(i32), [[UV2:%[0-9]+]]:_(i32) = G_UNMERGE_VALUES [[SEXT]](i96) ; CHECK-NEXT: $vgpr0 = COPY [[UV]](i32) @@ -327,7 +327,7 @@ define zeroext i65 @i65_zeroext_func_void() #0 { ; CHECK-LABEL: name: i65_zeroext_func_void ; CHECK: bb.1 (%ir-block.0): ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i65) = G_LOAD [[DEF]](p1) :: (load (i65) from `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i65) = G_LOAD [[DEF]](p1) :: (load (i65) from `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: [[ZEXT:%[0-9]+]]:_(i96) = G_ZEXT [[LOAD]](i65) ; CHECK-NEXT: [[UV:%[0-9]+]]:_(i32), [[UV1:%[0-9]+]]:_(i32), [[UV2:%[0-9]+]]:_(i32) = G_UNMERGE_VALUES [[ZEXT]](i96) ; CHECK-NEXT: $vgpr0 = COPY [[UV]](i32) @@ -1157,7 +1157,7 @@ define i1022 @i1022_func_void() #0 { ; CHECK-LABEL: name: i1022_func_void ; CHECK: bb.1 (%ir-block.0): ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i1022) = G_LOAD [[DEF]](p1) :: (load (i1022) from `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i1022) = G_LOAD [[DEF]](p1) :: (load (i1022) from `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: [[ANYEXT:%[0-9]+]]:_(i1024) = G_ANYEXT [[LOAD]](i1022) ; CHECK-NEXT: [[UV:%[0-9]+]]:_(i32), [[UV1:%[0-9]+]]:_(i32), [[UV2:%[0-9]+]]:_(i32), [[UV3:%[0-9]+]]:_(i32), [[UV4:%[0-9]+]]:_(i32), [[UV5:%[0-9]+]]:_(i32), [[UV6:%[0-9]+]]:_(i32), [[UV7:%[0-9]+]]:_(i32), [[UV8:%[0-9]+]]:_(i32), [[UV9:%[0-9]+]]:_(i32), [[UV10:%[0-9]+]]:_(i32), [[UV11:%[0-9]+]]:_(i32), [[UV12:%[0-9]+]]:_(i32), [[UV13:%[0-9]+]]:_(i32), [[UV14:%[0-9]+]]:_(i32), [[UV15:%[0-9]+]]:_(i32), [[UV16:%[0-9]+]]:_(i32), [[UV17:%[0-9]+]]:_(i32), [[UV18:%[0-9]+]]:_(i32), [[UV19:%[0-9]+]]:_(i32), [[UV20:%[0-9]+]]:_(i32), [[UV21:%[0-9]+]]:_(i32), [[UV22:%[0-9]+]]:_(i32), [[UV23:%[0-9]+]]:_(i32), [[UV24:%[0-9]+]]:_(i32), [[UV25:%[0-9]+]]:_(i32), [[UV26:%[0-9]+]]:_(i32), [[UV27:%[0-9]+]]:_(i32), [[UV28:%[0-9]+]]:_(i32), [[UV29:%[0-9]+]]:_(i32), [[UV30:%[0-9]+]]:_(i32), [[UV31:%[0-9]+]]:_(i32) = G_UNMERGE_VALUES [[ANYEXT]](i1024) ; CHECK-NEXT: $vgpr0 = COPY [[UV]](i32) @@ -1201,7 +1201,7 @@ define signext i1022 @i1022_signext_func_void() #0 { ; CHECK-LABEL: name: i1022_signext_func_void ; CHECK: bb.1 (%ir-block.0): ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i1022) = G_LOAD [[DEF]](p1) :: (load (i1022) from `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i1022) = G_LOAD [[DEF]](p1) :: (load (i1022) from `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: [[SEXT:%[0-9]+]]:_(i1024) = G_SEXT [[LOAD]](i1022) ; CHECK-NEXT: [[UV:%[0-9]+]]:_(i32), [[UV1:%[0-9]+]]:_(i32), [[UV2:%[0-9]+]]:_(i32), [[UV3:%[0-9]+]]:_(i32), [[UV4:%[0-9]+]]:_(i32), [[UV5:%[0-9]+]]:_(i32), [[UV6:%[0-9]+]]:_(i32), [[UV7:%[0-9]+]]:_(i32), [[UV8:%[0-9]+]]:_(i32), [[UV9:%[0-9]+]]:_(i32), [[UV10:%[0-9]+]]:_(i32), [[UV11:%[0-9]+]]:_(i32), [[UV12:%[0-9]+]]:_(i32), [[UV13:%[0-9]+]]:_(i32), [[UV14:%[0-9]+]]:_(i32), [[UV15:%[0-9]+]]:_(i32), [[UV16:%[0-9]+]]:_(i32), [[UV17:%[0-9]+]]:_(i32), [[UV18:%[0-9]+]]:_(i32), [[UV19:%[0-9]+]]:_(i32), [[UV20:%[0-9]+]]:_(i32), [[UV21:%[0-9]+]]:_(i32), [[UV22:%[0-9]+]]:_(i32), [[UV23:%[0-9]+]]:_(i32), [[UV24:%[0-9]+]]:_(i32), [[UV25:%[0-9]+]]:_(i32), [[UV26:%[0-9]+]]:_(i32), [[UV27:%[0-9]+]]:_(i32), [[UV28:%[0-9]+]]:_(i32), [[UV29:%[0-9]+]]:_(i32), [[UV30:%[0-9]+]]:_(i32), [[UV31:%[0-9]+]]:_(i32) = G_UNMERGE_VALUES [[SEXT]](i1024) ; CHECK-NEXT: $vgpr0 = COPY [[UV]](i32) @@ -1245,7 +1245,7 @@ define zeroext i1022 @i1022_zeroext_func_void() #0 { ; CHECK-LABEL: name: i1022_zeroext_func_void ; CHECK: bb.1 (%ir-block.0): ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i1022) = G_LOAD [[DEF]](p1) :: (load (i1022) from `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: [[LOAD:%[0-9]+]]:_(i1022) = G_LOAD [[DEF]](p1) :: (load (i1022) from `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: [[ZEXT:%[0-9]+]]:_(i1024) = G_ZEXT [[LOAD]](i1022) ; CHECK-NEXT: [[UV:%[0-9]+]]:_(i32), [[UV1:%[0-9]+]]:_(i32), [[UV2:%[0-9]+]]:_(i32), [[UV3:%[0-9]+]]:_(i32), [[UV4:%[0-9]+]]:_(i32), [[UV5:%[0-9]+]]:_(i32), [[UV6:%[0-9]+]]:_(i32), [[UV7:%[0-9]+]]:_(i32), [[UV8:%[0-9]+]]:_(i32), [[UV9:%[0-9]+]]:_(i32), [[UV10:%[0-9]+]]:_(i32), [[UV11:%[0-9]+]]:_(i32), [[UV12:%[0-9]+]]:_(i32), [[UV13:%[0-9]+]]:_(i32), [[UV14:%[0-9]+]]:_(i32), [[UV15:%[0-9]+]]:_(i32), [[UV16:%[0-9]+]]:_(i32), [[UV17:%[0-9]+]]:_(i32), [[UV18:%[0-9]+]]:_(i32), [[UV19:%[0-9]+]]:_(i32), [[UV20:%[0-9]+]]:_(i32), [[UV21:%[0-9]+]]:_(i32), [[UV22:%[0-9]+]]:_(i32), [[UV23:%[0-9]+]]:_(i32), [[UV24:%[0-9]+]]:_(i32), [[UV25:%[0-9]+]]:_(i32), [[UV26:%[0-9]+]]:_(i32), [[UV27:%[0-9]+]]:_(i32), [[UV28:%[0-9]+]]:_(i32), [[UV29:%[0-9]+]]:_(i32), [[UV30:%[0-9]+]]:_(i32), [[UV31:%[0-9]+]]:_(i32) = G_UNMERGE_VALUES [[ZEXT]](i1024) ; CHECK-NEXT: $vgpr0 = COPY [[UV]](i32) diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll index 25f952df63f29..c3f46dd877115 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll @@ -412,7 +412,7 @@ define void @void_func_i95(i95 %arg0) #0 { ; CHECK-NEXT: [[MV:%[0-9]+]]:_(i96) = G_MERGE_VALUES [[COPY]](i32), [[COPY1]](i32), [[COPY2]](i32) ; CHECK-NEXT: [[TRUNC:%[0-9]+]]:_(i95) = G_TRUNC [[MV]](i96) ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: G_STORE [[TRUNC]](i95), [[DEF]](p1) :: (store (i95) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: G_STORE [[TRUNC]](i95), [[DEF]](p1) :: (store (i95) into `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: SI_RETURN store i95 %arg0, ptr addrspace(1) poison ret void @@ -432,7 +432,7 @@ define void @void_func_i95_zeroext(i95 zeroext %arg0) #0 { ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF ; CHECK-NEXT: [[ZEXT:%[0-9]+]]:_(i96) = G_ZEXT [[TRUNC]](i95) ; CHECK-NEXT: [[ADD:%[0-9]+]]:_(i96) = G_ADD [[ZEXT]], [[C]] - ; CHECK-NEXT: G_STORE [[ADD]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: G_STORE [[ADD]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: SI_RETURN %ext = zext i95 %arg0 to i96 %add = add i96 %ext, 12 @@ -454,7 +454,7 @@ define void @void_func_i95_signext(i95 signext %arg0) #0 { ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF ; CHECK-NEXT: [[SEXT:%[0-9]+]]:_(i96) = G_SEXT [[TRUNC]](i95) ; CHECK-NEXT: [[ADD:%[0-9]+]]:_(i96) = G_ADD [[SEXT]], [[C]] - ; CHECK-NEXT: G_STORE [[ADD]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: G_STORE [[ADD]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: SI_RETURN %ext = sext i95 %arg0 to i96 %add = add i96 %ext, 12 @@ -472,7 +472,7 @@ define void @void_func_i96(i96 %arg0) #0 { ; CHECK-NEXT: [[COPY2:%[0-9]+]]:_(i32) = COPY $vgpr2 ; CHECK-NEXT: [[MV:%[0-9]+]]:_(i96) = G_MERGE_VALUES [[COPY]](i32), [[COPY1]](i32), [[COPY2]](i32) ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: G_STORE [[MV]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: G_STORE [[MV]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: SI_RETURN store i96 %arg0, ptr addrspace(1) poison ret void @@ -2875,7 +2875,7 @@ define void @void_func_i96_inreg(i96 inreg %arg0) #0 { ; CHECK-NEXT: [[COPY2:%[0-9]+]]:_(i32) = COPY $sgpr18 ; CHECK-NEXT: [[MV:%[0-9]+]]:_(i96) = G_MERGE_VALUES [[COPY]](i32), [[COPY1]](i32), [[COPY2]](i32) ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: G_STORE [[MV]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: G_STORE [[MV]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 16, addrspace 1) ; CHECK-NEXT: SI_RETURN store i96 %arg0, ptr addrspace(1) poison ret void @@ -2892,7 +2892,7 @@ define void @void_func_i128_inreg(i128 inreg %arg0) #0 { ; CHECK-NEXT: [[COPY3:%[0-9]+]]:_(i32) = COPY $sgpr19 ; CHECK-NEXT: [[MV:%[0-9]+]]:_(i128) = G_MERGE_VALUES [[COPY]](i32), [[COPY1]](i32), [[COPY2]](i32), [[COPY3]](i32) ; CHECK-NEXT: [[DEF:%[0-9]+]]:_(p1) = G_IMPLICIT_DEF - ; CHECK-NEXT: G_STORE [[MV]](i128), [[DEF]](p1) :: (store (i128) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; CHECK-NEXT: G_STORE [[MV]](i128), [[DEF]](p1) :: (store (i128) into `ptr addrspace(1) poison`, addrspace 1) ; CHECK-NEXT: SI_RETURN store i128 %arg0, ptr addrspace(1) poison ret void diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.s.buffer.load.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.s.buffer.load.ll index 8e985b16d322f..0f466f4272726 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.s.buffer.load.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.s.buffer.load.ll @@ -659,10 +659,10 @@ define amdgpu_ps void @s_buffer_load_i96_vgpr_offset(<4 x i32> inreg %rsrc, i32 ; GFX7-NEXT: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF ; GFX7-NEXT: [[C:%[0-9]+]]:sgpr(i32) = G_CONSTANT i32 0 ; GFX7-NEXT: [[C1:%[0-9]+]]:vgpr(i32) = G_CONSTANT i32 0 - ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(s128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s96), align 8) + ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(s128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s96), align 16) ; GFX7-NEXT: [[COPY5:%[0-9]+]]:vgpr(s128) = COPY [[AMDGPU_BUFFER_LOAD]](s128) ; GFX7-NEXT: [[TRUNC:%[0-9]+]]:vgpr(i96) = G_TRUNC [[COPY5]](s128) - ; GFX7-NEXT: G_STORE [[TRUNC]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; GFX7-NEXT: G_STORE [[TRUNC]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 16, addrspace 1) ; GFX7-NEXT: S_ENDPGM 0 ; ; GFX1200_1250-LABEL: name: s_buffer_load_i96_vgpr_offset @@ -678,9 +678,9 @@ define amdgpu_ps void @s_buffer_load_i96_vgpr_offset(<4 x i32> inreg %rsrc, i32 ; GFX1200_1250-NEXT: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF ; GFX1200_1250-NEXT: [[C:%[0-9]+]]:sgpr(i32) = G_CONSTANT i32 0 ; GFX1200_1250-NEXT: [[C1:%[0-9]+]]:vgpr(i32) = G_CONSTANT i32 0 - ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i96) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s96), align 8) + ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i96) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s96), align 16) ; GFX1200_1250-NEXT: [[COPY5:%[0-9]+]]:vgpr(i96) = COPY [[AMDGPU_BUFFER_LOAD]](i96) - ; GFX1200_1250-NEXT: G_STORE [[COPY5]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; GFX1200_1250-NEXT: G_STORE [[COPY5]](i96), [[DEF]](p1) :: (store (i96) into `ptr addrspace(1) poison`, align 16, addrspace 1) ; GFX1200_1250-NEXT: S_ENDPGM 0 %val = call i96 @llvm.amdgcn.s.buffer.load.i96(<4 x i32> %rsrc, i32 %soffset, i32 0) store i96 %val, ptr addrspace(1) poison @@ -702,14 +702,14 @@ define amdgpu_ps void @s_buffer_load_i256_vgpr_offset(<4 x i32> inreg %rsrc, i32 ; GFX7-NEXT: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF ; GFX7-NEXT: [[C:%[0-9]+]]:sgpr(i32) = G_CONSTANT i32 0 ; GFX7-NEXT: [[C1:%[0-9]+]]:vgpr(i32) = G_CONSTANT i32 0 - ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s128), align 8) - ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD1:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 16, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 16, align 8) + ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s128)) + ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD1:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 16, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 16) ; GFX7-NEXT: [[MV:%[0-9]+]]:vgpr(i256) = G_MERGE_VALUES [[AMDGPU_BUFFER_LOAD]](i128), [[AMDGPU_BUFFER_LOAD1]](i128) ; GFX7-NEXT: [[UV:%[0-9]+]]:vgpr(s128), [[UV1:%[0-9]+]]:vgpr(s128) = G_UNMERGE_VALUES [[MV]](i256) - ; GFX7-NEXT: G_STORE [[UV]](s128), [[DEF]](p1) :: (store (s128) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; GFX7-NEXT: G_STORE [[UV]](s128), [[DEF]](p1) :: (store (s128) into `ptr addrspace(1) poison`, addrspace 1) ; GFX7-NEXT: [[C2:%[0-9]+]]:sgpr(i64) = G_CONSTANT i64 16 ; GFX7-NEXT: [[PTR_ADD:%[0-9]+]]:sgpr(p1) = nuw inbounds G_PTR_ADD [[DEF]], [[C2]](i64) - ; GFX7-NEXT: G_STORE [[UV1]](s128), [[PTR_ADD]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 16, align 8, addrspace 1) + ; GFX7-NEXT: G_STORE [[UV1]](s128), [[PTR_ADD]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 16, addrspace 1) ; GFX7-NEXT: S_ENDPGM 0 ; ; GFX1200_1250-LABEL: name: s_buffer_load_i256_vgpr_offset @@ -725,14 +725,14 @@ define amdgpu_ps void @s_buffer_load_i256_vgpr_offset(<4 x i32> inreg %rsrc, i32 ; GFX1200_1250-NEXT: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF ; GFX1200_1250-NEXT: [[C:%[0-9]+]]:sgpr(i32) = G_CONSTANT i32 0 ; GFX1200_1250-NEXT: [[C1:%[0-9]+]]:vgpr(i32) = G_CONSTANT i32 0 - ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s128), align 8) - ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD1:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 16, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 16, align 8) + ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s128)) + ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD1:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 16, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 16) ; GFX1200_1250-NEXT: [[MV:%[0-9]+]]:vgpr(i256) = G_MERGE_VALUES [[AMDGPU_BUFFER_LOAD]](i128), [[AMDGPU_BUFFER_LOAD1]](i128) ; GFX1200_1250-NEXT: [[UV:%[0-9]+]]:vgpr(s128), [[UV1:%[0-9]+]]:vgpr(s128) = G_UNMERGE_VALUES [[MV]](i256) - ; GFX1200_1250-NEXT: G_STORE [[UV]](s128), [[DEF]](p1) :: (store (s128) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; GFX1200_1250-NEXT: G_STORE [[UV]](s128), [[DEF]](p1) :: (store (s128) into `ptr addrspace(1) poison`, addrspace 1) ; GFX1200_1250-NEXT: [[C2:%[0-9]+]]:sgpr(i64) = G_CONSTANT i64 16 ; GFX1200_1250-NEXT: [[PTR_ADD:%[0-9]+]]:sgpr(p1) = nuw inbounds G_PTR_ADD [[DEF]], [[C2]](i64) - ; GFX1200_1250-NEXT: G_STORE [[UV1]](s128), [[PTR_ADD]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 16, align 8, addrspace 1) + ; GFX1200_1250-NEXT: G_STORE [[UV1]](s128), [[PTR_ADD]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 16, addrspace 1) ; GFX1200_1250-NEXT: S_ENDPGM 0 %val = call i256 @llvm.amdgcn.s.buffer.load.i256(<4 x i32> %rsrc, i32 %soffset, i32 0) store i256 %val, ptr addrspace(1) poison @@ -754,22 +754,22 @@ define amdgpu_ps void @s_buffer_load_i512_vgpr_offset(<4 x i32> inreg %rsrc, i32 ; GFX7-NEXT: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF ; GFX7-NEXT: [[C:%[0-9]+]]:sgpr(i32) = G_CONSTANT i32 0 ; GFX7-NEXT: [[C1:%[0-9]+]]:vgpr(i32) = G_CONSTANT i32 0 - ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s128), align 8) - ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD1:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 16, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 16, align 8) - ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD2:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 32, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 32, align 8) - ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD3:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 48, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 48, align 8) + ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s128)) + ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD1:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 16, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 16) + ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD2:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 32, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 32) + ; GFX7-NEXT: [[AMDGPU_BUFFER_LOAD3:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 48, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 48) ; GFX7-NEXT: [[MV:%[0-9]+]]:vgpr(i512) = G_MERGE_VALUES [[AMDGPU_BUFFER_LOAD]](i128), [[AMDGPU_BUFFER_LOAD1]](i128), [[AMDGPU_BUFFER_LOAD2]](i128), [[AMDGPU_BUFFER_LOAD3]](i128) ; GFX7-NEXT: [[UV:%[0-9]+]]:vgpr(s128), [[UV1:%[0-9]+]]:vgpr(s128), [[UV2:%[0-9]+]]:vgpr(s128), [[UV3:%[0-9]+]]:vgpr(s128) = G_UNMERGE_VALUES [[MV]](i512) - ; GFX7-NEXT: G_STORE [[UV]](s128), [[DEF]](p1) :: (store (s128) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; GFX7-NEXT: G_STORE [[UV]](s128), [[DEF]](p1) :: (store (s128) into `ptr addrspace(1) poison`, addrspace 1) ; GFX7-NEXT: [[C2:%[0-9]+]]:sgpr(i64) = G_CONSTANT i64 16 ; GFX7-NEXT: [[PTR_ADD:%[0-9]+]]:sgpr(p1) = nuw inbounds G_PTR_ADD [[DEF]], [[C2]](i64) - ; GFX7-NEXT: G_STORE [[UV1]](s128), [[PTR_ADD]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 16, align 8, addrspace 1) + ; GFX7-NEXT: G_STORE [[UV1]](s128), [[PTR_ADD]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 16, addrspace 1) ; GFX7-NEXT: [[C3:%[0-9]+]]:sgpr(i64) = G_CONSTANT i64 32 ; GFX7-NEXT: [[PTR_ADD1:%[0-9]+]]:sgpr(p1) = nuw inbounds G_PTR_ADD [[DEF]], [[C3]](i64) - ; GFX7-NEXT: G_STORE [[UV2]](s128), [[PTR_ADD1]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 32, align 8, addrspace 1) + ; GFX7-NEXT: G_STORE [[UV2]](s128), [[PTR_ADD1]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 32, addrspace 1) ; GFX7-NEXT: [[C4:%[0-9]+]]:sgpr(i64) = G_CONSTANT i64 48 ; GFX7-NEXT: [[PTR_ADD2:%[0-9]+]]:sgpr(p1) = nuw inbounds G_PTR_ADD [[DEF]], [[C4]](i64) - ; GFX7-NEXT: G_STORE [[UV3]](s128), [[PTR_ADD2]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 48, align 8, addrspace 1) + ; GFX7-NEXT: G_STORE [[UV3]](s128), [[PTR_ADD2]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 48, addrspace 1) ; GFX7-NEXT: S_ENDPGM 0 ; ; GFX1200_1250-LABEL: name: s_buffer_load_i512_vgpr_offset @@ -785,22 +785,22 @@ define amdgpu_ps void @s_buffer_load_i512_vgpr_offset(<4 x i32> inreg %rsrc, i32 ; GFX1200_1250-NEXT: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF ; GFX1200_1250-NEXT: [[C:%[0-9]+]]:sgpr(i32) = G_CONSTANT i32 0 ; GFX1200_1250-NEXT: [[C1:%[0-9]+]]:vgpr(i32) = G_CONSTANT i32 0 - ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s128), align 8) - ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD1:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 16, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 16, align 8) - ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD2:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 32, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 32, align 8) - ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD3:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 48, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 48, align 8) + ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 0, 0, 0 :: (dereferenceable invariant load (s128)) + ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD1:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 16, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 16) + ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD2:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 32, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 32) + ; GFX1200_1250-NEXT: [[AMDGPU_BUFFER_LOAD3:%[0-9]+]]:vgpr(i128) = G_AMDGPU_BUFFER_LOAD [[BUILD_VECTOR]](<4 x i32>), [[C1]](i32), [[COPY4]], [[C]], 48, 0, 0 :: (dereferenceable invariant load (s128) from unknown-address + 48) ; GFX1200_1250-NEXT: [[MV:%[0-9]+]]:vgpr(i512) = G_MERGE_VALUES [[AMDGPU_BUFFER_LOAD]](i128), [[AMDGPU_BUFFER_LOAD1]](i128), [[AMDGPU_BUFFER_LOAD2]](i128), [[AMDGPU_BUFFER_LOAD3]](i128) ; GFX1200_1250-NEXT: [[UV:%[0-9]+]]:vgpr(s128), [[UV1:%[0-9]+]]:vgpr(s128), [[UV2:%[0-9]+]]:vgpr(s128), [[UV3:%[0-9]+]]:vgpr(s128) = G_UNMERGE_VALUES [[MV]](i512) - ; GFX1200_1250-NEXT: G_STORE [[UV]](s128), [[DEF]](p1) :: (store (s128) into `ptr addrspace(1) poison`, align 8, addrspace 1) + ; GFX1200_1250-NEXT: G_STORE [[UV]](s128), [[DEF]](p1) :: (store (s128) into `ptr addrspace(1) poison`, addrspace 1) ; GFX1200_1250-NEXT: [[C2:%[0-9]+]]:sgpr(i64) = G_CONSTANT i64 16 ; GFX1200_1250-NEXT: [[PTR_ADD:%[0-9]+]]:sgpr(p1) = nuw inbounds G_PTR_ADD [[DEF]], [[C2]](i64) - ; GFX1200_1250-NEXT: G_STORE [[UV1]](s128), [[PTR_ADD]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 16, align 8, addrspace 1) + ; GFX1200_1250-NEXT: G_STORE [[UV1]](s128), [[PTR_ADD]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 16, addrspace 1) ; GFX1200_1250-NEXT: [[C3:%[0-9]+]]:sgpr(i64) = G_CONSTANT i64 32 ; GFX1200_1250-NEXT: [[PTR_ADD1:%[0-9]+]]:sgpr(p1) = nuw inbounds G_PTR_ADD [[DEF]], [[C3]](i64) - ; GFX1200_1250-NEXT: G_STORE [[UV2]](s128), [[PTR_ADD1]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 32, align 8, addrspace 1) + ; GFX1200_1250-NEXT: G_STORE [[UV2]](s128), [[PTR_ADD1]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 32, addrspace 1) ; GFX1200_1250-NEXT: [[C4:%[0-9]+]]:sgpr(i64) = G_CONSTANT i64 48 ; GFX1200_1250-NEXT: [[PTR_ADD2:%[0-9]+]]:sgpr(p1) = nuw inbounds G_PTR_ADD [[DEF]], [[C4]](i64) - ; GFX1200_1250-NEXT: G_STORE [[UV3]](s128), [[PTR_ADD2]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 48, align 8, addrspace 1) + ; GFX1200_1250-NEXT: G_STORE [[UV3]](s128), [[PTR_ADD2]](p1) :: (store (s128) into `ptr addrspace(1) poison` + 48, addrspace 1) ; GFX1200_1250-NEXT: S_ENDPGM 0 %val = call i512 @llvm.amdgcn.s.buffer.load.i512(<4 x i32> %rsrc, i32 %soffset, i32 0) store i512 %val, ptr addrspace(1) poison diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/trunc-brc.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/trunc-brc.ll index 116719a6e3e7b..1fe8390ea639f 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/trunc-brc.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/trunc-brc.ll @@ -57,12 +57,13 @@ define amdgpu_ps void @s_trunc_i96_to_i64(ptr addrspace(1) inreg %src, ptr addrs ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:sreg_32 = COPY $sgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:sreg_32 = COPY $sgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64_xexec_xnull = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: [[S_LOAD_DWORDX2_IMM:%[0-9]+]]:sreg_64_xexec = S_LOAD_DWORDX2_IMM [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<2 x s32>) from %ir.src, addrspace 1) - ; GFX-950-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[REG_SEQUENCE]], 8, 0 :: ("amdgpu-noclobber" load (s32) from %ir.src + 8, align 8, addrspace 1) - ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX2_IMM]].sub0 - ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX2_IMM]].sub1 - ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:sgpr_96 = REG_SEQUENCE [[COPY4]], %subreg.sub0, [[COPY5]], %subreg.sub1, [[S_LOAD_DWORD_IMM]], %subreg.sub2 - ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:sreg_64 = COPY [[REG_SEQUENCE2]].sub0_sub1, debug-location !4 + ; GFX-950-NEXT: [[S_LOAD_DWORDX4_IMM:%[0-9]+]]:sgpr_128 = S_LOAD_DWORDX4_IMM [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x s32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub0 + ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub1 + ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub2 + ; GFX-950-NEXT: [[COPY7:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub3 + ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:sgpr_96 = REG_SEQUENCE [[COPY4]], %subreg.sub0, [[COPY5]], %subreg.sub1, [[COPY6]], %subreg.sub2 + ; GFX-950-NEXT: [[COPY8:%[0-9]+]]:sreg_64 = COPY [[REG_SEQUENCE2]].sub0_sub1, debug-location !4 %val = load i96, ptr addrspace(1) %src %trunc = trunc i96 %val to i64, !dbg !4 store i64 %trunc, ptr addrspace(1) %dst @@ -80,7 +81,7 @@ define void @v_trunc_i96_to_i64(ptr addrspace(1) %src, ptr addrspace(1) %dst) { ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY $vgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:vgpr_32 = COPY $vgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:vreg_64_align2 = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX3_:%[0-9]+]]:vreg_96_align2 = GLOBAL_LOAD_DWORDX3 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<3 x i32>) from %ir.src, align 8, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX3_:%[0-9]+]]:vreg_96_align2 = GLOBAL_LOAD_DWORDX3 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<3 x i32>) from %ir.src, align 16, addrspace 1) ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX3_]].sub0_sub1, debug-location !4 %val = load i96, ptr addrspace(1) %src %trunc = trunc i96 %val to i64, !dbg !4 @@ -101,8 +102,8 @@ define amdgpu_ps void @s_trunc_i128_to_i96(ptr addrspace(1) inreg %src, ptr addr ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:sreg_32 = COPY $sgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:sreg_32 = COPY $sgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64_xexec_xnull = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: early-clobber %9:sgpr_128 = S_LOAD_DWORDX4_IMM_ec [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sgpr_96 = COPY %9.sub0_sub1_sub2, debug-location !4 + ; GFX-950-NEXT: [[S_LOAD_DWORDX4_IMM:%[0-9]+]]:sgpr_128 = S_LOAD_DWORDX4_IMM [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sgpr_96 = COPY [[S_LOAD_DWORDX4_IMM]].sub0_sub1_sub2, debug-location !4 %val = load i128, ptr addrspace(1) %src %trunc = trunc i128 %val to i96, !dbg !4 store i96 %trunc, ptr addrspace(1) %dst @@ -120,7 +121,7 @@ define void @v_trunc_i128_to_i96(ptr addrspace(1) %src, ptr addrspace(1) %dst) { ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY $vgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:vgpr_32 = COPY $vgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:vreg_64_align2 = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, align 8, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, addrspace 1) ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:vreg_96_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub0_sub1_sub2, debug-location !4 %val = load i128, ptr addrspace(1) %src %trunc = trunc i128 %val to i96, !dbg !4 @@ -141,12 +142,12 @@ define amdgpu_ps void @s_trunc_i160_to_i128(ptr addrspace(1) inreg %src, ptr add ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:sreg_32 = COPY $sgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:sreg_32 = COPY $sgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64_xexec_xnull = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: early-clobber %10:sgpr_128 = S_LOAD_DWORDX4_IMM_ec [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[REG_SEQUENCE]], 16, 0 :: ("amdgpu-noclobber" load (i32) from %ir.src + 16, align 8, addrspace 1) - ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sreg_32 = COPY %10.sub0 - ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:sreg_32 = COPY %10.sub1 - ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:sreg_32 = COPY %10.sub2 - ; GFX-950-NEXT: [[COPY7:%[0-9]+]]:sreg_32 = COPY %10.sub3 + ; GFX-950-NEXT: [[S_LOAD_DWORDX4_IMM:%[0-9]+]]:sgpr_128 = S_LOAD_DWORDX4_IMM [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[REG_SEQUENCE]], 16, 0 :: ("amdgpu-noclobber" load (i32) from %ir.src + 16, align 16, addrspace 1) + ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub0 + ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub1 + ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub2 + ; GFX-950-NEXT: [[COPY7:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub3 ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:sgpr_160 = REG_SEQUENCE [[COPY4]], %subreg.sub0, [[COPY5]], %subreg.sub1, [[COPY6]], %subreg.sub2, [[COPY7]], %subreg.sub3, [[S_LOAD_DWORD_IMM]], %subreg.sub4 ; GFX-950-NEXT: [[COPY8:%[0-9]+]]:sgpr_128 = COPY [[REG_SEQUENCE2]].sub0_sub1_sub2_sub3, debug-location !4 %val = load i160, ptr addrspace(1) %src @@ -166,8 +167,8 @@ define void @v_trunc_i160_to_i128(ptr addrspace(1) %src, ptr addrspace(1) %dst) ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY $vgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:vgpr_32 = COPY $vgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:vreg_64_align2 = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORD:%[0-9]+]]:vgpr_32 = GLOBAL_LOAD_DWORD [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (i32) from %ir.src + 16, align 8, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORD:%[0-9]+]]:vgpr_32 = GLOBAL_LOAD_DWORD [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (i32) from %ir.src + 16, align 16, addrspace 1) ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub0 ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub1 ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub2 @@ -193,12 +194,12 @@ define amdgpu_ps void @s_trunc_i192_to_i160(ptr addrspace(1) inreg %src, ptr add ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:sreg_32 = COPY $sgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:sreg_32 = COPY $sgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64_xexec_xnull = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: early-clobber %18:sgpr_128 = S_LOAD_DWORDX4_IMM_ec [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[S_LOAD_DWORDX2_IMM:%[0-9]+]]:sreg_64_xexec = S_LOAD_DWORDX2_IMM [[REG_SEQUENCE]], 16, 0 :: ("amdgpu-noclobber" load (<2 x i32>) from %ir.src + 16, addrspace 1) - ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sreg_32 = COPY %18.sub0 - ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:sreg_32 = COPY %18.sub1 - ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:sreg_32 = COPY %18.sub2 - ; GFX-950-NEXT: [[COPY7:%[0-9]+]]:sreg_32 = COPY %18.sub3 + ; GFX-950-NEXT: [[S_LOAD_DWORDX4_IMM:%[0-9]+]]:sgpr_128 = S_LOAD_DWORDX4_IMM [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[S_LOAD_DWORDX2_IMM:%[0-9]+]]:sreg_64_xexec = S_LOAD_DWORDX2_IMM [[REG_SEQUENCE]], 16, 0 :: ("amdgpu-noclobber" load (<2 x i32>) from %ir.src + 16, align 16, addrspace 1) + ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub0 + ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub1 + ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub2 + ; GFX-950-NEXT: [[COPY7:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub3 ; GFX-950-NEXT: [[COPY8:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX2_IMM]].sub0 ; GFX-950-NEXT: [[COPY9:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX2_IMM]].sub1 ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:sgpr_192 = REG_SEQUENCE [[COPY4]], %subreg.sub0, [[COPY5]], %subreg.sub1, [[COPY6]], %subreg.sub2, [[COPY7]], %subreg.sub3, [[COPY8]], %subreg.sub4, [[COPY9]], %subreg.sub5 @@ -220,8 +221,8 @@ define void @v_trunc_i192_to_i160(ptr addrspace(1) %src, ptr addrspace(1) %dst) ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY $vgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:vgpr_32 = COPY $vgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:vreg_64_align2 = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:vreg_64_align2 = GLOBAL_LOAD_DWORDX2 [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (<2 x i32>) from %ir.src + 16, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:vreg_64_align2 = GLOBAL_LOAD_DWORDX2 [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (<2 x i32>) from %ir.src + 16, align 16, addrspace 1) ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub0 ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub1 ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub2 @@ -249,17 +250,18 @@ define amdgpu_ps void @s_trunc_i224_to_i192(ptr addrspace(1) inreg %src, ptr add ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:sreg_32 = COPY $sgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:sreg_32 = COPY $sgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64_xexec_xnull = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: early-clobber %16:sgpr_128 = S_LOAD_DWORDX4_IMM_ec [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[S_LOAD_DWORDX2_IMM:%[0-9]+]]:sreg_64_xexec = S_LOAD_DWORDX2_IMM [[REG_SEQUENCE]], 16, 0 :: ("amdgpu-noclobber" load (<2 x s32>) from %ir.src + 16, addrspace 1) - ; GFX-950-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[REG_SEQUENCE]], 24, 0 :: ("amdgpu-noclobber" load (s32) from %ir.src + 24, align 8, addrspace 1) - ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX2_IMM]].sub0 - ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX2_IMM]].sub1 - ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:sreg_32 = COPY %16.sub0 - ; GFX-950-NEXT: [[COPY7:%[0-9]+]]:sreg_32 = COPY %16.sub1 - ; GFX-950-NEXT: [[COPY8:%[0-9]+]]:sreg_32 = COPY %16.sub2 - ; GFX-950-NEXT: [[COPY9:%[0-9]+]]:sreg_32 = COPY %16.sub3 - ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:sgpr_224 = REG_SEQUENCE [[COPY6]], %subreg.sub0, [[COPY7]], %subreg.sub1, [[COPY8]], %subreg.sub2, [[COPY9]], %subreg.sub3, [[COPY4]], %subreg.sub4, [[COPY5]], %subreg.sub5, [[S_LOAD_DWORD_IMM]], %subreg.sub6 - ; GFX-950-NEXT: [[COPY10:%[0-9]+]]:sgpr_192 = COPY [[REG_SEQUENCE2]].sub0_sub1_sub2_sub3_sub4_sub5, debug-location !4 + ; GFX-950-NEXT: [[S_LOAD_DWORDX4_IMM:%[0-9]+]]:sgpr_128 = S_LOAD_DWORDX4_IMM [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[S_LOAD_DWORDX4_IMM1:%[0-9]+]]:sgpr_128 = S_LOAD_DWORDX4_IMM [[REG_SEQUENCE]], 16, 0 :: ("amdgpu-noclobber" load (<4 x s32>) from %ir.src + 16, addrspace 1) + ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM1]].sub0 + ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM1]].sub1 + ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM1]].sub2 + ; GFX-950-NEXT: [[COPY7:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM1]].sub3 + ; GFX-950-NEXT: [[COPY8:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub0 + ; GFX-950-NEXT: [[COPY9:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub1 + ; GFX-950-NEXT: [[COPY10:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub2 + ; GFX-950-NEXT: [[COPY11:%[0-9]+]]:sreg_32 = COPY [[S_LOAD_DWORDX4_IMM]].sub3 + ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:sgpr_224 = REG_SEQUENCE [[COPY8]], %subreg.sub0, [[COPY9]], %subreg.sub1, [[COPY10]], %subreg.sub2, [[COPY11]], %subreg.sub3, [[COPY4]], %subreg.sub4, [[COPY5]], %subreg.sub5, [[COPY6]], %subreg.sub6 + ; GFX-950-NEXT: [[COPY12:%[0-9]+]]:sgpr_192 = COPY [[REG_SEQUENCE2]].sub0_sub1_sub2_sub3_sub4_sub5, debug-location !4 %val = load i224, ptr addrspace(1) %src %trunc = trunc i224 %val to i192, !dbg !4 store i192 %trunc, ptr addrspace(1) %dst @@ -277,8 +279,8 @@ define void @v_trunc_i224_to_i192(ptr addrspace(1) %src, ptr addrspace(1) %dst) ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY $vgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:vgpr_32 = COPY $vgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:vreg_64_align2 = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX3_:%[0-9]+]]:vreg_96_align2 = GLOBAL_LOAD_DWORDX3 [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (<3 x i32>) from %ir.src + 16, align 8, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX3_:%[0-9]+]]:vreg_96_align2 = GLOBAL_LOAD_DWORDX3 [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (<3 x i32>) from %ir.src + 16, align 16, addrspace 1) ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub0 ; GFX-950-NEXT: [[COPY5:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub1 ; GFX-950-NEXT: [[COPY6:%[0-9]+]]:vgpr_32 = COPY [[GLOBAL_LOAD_DWORDX4_]].sub2 @@ -307,7 +309,7 @@ define amdgpu_ps void @s_trunc_i256_to_i224(ptr addrspace(1) inreg %src, ptr add ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:sreg_32 = COPY $sgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:sreg_32 = COPY $sgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64_xexec_xnull = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: early-clobber %20:sgpr_256 = S_LOAD_DWORDX8_IMM_ec [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<8 x i32>) from %ir.src, align 8, addrspace 1) + ; GFX-950-NEXT: early-clobber %20:sgpr_256 = S_LOAD_DWORDX8_IMM_ec [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<8 x i32>) from %ir.src, align 16, addrspace 1) ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sgpr_224 = COPY %20.lo16_hi16_sub1_lo16_sub1_hi16_sub2_lo16_sub2_hi16_sub3_lo16_sub3_hi16_sub4_lo16_sub4_hi16_sub5_lo16_sub5_hi16_sub6_lo16_sub6_hi16, debug-location !4 %val = load i256, ptr addrspace(1) %src %trunc = trunc i256 %val to i224, !dbg !4 @@ -326,8 +328,8 @@ define void @v_trunc_i256_to_i224(ptr addrspace(1) %src, ptr addrspace(1) %dst) ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY $vgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:vgpr_32 = COPY $vgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:vreg_64_align2 = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_1:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 16, align 8, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_1:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 16, addrspace 1) ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:vreg_256_align2 = REG_SEQUENCE [[GLOBAL_LOAD_DWORDX4_]], %subreg.sub0_sub1_sub2_sub3, [[GLOBAL_LOAD_DWORDX4_1]], %subreg.sub4_sub5_sub6_sub7 ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:vreg_224_align2 = COPY [[REG_SEQUENCE2]].lo16_hi16_sub1_lo16_sub1_hi16_sub2_lo16_sub2_hi16_sub3_lo16_sub3_hi16_sub4_lo16_sub4_hi16_sub5_lo16_sub5_hi16_sub6_lo16_sub6_hi16, debug-location !4 %val = load i256, ptr addrspace(1) %src @@ -349,8 +351,8 @@ define amdgpu_ps void @s_trunc_i1024_to_i512(ptr addrspace(1) inreg %src, ptr ad ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:sreg_32 = COPY $sgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:sreg_32 = COPY $sgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64_xexec_xnull = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: early-clobber %20:sgpr_512 = S_LOAD_DWORDX16_IMM_ec [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<16 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: early-clobber %23:sgpr_512 = S_LOAD_DWORDX16_IMM_ec [[REG_SEQUENCE]], 64, 0 :: ("amdgpu-noclobber" load (<16 x i32>) from %ir.src + 64, align 8, addrspace 1) + ; GFX-950-NEXT: early-clobber %20:sgpr_512 = S_LOAD_DWORDX16_IMM_ec [[REG_SEQUENCE]], 0, 0 :: ("amdgpu-noclobber" load (<16 x i32>) from %ir.src, align 16, addrspace 1) + ; GFX-950-NEXT: early-clobber %23:sgpr_512 = S_LOAD_DWORDX16_IMM_ec [[REG_SEQUENCE]], 64, 0 :: ("amdgpu-noclobber" load (<16 x i32>) from %ir.src + 64, align 16, addrspace 1) ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:sgpr_1024 = REG_SEQUENCE %20, %subreg.sub0_sub1_sub2_sub3_sub4_sub5_sub6_sub7_sub8_sub9_sub10_sub11_sub12_sub13_sub14_sub15, %23, %subreg.sub16_sub17_sub18_sub19_sub20_sub21_sub22_sub23_sub24_sub25_sub26_sub27_sub28_sub29_sub30_sub31 ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:sgpr_512 = COPY [[REG_SEQUENCE2]].sub0_sub1_sub2_sub3_sub4_sub5_sub6_sub7_sub8_sub9_sub10_sub11_sub12_sub13_sub14_sub15, debug-location !4 %val = load i1024, ptr addrspace(1) %src @@ -370,15 +372,15 @@ define void @v_trunc_i1024_to_i512(ptr addrspace(1) %src, ptr addrspace(1) %dst) ; GFX-950-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY $vgpr2 ; GFX-950-NEXT: [[COPY3:%[0-9]+]]:vgpr_32 = COPY $vgpr3 ; GFX-950-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:vreg_64_align2 = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[COPY3]], %subreg.sub1 - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_1:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 16, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_2:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 32, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 32, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_3:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 48, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 48, align 8, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 0, 0, implicit $exec :: (load (<4 x i32>) from %ir.src, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_1:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 16, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 16, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_2:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 32, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 32, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_3:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 48, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 48, addrspace 1) ; GFX-950-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:vreg_512_align2 = REG_SEQUENCE [[GLOBAL_LOAD_DWORDX4_]], %subreg.sub0_sub1_sub2_sub3, [[GLOBAL_LOAD_DWORDX4_1]], %subreg.sub4_sub5_sub6_sub7, [[GLOBAL_LOAD_DWORDX4_2]], %subreg.sub8_sub9_sub10_sub11, [[GLOBAL_LOAD_DWORDX4_3]], %subreg.sub12_sub13_sub14_sub15 - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_4:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 64, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 64, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_5:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 80, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 80, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_6:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 96, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 96, align 8, addrspace 1) - ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_7:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 112, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 112, align 8, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_4:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 64, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 64, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_5:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 80, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 80, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_6:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 96, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 96, addrspace 1) + ; GFX-950-NEXT: [[GLOBAL_LOAD_DWORDX4_7:%[0-9]+]]:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 [[REG_SEQUENCE]], 112, 0, implicit $exec :: (load (<4 x i32>) from %ir.src + 112, addrspace 1) ; GFX-950-NEXT: [[REG_SEQUENCE3:%[0-9]+]]:vreg_512_align2 = REG_SEQUENCE [[GLOBAL_LOAD_DWORDX4_4]], %subreg.sub0_sub1_sub2_sub3, [[GLOBAL_LOAD_DWORDX4_5]], %subreg.sub4_sub5_sub6_sub7, [[GLOBAL_LOAD_DWORDX4_6]], %subreg.sub8_sub9_sub10_sub11, [[GLOBAL_LOAD_DWORDX4_7]], %subreg.sub12_sub13_sub14_sub15 ; GFX-950-NEXT: [[REG_SEQUENCE4:%[0-9]+]]:vreg_1024_align2 = REG_SEQUENCE [[REG_SEQUENCE2]], %subreg.sub0_sub1_sub2_sub3_sub4_sub5_sub6_sub7_sub8_sub9_sub10_sub11_sub12_sub13_sub14_sub15, [[REG_SEQUENCE3]], %subreg.sub16_sub17_sub18_sub19_sub20_sub21_sub22_sub23_sub24_sub25_sub26_sub27_sub28_sub29_sub30_sub31 ; GFX-950-NEXT: [[COPY4:%[0-9]+]]:vreg_512_align2 = COPY [[REG_SEQUENCE4]].sub0_sub1_sub2_sub3_sub4_sub5_sub6_sub7_sub8_sub9_sub10_sub11_sub12_sub13_sub14_sub15, debug-location !4 diff --git a/llvm/test/CodeGen/AMDGPU/add_i128.ll b/llvm/test/CodeGen/AMDGPU/add_i128.ll index 5342ec89b0d69..094d7610fb529 100644 --- a/llvm/test/CodeGen/AMDGPU/add_i128.ll +++ b/llvm/test/CodeGen/AMDGPU/add_i128.ll @@ -90,7 +90,7 @@ define amdgpu_kernel void @sgpr_operand_reversed(ptr addrspace(1) noalias %out, define amdgpu_kernel void @test_sreg(ptr addrspace(1) noalias %out, i128 %a, i128 %b) { ; GCN-LABEL: test_sreg: ; GCN: ; %bb.0: -; GCN-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0xb +; GCN-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0xd ; GCN-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9 ; GCN-NEXT: s_mov_b32 s3, 0xf000 ; GCN-NEXT: s_mov_b32 s2, -1 diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll index 02eef5ab5e98b..c353a96fb3bfb 100644 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll +++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll @@ -764,7 +764,7 @@ define amdgpu_kernel void @kernel_uses_read_register_a56_59(ptr addrspace(1) %pt ; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_read_register_a56_59( ; CHECK-SAME: ptr addrspace(1) [[PTR:%.*]]) #[[ATTR20:[0-9]+]] { ; CHECK-NEXT: [[REG:%.*]] = call i128 @llvm.read_register.i128(metadata [[META3:![0-9]+]]) -; CHECK-NEXT: store i128 [[REG]], ptr addrspace(1) [[PTR]], align 8 +; CHECK-NEXT: store i128 [[REG]], ptr addrspace(1) [[PTR]], align 16 ; CHECK-NEXT: call void @use_most() ; CHECK-NEXT: ret void ; diff --git a/llvm/test/CodeGen/AMDGPU/ctpop64.ll b/llvm/test/CodeGen/AMDGPU/ctpop64.ll index edf9ef109a42d..2fa36b29506d0 100644 --- a/llvm/test/CodeGen/AMDGPU/ctpop64.ll +++ b/llvm/test/CodeGen/AMDGPU/ctpop64.ll @@ -572,7 +572,7 @@ endif: define amdgpu_kernel void @s_ctpop_i128(ptr addrspace(1) noalias %out, i128 %val) nounwind { ; SI-LABEL: s_ctpop_i128: ; SI: ; %bb.0: -; SI-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb +; SI-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xd ; SI-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9 ; SI-NEXT: s_mov_b32 s7, 0xf000 ; SI-NEXT: s_mov_b32 s6, -1 @@ -586,7 +586,7 @@ define amdgpu_kernel void @s_ctpop_i128(ptr addrspace(1) noalias %out, i128 %val ; ; VI-LABEL: s_ctpop_i128: ; VI: ; %bb.0: -; VI-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x2c +; VI-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34 ; VI-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x24 ; VI-NEXT: s_mov_b32 s7, 0xf000 ; VI-NEXT: s_mov_b32 s6, -1 @@ -601,7 +601,7 @@ define amdgpu_kernel void @s_ctpop_i128(ptr addrspace(1) noalias %out, i128 %val ; GFX12-LABEL: s_ctpop_i128: ; GFX12: ; %bb.0: ; GFX12-NEXT: s_clause 0x1 -; GFX12-NEXT: s_load_b128 s[0:3], s[4:5], 0x2c +; GFX12-NEXT: s_load_b128 s[0:3], s[4:5], 0x34 ; GFX12-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 ; GFX12-NEXT: v_mov_b32_e32 v1, 0 ; GFX12-NEXT: s_wait_kmcnt 0x0 @@ -621,52 +621,51 @@ define amdgpu_kernel void @s_ctpop_i128(ptr addrspace(1) noalias %out, i128 %val define amdgpu_kernel void @s_ctpop_i65(ptr addrspace(1) noalias %out, i65 %val) nounwind { ; SI-LABEL: s_ctpop_i65: ; SI: ; %bb.0: -; SI-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9 -; SI-NEXT: s_load_dword s8, s[4:5], 0xd -; SI-NEXT: s_mov_b32 s7, 0xf000 -; SI-NEXT: s_mov_b32 s6, -1 +; SI-NEXT: s_load_dword s6, s[4:5], 0xf +; SI-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9 +; SI-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0xd +; SI-NEXT: s_mov_b32 s3, 0xf000 +; SI-NEXT: s_mov_b32 s2, -1 ; SI-NEXT: s_waitcnt lgkmcnt(0) -; SI-NEXT: s_mov_b32 s4, s0 -; SI-NEXT: s_and_b32 s0, s8, 0xff -; SI-NEXT: s_mov_b32 s5, s1 -; SI-NEXT: s_bcnt1_i32_b32 s0, s0 -; SI-NEXT: s_bcnt1_i32_b64 s1, s[2:3] -; SI-NEXT: s_add_i32 s0, s1, s0 -; SI-NEXT: v_mov_b32_e32 v0, s0 -; SI-NEXT: buffer_store_dword v0, off, s[4:7], 0 +; SI-NEXT: s_and_b32 s6, s6, 0xff +; SI-NEXT: s_bcnt1_i32_b32 s6, s6 +; SI-NEXT: s_bcnt1_i32_b64 s4, s[4:5] +; SI-NEXT: s_add_i32 s4, s4, s6 +; SI-NEXT: v_mov_b32_e32 v0, s4 +; SI-NEXT: buffer_store_dword v0, off, s[0:3], 0 ; SI-NEXT: s_endpgm ; ; VI-LABEL: s_ctpop_i65: ; VI: ; %bb.0: -; VI-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24 -; VI-NEXT: s_load_dword s8, s[4:5], 0x34 -; VI-NEXT: s_mov_b32 s7, 0xf000 -; VI-NEXT: s_mov_b32 s6, -1 +; VI-NEXT: s_load_dword s6, s[4:5], 0x3c +; VI-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24 +; VI-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x34 +; VI-NEXT: s_mov_b32 s3, 0xf000 +; VI-NEXT: s_mov_b32 s2, -1 ; VI-NEXT: s_waitcnt lgkmcnt(0) -; VI-NEXT: s_mov_b32 s4, s0 -; VI-NEXT: s_and_b32 s0, s8, 0xff -; VI-NEXT: s_mov_b32 s5, s1 -; VI-NEXT: s_bcnt1_i32_b32 s0, s0 -; VI-NEXT: s_bcnt1_i32_b64 s1, s[2:3] -; VI-NEXT: s_add_i32 s0, s1, s0 -; VI-NEXT: v_mov_b32_e32 v0, s0 -; VI-NEXT: buffer_store_dword v0, off, s[4:7], 0 +; VI-NEXT: s_and_b32 s6, s6, 0xff +; VI-NEXT: s_bcnt1_i32_b32 s6, s6 +; VI-NEXT: s_bcnt1_i32_b64 s4, s[4:5] +; VI-NEXT: s_add_i32 s4, s4, s6 +; VI-NEXT: v_mov_b32_e32 v0, s4 +; VI-NEXT: buffer_store_dword v0, off, s[0:3], 0 ; VI-NEXT: s_endpgm ; ; GFX12-LABEL: s_ctpop_i65: ; GFX12: ; %bb.0: -; GFX12-NEXT: s_clause 0x1 -; GFX12-NEXT: s_load_u8 s6, s[4:5], 0x34 -; GFX12-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 +; GFX12-NEXT: s_clause 0x2 +; GFX12-NEXT: s_load_u8 s0, s[4:5], 0x3c +; GFX12-NEXT: s_load_b64 s[2:3], s[4:5], 0x34 +; GFX12-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 ; GFX12-NEXT: v_mov_b32_e32 v1, 0 ; GFX12-NEXT: s_wait_kmcnt 0x0 -; GFX12-NEXT: s_and_b64 s[4:5], s[6:7], 1 +; GFX12-NEXT: s_and_b64 s[0:1], s[0:1], 1 ; GFX12-NEXT: s_bcnt1_i32_b64 s2, s[2:3] -; GFX12-NEXT: s_bcnt1_i32_b64 s3, s[4:5] +; GFX12-NEXT: s_bcnt1_i32_b64 s0, s[0:1] ; GFX12-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1) -; GFX12-NEXT: s_add_co_i32 s2, s3, s2 -; GFX12-NEXT: v_mov_b32_e32 v0, s2 -; GFX12-NEXT: global_store_b32 v1, v0, s[0:1] +; GFX12-NEXT: s_add_co_i32 s0, s0, s2 +; GFX12-NEXT: v_mov_b32_e32 v0, s0 +; GFX12-NEXT: global_store_b32 v1, v0, s[4:5] ; GFX12-NEXT: s_endpgm %ctpop = call i65 @llvm.ctpop.i65(i65 %val) nounwind readnone %truncctpop = trunc i65 %ctpop to i32 diff --git a/llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll b/llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll new file mode 100644 index 0000000000000..aa7d33fb4fd77 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll @@ -0,0 +1,82 @@ +; RUN: llc -mtriple=amdgpu11.00-amd-amdhsa < %s | FileCheck %s + +; The AMDGPU data layout has to give i128 an ABI alignment of 16, matching the +; ABI implemented by Clang, whose AMDGPUTargetInfo leaves Int128Align at its +; 128-bit default. Without an explicit entry the alignment would be inherited +; from the i64:64 entry, and the layout LLVM computes for an aggregate would +; then disagree with the one a frontend used when it emitted the field offsets. +; +; For kernel arguments that disagreement is an ABI break rather than a missed +; optimization: the kernarg slot is sized from the data layout, so the tail of +; the argument is never copied into the kernarg segment and loads of it run off +; the end of the segment. + +; struct S { i64 a; i128 b; }: ABI align 16, size 32, offsetof(b) == 16. +; The byref slot must be 32 bytes, not 24, and `b` must be loaded from 0x20 +; (kernarg base 16 plus a field offset of 16) which stays inside the segment. +; CHECK-LABEL: {{^}}kernarg_i128: +; CHECK: s_load_b128 s[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], 0x20 + +; An i128 following a smaller member is padded out to offset 16 rather than +; packed at offset 8. +; CHECK-LABEL: {{^}}kernarg_i128_after_i8: +; CHECK: s_load_b128 s[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], 0x20 + +; A bare i128 kernel argument is 16-byte aligned in the kernarg segment, so it +; starts at 16 (not 8) and the argument after it at 32 (not 24). +; CHECK-LABEL: {{^}}kernarg_i128_scalar: +; CHECK: s_load_b128 s[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], 0x10 + +; The kernel metadata is emitted once, after every function, so the per-kernel +; argument offsets are checked here in order rather than under each label. +; +; .kernarg_segment_size is deliberately not checked: it also covers the 256 +; bytes of hidden implicit arguments, which is noise for what this test pins +; down. The per-argument .offset/.size entries and .kernarg_segment_align are +; the layout facts that matter. +; CHECK: .amdgpu_metadata + +; CHECK: .name: s +; CHECK-NEXT: .offset: 16 +; CHECK-NEXT: .size: 32 +; CHECK: .kernarg_segment_align: 16 +; CHECK: .name: kernarg_i128 +; +; CHECK: .name: s +; CHECK-NEXT: .offset: 16 +; CHECK-NEXT: .size: 32 +; CHECK: .name: kernarg_i128_after_i8 +; +; CHECK: .name: a +; CHECK-NEXT: .offset: 16 +; CHECK-NEXT: .size: 16 +; CHECK: .name: b +; CHECK-NEXT: .offset: 32 +; CHECK-NEXT: .size: 8 +; CHECK: .name: kernarg_i128_scalar + +define amdgpu_kernel void @kernarg_i128(ptr addrspace(1) %out, + ptr addrspace(4) byref({ i64, i128 }) align 16 %s) { + %pb = getelementptr inbounds i8, ptr addrspace(4) %s, i64 16 + %b = load i128, ptr addrspace(4) %pb, align 16 + store i128 %b, ptr addrspace(1) %out, align 16 + ret void +} + +define amdgpu_kernel void @kernarg_i128_after_i8(ptr addrspace(1) %out, + ptr addrspace(4) byref({ i8, i128 }) align 16 %s) { + %pb = getelementptr inbounds i8, ptr addrspace(4) %s, i64 16 + %b = load i128, ptr addrspace(4) %pb, align 16 + store i128 %b, ptr addrspace(1) %out, align 16 + ret void +} + +define amdgpu_kernel void @kernarg_i128_scalar(ptr addrspace(1) %out, i128 %a, i64 %b) { + %ext = zext i64 %b to i128 + %sum = add i128 %a, %ext + store i128 %sum, ptr addrspace(1) %out, align 16 + ret void +} + +!llvm.module.flags = !{!0} +!0 = !{i32 1, !"amdhsa_code_object_version", i32 500} diff --git a/llvm/test/CodeGen/AMDGPU/kernel-args.ll b/llvm/test/CodeGen/AMDGPU/kernel-args.ll index 5b95ef59499ac..9bd410d40da08 100644 --- a/llvm/test/CodeGen/AMDGPU/kernel-args.ll +++ b/llvm/test/CodeGen/AMDGPU/kernel-args.ll @@ -4459,53 +4459,54 @@ entry: define amdgpu_kernel void @i65_arg(ptr addrspace(1) nocapture %out, i65 %in) nounwind { ; SI-LABEL: i65_arg: ; SI: ; %bb.0: ; %entry -; SI-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9 -; SI-NEXT: s_load_dword s8, s[4:5], 0xd -; SI-NEXT: s_mov_b32 s7, 0xf000 -; SI-NEXT: s_mov_b32 s6, -1 +; SI-NEXT: s_load_dword s6, s[4:5], 0xf +; SI-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9 +; SI-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0xd +; SI-NEXT: s_mov_b32 s3, 0xf000 +; SI-NEXT: s_mov_b32 s2, -1 ; SI-NEXT: s_waitcnt lgkmcnt(0) -; SI-NEXT: s_mov_b32 s4, s0 -; SI-NEXT: s_and_b32 s0, s8, 1 -; SI-NEXT: s_mov_b32 s5, s1 -; SI-NEXT: v_mov_b32_e32 v0, s0 -; SI-NEXT: buffer_store_byte v0, off, s[4:7], 0 offset:8 +; SI-NEXT: s_and_b32 s6, s6, 1 +; SI-NEXT: v_mov_b32_e32 v0, s6 +; SI-NEXT: buffer_store_byte v0, off, s[0:3], 0 offset:8 ; SI-NEXT: s_waitcnt expcnt(0) -; SI-NEXT: v_mov_b32_e32 v0, s2 -; SI-NEXT: v_mov_b32_e32 v1, s3 -; SI-NEXT: buffer_store_dwordx2 v[0:1], off, s[4:7], 0 +; SI-NEXT: v_mov_b32_e32 v0, s4 +; SI-NEXT: v_mov_b32_e32 v1, s5 +; SI-NEXT: buffer_store_dwordx2 v[0:1], off, s[0:3], 0 ; SI-NEXT: s_endpgm ; ; VI-LABEL: i65_arg: ; VI: ; %bb.0: ; %entry -; VI-NEXT: s_load_dword s6, s[4:5], 0x34 -; VI-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24 +; VI-NEXT: s_load_dword s6, s[4:5], 0x3c +; VI-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24 +; VI-NEXT: s_load_dwordx2 s[2:3], s[4:5], 0x34 ; VI-NEXT: s_waitcnt lgkmcnt(0) ; VI-NEXT: s_and_b32 s4, s6, 1 ; VI-NEXT: v_mov_b32_e32 v0, s0 ; VI-NEXT: v_mov_b32_e32 v1, s1 ; VI-NEXT: s_add_u32 s0, s0, 8 ; VI-NEXT: s_addc_u32 s1, s1, 0 -; VI-NEXT: v_mov_b32_e32 v6, s4 -; VI-NEXT: v_mov_b32_e32 v5, s1 -; VI-NEXT: v_mov_b32_e32 v4, s0 +; VI-NEXT: v_mov_b32_e32 v4, s4 +; VI-NEXT: v_mov_b32_e32 v3, s1 +; VI-NEXT: v_mov_b32_e32 v2, s0 +; VI-NEXT: flat_store_byte v[2:3], v4 ; VI-NEXT: v_mov_b32_e32 v2, s2 ; VI-NEXT: v_mov_b32_e32 v3, s3 -; VI-NEXT: flat_store_byte v[4:5], v6 ; VI-NEXT: flat_store_dwordx2 v[0:1], v[2:3] ; VI-NEXT: s_endpgm ; ; GFX9-LABEL: i65_arg: ; GFX9: ; %bb.0: ; %entry -; GFX9-NEXT: s_load_dword s4, s[8:9], 0x10 -; GFX9-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0 +; GFX9-NEXT: s_load_dword s4, s[8:9], 0x18 +; GFX9-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x10 +; GFX9-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x0 ; GFX9-NEXT: v_mov_b32_e32 v2, 0 ; GFX9-NEXT: s_waitcnt lgkmcnt(0) ; GFX9-NEXT: s_and_b32 s4, s4, 1 ; GFX9-NEXT: v_mov_b32_e32 v3, s4 -; GFX9-NEXT: v_mov_b32_e32 v0, s2 -; GFX9-NEXT: v_mov_b32_e32 v1, s3 -; GFX9-NEXT: global_store_byte v2, v3, s[0:1] offset:8 -; GFX9-NEXT: global_store_dwordx2 v2, v[0:1], s[0:1] +; GFX9-NEXT: v_mov_b32_e32 v0, s0 +; GFX9-NEXT: v_mov_b32_e32 v1, s1 +; GFX9-NEXT: global_store_byte v2, v3, s[2:3] offset:8 +; GFX9-NEXT: global_store_dwordx2 v2, v[0:1], s[2:3] ; GFX9-NEXT: s_endpgm ; ; EG-LABEL: i65_arg: diff --git a/llvm/test/CodeGen/AMDGPU/kernel-argument-dag-lowering.ll b/llvm/test/CodeGen/AMDGPU/kernel-argument-dag-lowering.ll index e9f4ccda47905..6e12a24927b22 100644 --- a/llvm/test/CodeGen/AMDGPU/kernel-argument-dag-lowering.ll +++ b/llvm/test/CodeGen/AMDGPU/kernel-argument-dag-lowering.ll @@ -180,22 +180,23 @@ define amdgpu_kernel void @v6i32_arg(<6 x i32> %in) nounwind { define amdgpu_kernel void @i65_arg(ptr addrspace(1) nocapture %out, i65 %in) #0 { ; GCN-LABEL: i65_arg: ; GCN: ; %bb.0: ; %entry -; GCN-NEXT: s_load_dword s4, s[8:9], 0x10 -; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0 +; GCN-NEXT: s_load_dword s4, s[8:9], 0x18 +; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x10 +; GCN-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x0 ; GCN-NEXT: v_mov_b32_e32 v2, 0 ; GCN-NEXT: s_waitcnt lgkmcnt(0) ; GCN-NEXT: s_and_b32 s4, s4, 1 ; GCN-NEXT: v_mov_b32_e32 v3, s4 -; GCN-NEXT: v_mov_b32_e32 v0, s2 -; GCN-NEXT: v_mov_b32_e32 v1, s3 -; GCN-NEXT: global_store_byte v2, v3, s[0:1] offset:8 -; GCN-NEXT: global_store_dwordx2 v2, v[0:1], s[0:1] +; GCN-NEXT: v_mov_b32_e32 v0, s0 +; GCN-NEXT: v_mov_b32_e32 v1, s1 +; GCN-NEXT: global_store_byte v2, v3, s[2:3] offset:8 +; GCN-NEXT: global_store_dwordx2 v2, v[0:1], s[2:3] ; GCN-NEXT: s_endpgm entry: store i65 %in, ptr addrspace(1) %out, align 4 ret void } -; GCN: .amdhsa_kernarg_size 24 +; GCN: .amdhsa_kernarg_size 32 define amdgpu_kernel void @empty_struct_arg({} %in) #0 { ; GCN-LABEL: empty_struct_arg: diff --git a/llvm/test/CodeGen/AMDGPU/lower-buffer-fat-pointers-constants.ll b/llvm/test/CodeGen/AMDGPU/lower-buffer-fat-pointers-constants.ll index e471178394303..970535c96745f 100644 --- a/llvm/test/CodeGen/AMDGPU/lower-buffer-fat-pointers-constants.ll +++ b/llvm/test/CodeGen/AMDGPU/lower-buffer-fat-pointers-constants.ll @@ -117,7 +117,7 @@ define ptr @gep_of_p7_struct() { define ptr addrspace(7) @gep_p7_from_p7() { ; CHECK-LABEL: define { ptr addrspace(8), i32 } @gep_p7_from_p7() { -; CHECK-NEXT: ret { ptr addrspace(8), i32 } { ptr addrspace(8) @buf, i32 48 } +; CHECK-NEXT: ret { ptr addrspace(8), i32 } { ptr addrspace(8) @buf, i32 64 } ; ret ptr addrspace(7) getelementptr (ptr addrspace(7), ptr addrspace(7) addrspacecast (ptr addrspace(8) @buf to ptr addrspace(7)), diff --git a/llvm/test/CodeGen/AMDGPU/lower-buffer-fat-pointers-contents-legalization.ll b/llvm/test/CodeGen/AMDGPU/lower-buffer-fat-pointers-contents-legalization.ll index 5efed9e569f2f..91b0c643c6e23 100644 --- a/llvm/test/CodeGen/AMDGPU/lower-buffer-fat-pointers-contents-legalization.ll +++ b/llvm/test/CodeGen/AMDGPU/lower-buffer-fat-pointers-contents-legalization.ll @@ -100,7 +100,7 @@ define void @store_i64(i64 %data, ptr addrspace(8) inreg %buf) { define i128 @load_i128(ptr addrspace(8) inreg %buf) { ; CHECK-LABEL: define i128 @load_i128( ; CHECK-SAME: ptr addrspace(8) inreg [[BUF:%.*]]) { -; CHECK-NEXT: [[RET:%.*]] = call i128 @llvm.amdgcn.raw.ptr.buffer.load.i128(ptr addrspace(8) align 8 [[BUF]], i32 0, i32 0, i32 0) +; CHECK-NEXT: [[RET:%.*]] = call i128 @llvm.amdgcn.raw.ptr.buffer.load.i128(ptr addrspace(8) align 16 [[BUF]], i32 0, i32 0, i32 0) ; CHECK-NEXT: ret i128 [[RET]] ; %p = addrspacecast ptr addrspace(8) %buf to ptr addrspace(7) @@ -111,7 +111,7 @@ define i128 @load_i128(ptr addrspace(8) inreg %buf) { define void @store_i128(i128 %data, ptr addrspace(8) inreg %buf) { ; CHECK-LABEL: define void @store_i128( ; CHECK-SAME: i128 [[DATA:%.*]], ptr addrspace(8) inreg [[BUF:%.*]]) { -; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.i128(i128 [[DATA]], ptr addrspace(8) align 8 [[BUF]], i32 0, i32 0, i32 0) +; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.i128(i128 [[DATA]], ptr addrspace(8) align 16 [[BUF]], i32 0, i32 0, i32 0) ; CHECK-NEXT: ret void ; %p = addrspacecast ptr addrspace(8) %buf to ptr addrspace(7) @@ -1720,7 +1720,7 @@ define void @store_i40(i40 %data, ptr addrspace(8) inreg %buf) { define i96 @load_i96(ptr addrspace(8) inreg %buf) { ; CHECK-LABEL: define i96 @load_i96( ; CHECK-SAME: ptr addrspace(8) inreg [[BUF:%.*]]) { -; CHECK-NEXT: [[RET_LOADABLE:%.*]] = call <3 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v3i32(ptr addrspace(8) align 8 [[BUF]], i32 0, i32 0, i32 0) +; CHECK-NEXT: [[RET_LOADABLE:%.*]] = call <3 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v3i32(ptr addrspace(8) align 16 [[BUF]], i32 0, i32 0, i32 0) ; CHECK-NEXT: [[RET:%.*]] = bitcast <3 x i32> [[RET_LOADABLE]] to i96 ; CHECK-NEXT: ret i96 [[RET]] ; @@ -1733,7 +1733,7 @@ define void @store_i96(i96 %data, ptr addrspace(8) inreg %buf) { ; CHECK-LABEL: define void @store_i96( ; CHECK-SAME: i96 [[DATA:%.*]], ptr addrspace(8) inreg [[BUF:%.*]]) { ; CHECK-NEXT: [[DATA_LEGAL:%.*]] = bitcast i96 [[DATA]] to <3 x i32> -; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.v3i32(<3 x i32> [[DATA_LEGAL]], ptr addrspace(8) align 8 [[BUF]], i32 0, i32 0, i32 0) +; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.v3i32(<3 x i32> [[DATA_LEGAL]], ptr addrspace(8) align 16 [[BUF]], i32 0, i32 0, i32 0) ; CHECK-NEXT: ret void ; %p = addrspacecast ptr addrspace(8) %buf to ptr addrspace(7) @@ -1744,10 +1744,10 @@ define void @store_i96(i96 %data, ptr addrspace(8) inreg %buf) { define i160 @load_i160(ptr addrspace(8) inreg %buf) { ; CHECK-LABEL: define i160 @load_i160( ; CHECK-SAME: ptr addrspace(8) inreg [[BUF:%.*]]) { -; CHECK-NEXT: [[RET_OFF_0:%.*]] = call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) align 8 [[BUF]], i32 0, i32 0, i32 0) +; CHECK-NEXT: [[RET_OFF_0:%.*]] = call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) align 16 [[BUF]], i32 0, i32 0, i32 0) ; CHECK-NEXT: [[RET_EXT_0:%.*]] = shufflevector <4 x i32> [[RET_OFF_0]], <4 x i32> poison, <5 x i32> ; CHECK-NEXT: [[RET_PARTS_0:%.*]] = shufflevector <5 x i32> poison, <5 x i32> [[RET_EXT_0]], <5 x i32> -; CHECK-NEXT: [[RET_OFF_16:%.*]] = call i32 @llvm.amdgcn.raw.ptr.buffer.load.i32(ptr addrspace(8) align 8 [[BUF]], i32 16, i32 0, i32 0) +; CHECK-NEXT: [[RET_OFF_16:%.*]] = call i32 @llvm.amdgcn.raw.ptr.buffer.load.i32(ptr addrspace(8) align 16 [[BUF]], i32 16, i32 0, i32 0) ; CHECK-NEXT: [[RET_SLICE_4:%.*]] = insertelement <5 x i32> [[RET_PARTS_0]], i32 [[RET_OFF_16]], i64 4 ; CHECK-NEXT: [[RET:%.*]] = bitcast <5 x i32> [[RET_SLICE_4]] to i160 ; CHECK-NEXT: ret i160 [[RET]] @@ -1762,9 +1762,9 @@ define void @store_i160(i160 %data, ptr addrspace(8) inreg %buf) { ; CHECK-SAME: i160 [[DATA:%.*]], ptr addrspace(8) inreg [[BUF:%.*]]) { ; CHECK-NEXT: [[DATA_LEGAL:%.*]] = bitcast i160 [[DATA]] to <5 x i32> ; CHECK-NEXT: [[DATA_SLICE_0:%.*]] = shufflevector <5 x i32> [[DATA_LEGAL]], <5 x i32> poison, <4 x i32> -; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[DATA_SLICE_0]], ptr addrspace(8) align 8 [[BUF]], i32 0, i32 0, i32 0) +; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[DATA_SLICE_0]], ptr addrspace(8) align 16 [[BUF]], i32 0, i32 0, i32 0) ; CHECK-NEXT: [[DATA_SLICE_4:%.*]] = extractelement <5 x i32> [[DATA_LEGAL]], i64 4 -; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[DATA_SLICE_4]], ptr addrspace(8) align 8 [[BUF]], i32 16, i32 0, i32 0) +; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[DATA_SLICE_4]], ptr addrspace(8) align 16 [[BUF]], i32 16, i32 0, i32 0) ; CHECK-NEXT: ret void ; %p = addrspacecast ptr addrspace(8) %buf to ptr addrspace(7) @@ -1775,10 +1775,10 @@ define void @store_i160(i160 %data, ptr addrspace(8) inreg %buf) { define i256 @load_i256(ptr addrspace(8) inreg %buf) { ; CHECK-LABEL: define i256 @load_i256( ; CHECK-SAME: ptr addrspace(8) inreg [[BUF:%.*]]) { -; CHECK-NEXT: [[RET_OFF_0:%.*]] = call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) align 8 [[BUF]], i32 0, i32 0, i32 0) +; CHECK-NEXT: [[RET_OFF_0:%.*]] = call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) align 16 [[BUF]], i32 0, i32 0, i32 0) ; CHECK-NEXT: [[RET_EXT_0:%.*]] = shufflevector <4 x i32> [[RET_OFF_0]], <4 x i32> poison, <8 x i32> ; CHECK-NEXT: [[RET_PARTS_0:%.*]] = shufflevector <8 x i32> poison, <8 x i32> [[RET_EXT_0]], <8 x i32> -; CHECK-NEXT: [[RET_OFF_16:%.*]] = call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) align 8 [[BUF]], i32 16, i32 0, i32 0) +; CHECK-NEXT: [[RET_OFF_16:%.*]] = call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) align 16 [[BUF]], i32 16, i32 0, i32 0) ; CHECK-NEXT: [[RET_EXT_4:%.*]] = shufflevector <4 x i32> [[RET_OFF_16]], <4 x i32> poison, <8 x i32> ; CHECK-NEXT: [[RET_PARTS_4:%.*]] = shufflevector <8 x i32> [[RET_PARTS_0]], <8 x i32> [[RET_EXT_4]], <8 x i32> ; CHECK-NEXT: [[RET:%.*]] = bitcast <8 x i32> [[RET_PARTS_4]] to i256 @@ -1794,9 +1794,9 @@ define void @store_i256(i256 %data, ptr addrspace(8) inreg %buf) { ; CHECK-SAME: i256 [[DATA:%.*]], ptr addrspace(8) inreg [[BUF:%.*]]) { ; CHECK-NEXT: [[DATA_LEGAL:%.*]] = bitcast i256 [[DATA]] to <8 x i32> ; CHECK-NEXT: [[DATA_SLICE_0:%.*]] = shufflevector <8 x i32> [[DATA_LEGAL]], <8 x i32> poison, <4 x i32> -; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[DATA_SLICE_0]], ptr addrspace(8) align 8 [[BUF]], i32 0, i32 0, i32 0) +; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[DATA_SLICE_0]], ptr addrspace(8) align 16 [[BUF]], i32 0, i32 0, i32 0) ; CHECK-NEXT: [[DATA_SLICE_4:%.*]] = shufflevector <8 x i32> [[DATA_LEGAL]], <8 x i32> poison, <4 x i32> -; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[DATA_SLICE_4]], ptr addrspace(8) align 8 [[BUF]], i32 16, i32 0, i32 0) +; CHECK-NEXT: call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[DATA_SLICE_4]], ptr addrspace(8) align 16 [[BUF]], i32 16, i32 0, i32 0) ; CHECK-NEXT: ret void ; %p = addrspacecast ptr addrspace(8) %buf to ptr addrspace(7) diff --git a/llvm/test/CodeGen/AMDGPU/mul.ll b/llvm/test/CodeGen/AMDGPU/mul.ll index d26a899fa57e5..8e4c092a9ae60 100644 --- a/llvm/test/CodeGen/AMDGPU/mul.ll +++ b/llvm/test/CodeGen/AMDGPU/mul.ll @@ -3393,8 +3393,8 @@ endif: define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a, [8 x i32], i128 %b) nounwind #0 { ; SI-LABEL: s_mul_i128: ; SI: ; %bb.0: ; %entry -; SI-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x13 -; SI-NEXT: s_load_dwordx4 s[12:15], s[4:5], 0x1f +; SI-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x15 +; SI-NEXT: s_load_dwordx4 s[12:15], s[4:5], 0x21 ; SI-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9 ; SI-NEXT: s_mov_b32 s3, 0xf000 ; SI-NEXT: s_mov_b32 s2, -1 @@ -3442,8 +3442,8 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a, ; ; VI-LABEL: s_mul_i128: ; VI: ; %bb.0: ; %entry -; VI-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x4c -; VI-NEXT: s_load_dwordx4 s[12:15], s[4:5], 0x7c +; VI-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x54 +; VI-NEXT: s_load_dwordx4 s[12:15], s[4:5], 0x84 ; VI-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24 ; VI-NEXT: s_mov_b32 s3, 0xf000 ; VI-NEXT: s_mov_b32 s2, -1 @@ -3477,8 +3477,8 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a, ; ; GFX9-LABEL: s_mul_i128: ; GFX9: ; %bb.0: ; %entry -; GFX9-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x7c -; GFX9-NEXT: s_load_dwordx4 s[12:15], s[4:5], 0x4c +; GFX9-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x84 +; GFX9-NEXT: s_load_dwordx4 s[12:15], s[4:5], 0x54 ; GFX9-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24 ; GFX9-NEXT: s_mov_b32 s3, 0xf000 ; GFX9-NEXT: s_mov_b32 s2, -1 @@ -3528,8 +3528,8 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a, ; GFX10-LABEL: s_mul_i128: ; GFX10: ; %bb.0: ; %entry ; GFX10-NEXT: s_clause 0x2 -; GFX10-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x4c -; GFX10-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x7c +; GFX10-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x54 +; GFX10-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x84 ; GFX10-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x24 ; GFX10-NEXT: s_mov_b32 s6, 0 ; GFX10-NEXT: s_mov_b32 s5, s6 @@ -3579,8 +3579,8 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a, ; GFX11-LABEL: s_mul_i128: ; GFX11: ; %bb.0: ; %entry ; GFX11-NEXT: s_clause 0x2 -; GFX11-NEXT: s_load_b128 s[0:3], s[4:5], 0x4c -; GFX11-NEXT: s_load_b128 s[8:11], s[4:5], 0x7c +; GFX11-NEXT: s_load_b128 s[0:3], s[4:5], 0x54 +; GFX11-NEXT: s_load_b128 s[8:11], s[4:5], 0x84 ; GFX11-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 ; GFX11-NEXT: s_mov_b32 s6, 0 ; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) @@ -3630,8 +3630,8 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a, ; GFX12-LABEL: s_mul_i128: ; GFX12: ; %bb.0: ; %entry ; GFX12-NEXT: s_clause 0x1 -; GFX12-NEXT: s_load_b128 s[8:11], s[4:5], 0x7c -; GFX12-NEXT: s_load_b128 s[12:15], s[4:5], 0x4c +; GFX12-NEXT: s_load_b128 s[8:11], s[4:5], 0x84 +; GFX12-NEXT: s_load_b128 s[12:15], s[4:5], 0x54 ; GFX12-NEXT: s_mov_b32 s3, 0 ; GFX12-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 ; GFX12-NEXT: s_mov_b32 s7, s3 @@ -3676,8 +3676,8 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a, ; 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_clause 0x2 -; GFX1250-NEXT: s_load_b128 s[8:11], s[4:5], 0x7c nv -; GFX1250-NEXT: s_load_b128 s[12:15], s[4:5], 0x4c nv +; GFX1250-NEXT: s_load_b128 s[8:11], s[4:5], 0x84 nv +; GFX1250-NEXT: s_load_b128 s[12:15], s[4:5], 0x54 nv ; GFX1250-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv ; GFX1250-NEXT: s_wait_xcnt 0x0 ; GFX1250-NEXT: s_mov_b64 s[4:5], 0xffffffff @@ -3722,8 +3722,8 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a, ; GFX13-LABEL: s_mul_i128: ; GFX13: ; %bb.0: ; %entry ; GFX13-NEXT: s_clause 0x2 -; GFX13-NEXT: s_load_b128 s[8:11], s[4:5], 0x7c nv -; GFX13-NEXT: s_load_b128 s[12:15], s[4:5], 0x4c nv +; GFX13-NEXT: s_load_b128 s[8:11], s[4:5], 0x84 nv +; GFX13-NEXT: s_load_b128 s[12:15], s[4:5], 0x54 nv ; GFX13-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv ; GFX13-NEXT: s_mov_b64 s[4:5], 0xffffffff ; GFX13-NEXT: s_mov_b32 s3, 0 diff --git a/llvm/test/CodeGen/AMDGPU/opencl-printf.ll b/llvm/test/CodeGen/AMDGPU/opencl-printf.ll index 292fbe2331b95..ac0dce92df408 100644 --- a/llvm/test/CodeGen/AMDGPU/opencl-printf.ll +++ b/llvm/test/CodeGen/AMDGPU/opencl-printf.ll @@ -292,9 +292,9 @@ define amdgpu_kernel void @format_str_d(i1 %i1, i4 %i4, i8 %i8, i24 %i24, i16 %i ; GCN-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 4 ; GCN-NEXT: store i64 [[I64:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8 ; GCN-NEXT: [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8 -; GCN-NEXT: store i96 [[I96:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8 +; GCN-NEXT: store i96 [[I96:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 16 ; GCN-NEXT: [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 16 -; GCN-NEXT: store i128 [[I128:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8 +; GCN-NEXT: store i128 [[I128:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 16 ; GCN-NEXT: [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 16 ; GCN-NEXT: store i32 1234, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 4 ; GCN-NEXT: br label [[TMP7]] @@ -335,9 +335,9 @@ define amdgpu_kernel void @format_str_u(i1 %i1, i4 %i4, i8 %i8, i24 %i24, i16 %i ; GCN-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 4 ; GCN-NEXT: store i64 [[I64:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8 ; GCN-NEXT: [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8 -; GCN-NEXT: store i96 [[I96:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8 +; GCN-NEXT: store i96 [[I96:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 16 ; GCN-NEXT: [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 16 -; GCN-NEXT: store i128 [[I128:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8 +; GCN-NEXT: store i128 [[I128:%.*]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 16 ; GCN-NEXT: [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 16 ; GCN-NEXT: store i32 1234, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 4 ; GCN-NEXT: br label [[TMP7]] diff --git a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs-IR-lowering.ll b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs-IR-lowering.ll index 86b17eac16344..fd8b44f138d08 100644 --- a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs-IR-lowering.ll +++ b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs-IR-lowering.ll @@ -30,7 +30,7 @@ define amdgpu_kernel void @preload_block_count_x(ptr addrspace(1) %out) { define amdgpu_kernel void @no_free_sgprs_block_count_x(ptr addrspace(1) %out, i512) { ; NO-PRELOAD-LABEL: define amdgpu_kernel void @no_free_sgprs_block_count_x( ; NO-PRELOAD-SAME: ptr addrspace(1) [[OUT:%.*]], i512 [[TMP0:%.*]]) #[[ATTR0]] { -; NO-PRELOAD-NEXT: [[NO_FREE_SGPRS_BLOCK_COUNT_X_KERNARG_SEGMENT:%.*]] = call nonnull align 16 dereferenceable(328) ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr() +; NO-PRELOAD-NEXT: [[NO_FREE_SGPRS_BLOCK_COUNT_X_KERNARG_SEGMENT:%.*]] = call nonnull align 16 dereferenceable(336) ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr() ; NO-PRELOAD-NEXT: [[OUT_KERNARG_OFFSET:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[NO_FREE_SGPRS_BLOCK_COUNT_X_KERNARG_SEGMENT]], i64 0 ; NO-PRELOAD-NEXT: [[OUT_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT_KERNARG_OFFSET]], align 16, !invariant.load [[META0]] ; NO-PRELOAD-NEXT: [[IMP_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() @@ -40,7 +40,7 @@ define amdgpu_kernel void @no_free_sgprs_block_count_x(ptr addrspace(1) %out, i5 ; ; PRELOAD-LABEL: define amdgpu_kernel void @no_free_sgprs_block_count_x( ; PRELOAD-SAME: ptr addrspace(1) inreg [[OUT:%.*]], i512 [[TMP0:%.*]]) #[[ATTR0]] { -; PRELOAD-NEXT: [[NO_FREE_SGPRS_BLOCK_COUNT_X_KERNARG_SEGMENT:%.*]] = call nonnull align 16 dereferenceable(328) ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr() +; PRELOAD-NEXT: [[NO_FREE_SGPRS_BLOCK_COUNT_X_KERNARG_SEGMENT:%.*]] = call nonnull align 16 dereferenceable(336) ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr() ; PRELOAD-NEXT: [[IMP_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() ; PRELOAD-NEXT: [[LOAD:%.*]] = load i32, ptr addrspace(4) [[IMP_ARG_PTR]], align 4 ; PRELOAD-NEXT: store i32 [[LOAD]], ptr addrspace(1) [[OUT]], align 4 diff --git a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll index db7f998270f16..e02a4430bb8b2 100644 --- a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll +++ b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll @@ -100,7 +100,7 @@ define amdgpu_kernel void @no_free_sgprs_block_count_x(ptr addrspace(1) inreg %o ; GFX942-NEXT: .p2align 8 ; GFX942-NEXT: ; %bb.2: ; GFX942-NEXT: .LBB2_0: -; GFX942-NEXT: s_load_dword s0, s[4:5], 0x28 +; GFX942-NEXT: s_load_dword s0, s[4:5], 0x30 ; GFX942-NEXT: v_mov_b32_e32 v0, 0 ; GFX942-NEXT: s_waitcnt lgkmcnt(0) ; GFX942-NEXT: v_mov_b32_e32 v1, s0 @@ -115,7 +115,7 @@ define amdgpu_kernel void @no_free_sgprs_block_count_x(ptr addrspace(1) inreg %o ; GFX90a-NEXT: .p2align 8 ; GFX90a-NEXT: ; %bb.2: ; GFX90a-NEXT: .LBB2_0: -; GFX90a-NEXT: s_load_dword s0, s[8:9], 0x28 +; GFX90a-NEXT: s_load_dword s0, s[8:9], 0x30 ; GFX90a-NEXT: v_mov_b32_e32 v0, 0 ; GFX90a-NEXT: s_waitcnt lgkmcnt(0) ; GFX90a-NEXT: v_mov_b32_e32 v1, s0 @@ -127,7 +127,7 @@ define amdgpu_kernel void @no_free_sgprs_block_count_x(ptr addrspace(1) inreg %o ; GFX1250-NEXT: global_prefetch_b8 v0, s[0:1] scope:SCOPE_SE ; 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: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s18 +; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s20 ; GFX1250-NEXT: global_store_b32 v0, v1, s[8:9] ; GFX1250-NEXT: s_endpgm %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() @@ -1048,14 +1048,17 @@ define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg % ; GFX942: ; %bb.1: ; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0 ; GFX942-NEXT: s_load_dwordx8 s[4:11], s[0:1], 0x8 -; GFX942-NEXT: s_load_dword s12, s[0:1], 0x28 +; GFX942-NEXT: s_load_dwordx2 s[12:13], s[0:1], 0x28 +; GFX942-NEXT: s_load_dword s14, s[0:1], 0x30 ; GFX942-NEXT: s_waitcnt lgkmcnt(0) ; GFX942-NEXT: s_branch .LBB21_0 ; GFX942-NEXT: .p2align 8 ; GFX942-NEXT: ; %bb.2: ; GFX942-NEXT: .LBB21_0: +; GFX942-NEXT: s_load_dword s0, s[0:1], 0x38 ; GFX942-NEXT: v_mov_b32_e32 v0, 0 -; GFX942-NEXT: v_mov_b32_e32 v1, s12 +; GFX942-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-NEXT: v_mov_b32_e32 v1, s0 ; GFX942-NEXT: global_store_dword v0, v1, s[2:3] ; GFX942-NEXT: s_endpgm ; @@ -1067,7 +1070,7 @@ define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg % ; GFX90a-NEXT: .p2align 8 ; GFX90a-NEXT: ; %bb.2: ; GFX90a-NEXT: .LBB21_0: -; GFX90a-NEXT: s_load_dword s0, s[4:5], 0x28 +; GFX90a-NEXT: s_load_dword s0, s[4:5], 0x38 ; GFX90a-NEXT: v_mov_b32_e32 v0, 0 ; GFX90a-NEXT: s_waitcnt lgkmcnt(0) ; GFX90a-NEXT: v_mov_b32_e32 v1, s0 @@ -1079,7 +1082,7 @@ define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg % ; GFX1250-NEXT: global_prefetch_b8 v0, s[0:1] scope:SCOPE_SE ; 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: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s12 +; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s16 ; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3] ; GFX1250-NEXT: s_endpgm %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() diff --git a/llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll b/llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll index 5fd6b75662efb..82abb64c218c2 100644 --- a/llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll +++ b/llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll @@ -252,9 +252,9 @@ define amdgpu_kernel void @local_store_i48(ptr addrspace(3) %ptr, i48 %arg) #0 { define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 { ; HAWAII-LABEL: local_store_i65: ; HAWAII: ; %bb.0: -; HAWAII-NEXT: s_load_dword s2, s[8:9], 0x4 +; HAWAII-NEXT: s_load_dword s2, s[8:9], 0x6 ; HAWAII-NEXT: s_load_dword s3, s[8:9], 0x0 -; HAWAII-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x2 +; HAWAII-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x4 ; HAWAII-NEXT: s_mov_b32 m0, -1 ; HAWAII-NEXT: s_waitcnt lgkmcnt(0) ; HAWAII-NEXT: s_and_b32 s2, s2, 1 @@ -268,9 +268,9 @@ define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 { ; ; FIJI-LABEL: local_store_i65: ; FIJI: ; %bb.0: -; FIJI-NEXT: s_load_dword s2, s[8:9], 0x10 +; FIJI-NEXT: s_load_dword s2, s[8:9], 0x18 ; FIJI-NEXT: s_load_dword s3, s[8:9], 0x0 -; FIJI-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x8 +; FIJI-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x10 ; FIJI-NEXT: s_mov_b32 m0, -1 ; FIJI-NEXT: s_waitcnt lgkmcnt(0) ; FIJI-NEXT: s_and_b32 s2, s2, 1 @@ -284,9 +284,9 @@ define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 { ; ; GFX9-LABEL: local_store_i65: ; GFX9: ; %bb.0: -; GFX9-NEXT: s_load_dword s2, s[8:9], 0x10 +; GFX9-NEXT: s_load_dword s2, s[8:9], 0x18 ; GFX9-NEXT: s_load_dword s3, s[8:9], 0x0 -; GFX9-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x8 +; GFX9-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x10 ; GFX9-NEXT: s_waitcnt lgkmcnt(0) ; GFX9-NEXT: s_and_b32 s2, s2, 1 ; GFX9-NEXT: v_mov_b32_e32 v2, s3 @@ -300,9 +300,9 @@ define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 { ; GFX10-LABEL: local_store_i65: ; GFX10: ; %bb.0: ; GFX10-NEXT: s_clause 0x2 -; GFX10-NEXT: s_load_dword s2, s[8:9], 0x10 +; GFX10-NEXT: s_load_dword s2, s[8:9], 0x18 ; GFX10-NEXT: s_load_dword s3, s[8:9], 0x0 -; GFX10-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x8 +; GFX10-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x10 ; GFX10-NEXT: s_waitcnt lgkmcnt(0) ; GFX10-NEXT: s_and_b32 s2, s2, 1 ; GFX10-NEXT: v_mov_b32_e32 v2, s3 @@ -316,9 +316,9 @@ define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 { ; GFX11-LABEL: local_store_i65: ; GFX11: ; %bb.0: ; GFX11-NEXT: s_clause 0x2 -; GFX11-NEXT: s_load_b32 s2, s[4:5], 0x10 +; GFX11-NEXT: s_load_b32 s2, s[4:5], 0x18 ; GFX11-NEXT: s_load_b32 s3, s[4:5], 0x0 -; GFX11-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 +; GFX11-NEXT: s_load_b64 s[0:1], s[4:5], 0x10 ; GFX11-NEXT: s_waitcnt lgkmcnt(0) ; GFX11-NEXT: s_and_b32 s2, s2, 1 ; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) diff --git a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp index a082adbf6565e..e57461ee48ac3 100644 --- a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp +++ b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp @@ -43,18 +43,27 @@ 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-i128:128"); 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-i128:128"); // 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-i128:128"); // but that r600 does not. EXPECT_EQ(UpgradeDataLayoutString("e-p:32:32-G1", "r600"), "m:e-e-p:32:32-G1"); + // Check that AMDGCN targets don't add an already declared i128 alignment, + // and that r600 never gains one. + EXPECT_EQ( + UpgradeDataLayoutString("e-p:64:64-i128:64-G1", "amdgcn"), + "m:e-e-p:64:64-i128:64-G1-ni:7:8:9-p7:160:256:256:32-p8:128:128:128:" + "48-p9:192:256:256:32"); + EXPECT_EQ(UpgradeDataLayoutString("e-p:32:32-i64:64-G1", "r600"), + "m:e-e-p:32:32-i64:64-G1"); + // Ensure that the non-integral direction for address space 8 doesn't get // added in to pointer declarations. EXPECT_EQ( @@ -66,7 +75,7 @@ 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-i128:128"); // Check that SystemZ adds -S64 if needed. EXPECT_EQ(UpgradeDataLayoutString( @@ -158,24 +167,27 @@ 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-i128:128"); 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-i128:128"); 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-i128:128"); // 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-" + "i128:128"); 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-" + "i128:128"); 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-" + "i128:128"); // Check that SPIR & SPIRV targets don't add -G1 if there is already a -G // flag. @@ -218,7 +230,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-" + "i128:128"); // Check that SPIR & SPIRV targets add G1 if it's not present. EXPECT_EQ(UpgradeDataLayoutString("", "spir"), "G1"); diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp index a3819df4f8a84..c8bcf259d24d7 100644 --- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp +++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp @@ -170,7 +170,8 @@ static Value getKnownOrOcklDim(RewriterBase &rewriter, static constexpr StringLiteral amdgcnDataLayout = "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:" + "-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-i128:128-" + "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/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir index 68a5328b8eb77..396f9bb63d84f 100755 --- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir +++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir @@ -3,7 +3,7 @@ // RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx950 index-bitwidth=32' -split-input-file | FileCheck --check-prefix=CHECK32 %s // CHECK-LABEL: @test_module -// CHECK-SAME: llvm.data_layout = "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-SAME: llvm.data_layout = "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-i128:128-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" gpu.module @test_module { // CHECK-LABEL: func @gpu_index_ops()