[clang] [llvm] [mlir] [AMDGPU] fix alignment of i128 (PR #213523)
Simeon David Schaub via cfe-commits
cfe-commits at lists.llvm.org
Sun Aug 2 04:57:56 PDT 2026
https://github.com/simeonschaub updated https://github.com/llvm/llvm-project/pull/213523
>From 9d08f060ca5c8b42b08b1990fc839a400da354c0 Mon Sep 17 00:00:00 2001
From: Simeon David Schaub <simeon at schaub.rocks>
Date: Sun, 2 Aug 2026 10:01:56 +0000
Subject: [PATCH 1/3] [AMDGPU] Give i128 an ABI alignment of 16 in the data
layout
The AMDGPU data layout has no entry for i128, so `DataLayout` falls back
to the next-lower integer entry (`i64:64`) and gives i128 an ABI
alignment of 8. Clang, however, aligns `__int128` to 16 bytes on AMDGPU:
`AMDGPUTargetInfo` never overrides `Int128Align`, so it keeps the
128-bit default from `TargetInfo`. NVPTX already spells `i128:128` out
in `computeNVPTXDataLayout` for the same reason.
Any frontend that computes aggregate offsets itself and then emits an
LLVM struct type without explicit padding therefore ends up with a
layout that LLVM disagrees with. For kernel arguments that is an ABI
break rather than a missed optimization: `getExplicitKernArgSize()`
sizes the kernarg slot from the data layout, so the tail of an argument
containing an i128 is never copied into the kernarg segment, and loads
of it read past the end of the segment. A kernel taking
`struct { long a; __int128 b; }` by reference gets a 24-byte slot for a
32-byte object, and the load of `b` runs 8 bytes off the end.
Add the missing `i128:128` entry, and auto-upgrade existing AMDGCN data
layout strings that lack one. r600 is left alone: its pointers are
32-bit, so `TargetInfo::hasInt128Type()` is false there and Clang never
forms an `__int128`.
Assisted-by: Claude Code (claude-opus-5)
---
clang/test/CodeGen/target-data.c | 4 +-
clang/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl | 2 +-
llvm/lib/IR/AutoUpgrade.cpp | 6 ++
llvm/lib/TargetParser/TargetDataLayout.cpp | 10 ++-
.../CodeGen/AMDGPU/kernarg-i128-alignment.ll | 82 +++++++++++++++++++
.../Bitcode/DataLayoutUpgradeTest.cpp | 34 +++++---
.../Conversion/GPUToROCDL/gpu-to-rocdl.mlir | 2 +-
7 files changed, 123 insertions(+), 17 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll
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/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/CodeGen/AMDGPU/kernarg-i128-alignment.ll b/llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll
new file mode 100644
index 0000000000000..32e4dcb90e329
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll
@@ -0,0 +1,82 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1100 < %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:1], 0x20
+; CHECK: .amdhsa_kernarg_size 48
+
+; 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:1], 0x20
+; CHECK: .amdhsa_kernarg_size 48
+
+; 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:1], 0x10
+; CHECK: .amdhsa_kernarg_size 40
+
+; 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.
+; CHECK: .amdgpu_metadata
+
+; CHECK: .name: s
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .size: 32
+; CHECK: .kernarg_segment_size: 48
+; CHECK: .name: kernarg_i128
+;
+; CHECK: .name: s
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .size: 32
+; CHECK: .kernarg_segment_size: 48
+; 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: .kernarg_segment_size: 40
+; 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/unittests/Bitcode/DataLayoutUpgradeTest.cpp b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
index a082adbf6565e..49c5055076a59 100644
--- a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
+++ b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
@@ -43,18 +43,26 @@ 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 +74,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 +166,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 +229,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/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()
>From 64a7c62de9512873f7539efbd8545d1564442fb7 Mon Sep 17 00:00:00 2001
From: Simeon David Schaub <simeon at schaub.rocks>
Date: Sun, 2 Aug 2026 12:27:40 +0200
Subject: [PATCH 2/3] fix formatting
---
llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp | 7 ++++---
1 file changed, 4 insertions(+), 3 deletions(-)
diff --git a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
index 49c5055076a59..e57461ee48ac3 100644
--- a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
+++ b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
@@ -57,9 +57,10 @@ TEST(DataLayoutUpgradeTest, ValidDataLayoutUpgrade) {
// 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: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");
>From 93f20243bea32641a9e6ddff3320836d1ce639d8 Mon Sep 17 00:00:00 2001
From: Simeon David Schaub <simeon at schaub.rocks>
Date: Sun, 2 Aug 2026 11:56:30 +0000
Subject: [PATCH 3/3] [AMDGPU] Update MLIR and test expectations for the i128
alignment change
MLIR's GPUToROCDL conversion keeps its own copy of the AMDGPU data layout
string instead of going through computeAMDDataLayout(), so it needs the
i128 entry too.
The rest is test fallout. Kernarg offsets shift, printed load and store
alignments go from 8 to 16, and the folded offset in
lower-buffer-fat-pointers-constants.ll grows from 48 to 64. Note that the
change reaches beyond i128: DataLayout resolves an integer type with no
entry of its own to the next larger one, so i65..i127 pick up the new
16-byte alignment as well. That is why kernel-argument-dag-lowering.ll
grows i65_arg's kernarg segment from 24 to 32 (a hand-written check that
update_llc_test_checks.py does not regenerate) and why opencl-printf.ll
and the HIP printf test see their i96 stores change.
Two changes are more than mechanical. In trunc-brc.ll the i96/i128/i160
loads carry no explicit alignment, so they now inherit 16 and the backend
widens S_LOAD_DWORDX2_IMM to S_LOAD_DWORDX4_IMM;
regbankselect-amdgcn.s.buffer.load.ll likewise stops splitting an i128
buffer load. Both are improvements.
kernarg-i128-alignment.ll checks .kernarg_segment_align rather than
.kernarg_segment_size, since the latter also covers the 256 bytes of
hidden implicit arguments and says nothing about the layout under test.
Assisted-by: Claude Code (claude-opus-5)
---
clang/test/CodeGenHIP/printf_nonhostcall.cpp | 4 +-
.../CostModel/AMDGPU/load-to-trunc.ll | 4 +-
.../AMDGPU/GlobalISel/function-returns.ll | 12 +-
.../GlobalISel/irtranslator-function-args.ll | 12 +-
.../regbankselect-amdgcn.s.buffer.load.ll | 56 ++++-----
.../CodeGen/AMDGPU/GlobalISel/trunc-brc.ll | 106 +++++++++---------
llvm/test/CodeGen/AMDGPU/add_i128.ll | 2 +-
.../amdgpu-attributor-min-agpr-alloc.ll | 2 +-
llvm/test/CodeGen/AMDGPU/ctpop64.ll | 69 ++++++------
.../CodeGen/AMDGPU/kernarg-i128-alignment.ll | 18 +--
llvm/test/CodeGen/AMDGPU/kernel-args.ll | 49 ++++----
.../AMDGPU/kernel-argument-dag-lowering.ll | 15 +--
.../lower-buffer-fat-pointers-constants.ll | 2 +-
...ffer-fat-pointers-contents-legalization.ll | 24 ++--
llvm/test/CodeGen/AMDGPU/mul.ll | 32 +++---
llvm/test/CodeGen/AMDGPU/opencl-printf.ll | 8 +-
.../preload-implicit-kernargs-IR-lowering.ll | 4 +-
.../AMDGPU/preload-implicit-kernargs.ll | 17 +--
llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll | 20 ++--
.../GPUToROCDL/LowerGpuOpsToROCDLOps.cpp | 3 +-
20 files changed, 233 insertions(+), 226 deletions(-)
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/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
index 32e4dcb90e329..fc919a72a9a38 100644
--- a/llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll
+++ b/llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll
@@ -15,35 +15,36 @@
; 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:1], 0x20
-; CHECK: .amdhsa_kernarg_size 48
+; 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:1], 0x20
-; CHECK: .amdhsa_kernarg_size 48
+; 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:1], 0x10
-; CHECK: .amdhsa_kernarg_size 40
+; 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_size: 48
+; CHECK: .kernarg_segment_align: 16
; CHECK: .name: kernarg_i128
;
; CHECK: .name: s
; CHECK-NEXT: .offset: 16
; CHECK-NEXT: .size: 32
-; CHECK: .kernarg_segment_size: 48
; CHECK: .name: kernarg_i128_after_i8
;
; CHECK: .name: a
@@ -52,7 +53,6 @@
; CHECK: .name: b
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8
-; CHECK: .kernarg_segment_size: 40
; CHECK: .name: kernarg_i128_scalar
define amdgpu_kernel void @kernarg_i128(ptr addrspace(1) %out,
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> <i32 0, i32 1, i32 2, i32 3, i32 poison>
; CHECK-NEXT: [[RET_PARTS_0:%.*]] = shufflevector <5 x i32> poison, <5 x i32> [[RET_EXT_0]], <5 x i32> <i32 5, i32 6, i32 7, i32 8, i32 4>
-; 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> <i32 0, i32 1, i32 2, i32 3>
-; 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> <i32 0, i32 1, i32 2, i32 3, i32 poison, i32 poison, i32 poison, i32 poison>
; CHECK-NEXT: [[RET_PARTS_0:%.*]] = shufflevector <8 x i32> poison, <8 x i32> [[RET_EXT_0]], <8 x i32> <i32 8, i32 9, i32 10, i32 11, i32 4, i32 5, i32 6, i32 7>
-; 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> <i32 0, i32 1, i32 2, i32 3, i32 poison, i32 poison, i32 poison, i32 poison>
; CHECK-NEXT: [[RET_PARTS_4:%.*]] = shufflevector <8 x i32> [[RET_PARTS_0]], <8 x i32> [[RET_EXT_4]], <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 8, i32 9, i32 10, i32 11>
; 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> <i32 0, i32 1, i32 2, i32 3>
-; 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> <i32 4, i32 5, i32 6, i32 7>
-; 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/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";
More information about the cfe-commits
mailing list