[clang] [llvm] [IR] change i128 datalayout default to match clang (PR #214518)
Simeon David Schaub via cfe-commits
cfe-commits at lists.llvm.org
Fri Aug 7 00:04:10 PDT 2026
https://github.com/simeonschaub updated https://github.com/llvm/llvm-project/pull/214518
>From 75840527fdc3254a60e0d9f539e4aa55001ba8a0 Mon Sep 17 00:00:00 2001
From: Simeon David Schaub <simeon at schaub.rocks>
Date: Thu, 6 Aug 2026 15:48:16 +0000
Subject: [PATCH 1/2] [IR] Give i128 a default ABI alignment of 16 in the data
layout
Data layout strings without an explicit i128 entry used to fall back to
the alignment of the next smaller integer entry, typically i64:64. Clang
however defaults Int128Align to 128 bits, so on every target that didn't
spell out an i128 entry, the struct layouts computed by LLVM disagreed
with the ABI implemented by the frontend. On AMDGPU this caused
miscompiles of kernels taking i128 arguments, since the kernarg segment
is laid out with the LLVM rules (see
https://github.com/JuliaGPU/AMDGPU.jl/issues/1002).
Fix this by adding i128:128:128 to the default integer specifications,
matching the Clang default. SystemZ is the only target whose ABI aligns
__int128 to only 8 bytes, so its data layout now spells out i128:64
explicitly, and data layouts of existing SystemZ IR are upgraded
accordingly. All other targets either already declare i128:128 or
inherit the new default.
Assisted-by: Claude Code (claude-fable-5)
---
clang/test/CodeGen/target-data.c | 6 +-
llvm/docs/LangRef.md | 4 +-
llvm/docs/ReleaseNotes.md | 8 ++
llvm/lib/IR/AutoUpgrade.cpp | 10 ++-
llvm/lib/IR/DataLayout.cpp | 9 +-
llvm/lib/TargetParser/TargetDataLayout.cpp | 4 +
llvm/test/Bitcode/upgrade-datalayout6.ll | 9 ++
.../CodeGen/AMDGPU/kernarg-i128-alignment.ll | 83 +++++++++++++++++++
.../Bitcode/DataLayoutUpgradeTest.cpp | 11 ++-
9 files changed, 132 insertions(+), 12 deletions(-)
create mode 100644 llvm/test/Bitcode/upgrade-datalayout6.ll
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..2373eefa2e4b2 100644
--- a/clang/test/CodeGen/target-data.c
+++ b/clang/test/CodeGen/target-data.c
@@ -207,11 +207,11 @@
// RUN: FileCheck %s -check-prefix=SYSTEMZ
// RUN: %clang_cc1 -triple s390x-unknown -target-cpu z13 -target-feature +soft-float -o - -emit-llvm %s | \
// RUN: FileCheck %s -check-prefix=SYSTEMZ
-// SYSTEMZ: target datalayout = "E-S64-m:e-i1:8:16-i8:8:16-i64:64-f128:64-v128:64-a:8:16-n32:64"
+// SYSTEMZ: target datalayout = "E-S64-m:e-i1:8:16-i8:8:16-i64:64-i128:64-f128:64-v128:64-a:8:16-n32:64"
// RUN: %clang_cc1 -triple s390x-unknown -target-cpu z13 -o - -emit-llvm %s | \
// RUN: FileCheck %s -check-prefix=SYSTEMZ-VECTOR
-// SYSTEMZ-VECTOR: target datalayout = "E-S64-m:e-i1:8:16-i8:8:16-i64:64-f128:64-v128:64-a:8:16-n32:64"
+// SYSTEMZ-VECTOR: target datalayout = "E-S64-m:e-i1:8:16-i8:8:16-i64:64-i128:64-f128:64-v128:64-a:8:16-n32:64"
// RUN: %clang_cc1 -triple s390x-none-zos -target-cpu z10 -o - -emit-llvm %s | \
// RUN: FileCheck %s -check-prefix=ZOS
@@ -219,7 +219,7 @@
// RUN: FileCheck %s -check-prefix=ZOS
// RUN: %clang_cc1 -triple s390x-none-zos -target-cpu z13 -o - -emit-llvm %s | \
// RUN: FileCheck %s -check-prefix=ZOS
-// ZOS: target datalayout = "E-S64-m:l-p1:32:32-i1:8:16-i8:8:16-i64:64-f128:64-v128:64-a:8:16-n32:64"
+// ZOS: target datalayout = "E-S64-m:l-p1:32:32-i1:8:16-i8:8:16-i64:64-i128:64-f128:64-v128:64-a:8:16-n32:64"
// RUN: %clang_cc1 -triple msp430-unknown -o - -emit-llvm %s | \
// RUN: FileCheck %s -check-prefix=MSP430
diff --git a/llvm/docs/LangRef.md b/llvm/docs/LangRef.md
index 6a3194271b838..02de32e247381 100644
--- a/llvm/docs/LangRef.md
+++ b/llvm/docs/LangRef.md
@@ -3649,6 +3649,7 @@ specifications are given in this list:
- `i32:32:32` - i32 is 32-bit aligned
- `i64:32:64` - i64 has ABI alignment of 32-bits but preferred
alignment of 64-bits
+- `i128:128:128` - i128 is 128-bit aligned
- `f16:16:16` - half is 16-bit aligned
- `f32:32:32` - float is 32-bit aligned
- `f64:64:64` - double is 64-bit aligned
@@ -3668,7 +3669,8 @@ following rules:
the bitwidth then the largest integer type is used. For example,
given the default specifications above, the i7 type will use the
alignment of i8 (next largest) while both i65 and i256 will use the
- alignment of i64 (largest specified).
+ alignment of i128 (i65 because it is the next largest, i256 because
+ i128 is the largest specified).
The function of the data layout string may not be what you expect.
Notably, this is not a specification from the frontend of what alignment
diff --git a/llvm/docs/ReleaseNotes.md b/llvm/docs/ReleaseNotes.md
index f8fce847b62f3..58d77dac0421b 100644
--- a/llvm/docs/ReleaseNotes.md
+++ b/llvm/docs/ReleaseNotes.md
@@ -52,6 +52,14 @@ Makes programs 10x faster by doing Special New Thing.
### Changes to the LLVM IR
+* The default ABI and preferred alignment of `i128` in the data layout is now
+ 128 bits, matching the `__int128` ABI implemented by Clang. Previously, data
+ layout strings without an explicit `i128` entry fell back to the alignment of
+ the next smaller specified integer type, typically `i64`. Targets whose ABI
+ aligns `i128` to fewer than 128 bits (currently only SystemZ) now spell this
+ out explicitly in their data layout string, and data layouts of existing
+ SystemZ IR are upgraded accordingly.
+
### Changes to LLVM infrastructure
### Changes to building LLVM
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index d2216e4072337..bbc524afc218c 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -7270,10 +7270,16 @@ std::string llvm::UpgradeDataLayoutString(StringRef DL, StringRef TT) {
}
if (T.isSystemZ() && !DL.empty()) {
+ Res = DL.str();
+ // i128 used to be aligned only to 64 bits by way of the i64:64 entry.
+ // Now that the default i128 alignment is 128 bits, this needs to be
+ // spelled out explicitly.
+ if (!DL.contains("-i128") && !DL.starts_with("i128"))
+ Res.append("-i128:64");
// Make sure the stack alignment is present.
if (!DL.contains("-S64"))
- return "E-S64" + DL.drop_front(1).str();
- return DL.str();
+ Res = "E-S64" + StringRef(Res).drop_front(1).str();
+ return Res;
}
auto AddPtr32Ptr64AddrSpaces = [&DL, &Res]() {
diff --git a/llvm/lib/IR/DataLayout.cpp b/llvm/lib/IR/DataLayout.cpp
index 82b33887b81f2..a480768947d34 100644
--- a/llvm/lib/IR/DataLayout.cpp
+++ b/llvm/lib/IR/DataLayout.cpp
@@ -178,10 +178,11 @@ struct LessPointerAddrSpace {
// Default primitive type specifications.
// NOTE: These arrays must be sorted by type bit width.
constexpr DataLayout::PrimitiveSpec DefaultIntSpecs[] = {
- {8, Align::Constant<1>(), Align::Constant<1>()}, // i8:8:8
- {16, Align::Constant<2>(), Align::Constant<2>()}, // i16:16:16
- {32, Align::Constant<4>(), Align::Constant<4>()}, // i32:32:32
- {64, Align::Constant<4>(), Align::Constant<8>()}, // i64:32:64
+ {8, Align::Constant<1>(), Align::Constant<1>()}, // i8:8:8
+ {16, Align::Constant<2>(), Align::Constant<2>()}, // i16:16:16
+ {32, Align::Constant<4>(), Align::Constant<4>()}, // i32:32:32
+ {64, Align::Constant<4>(), Align::Constant<8>()}, // i64:32:64
+ {128, Align::Constant<16>(), Align::Constant<16>()}, // i128:128:128
};
constexpr DataLayout::PrimitiveSpec DefaultFloatSpecs[] = {
{16, Align::Constant<2>(), Align::Constant<2>()}, // f16:16:16
diff --git a/llvm/lib/TargetParser/TargetDataLayout.cpp b/llvm/lib/TargetParser/TargetDataLayout.cpp
index 8b6f46642e4fa..b0f0348ce4e1d 100644
--- a/llvm/lib/TargetParser/TargetDataLayout.cpp
+++ b/llvm/lib/TargetParser/TargetDataLayout.cpp
@@ -389,6 +389,10 @@ static std::string computeSystemZDataLayout(const Triple &TT) {
// 64-bit integers are naturally aligned.
Ret += "-i64:64";
+ // 128-bit integers are aligned only to 64 bits, unlike the 128-bit default
+ // of the data layout.
+ Ret += "-i128:64";
+
// 128-bit floats are aligned only to 64 bits.
Ret += "-f128:64";
diff --git a/llvm/test/Bitcode/upgrade-datalayout6.ll b/llvm/test/Bitcode/upgrade-datalayout6.ll
new file mode 100644
index 0000000000000..8bbcfdcd77d22
--- /dev/null
+++ b/llvm/test/Bitcode/upgrade-datalayout6.ll
@@ -0,0 +1,9 @@
+; Test that an explicit i128:64 entry is added to SystemZ data layouts, which
+; predate the change of the default i128 alignment to 128 bits.
+;
+; RUN: llvm-as %s -o - | llvm-dis - | FileCheck %s
+
+target datalayout = "E-m:e-i1:8:16-i8:8:16-i64:64-f128:64-v128:64-a:8:16-n32:64-S64"
+target triple = "s390x-unknown-linux-gnu"
+
+; CHECK: target datalayout = "E-m:e-i1:8:16-i8:8:16-i64:64-f128:64-v128:64-a:8:16-n32:64-S64-i128:64"
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..dacbaedcaa37f
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/kernarg-i128-alignment.ll
@@ -0,0 +1,83 @@
+; RUN: llc -mtriple=amdgpu11.00-amd-amdhsa < %s | FileCheck %s
+
+; The 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. This now comes from the default i128:128 entry of the data
+; layout; it used to be inherited from the i64:64 entry, and the layout LLVM
+; computed for an aggregate would then disagree with the one a frontend used
+; when it emitted the field offsets.
+;
+; For kernel arguments that disagreement is an ABI break rather than a missed
+; optimization: the kernarg slot is sized from the data layout, so the tail of
+; the argument is never copied into the kernarg segment and loads of it run off
+; the end of the segment.
+
+; struct S { i64 a; i128 b; }: ABI align 16, size 32, offsetof(b) == 16.
+; The byref slot must be 32 bytes, not 24, and `b` must be loaded from 0x20
+; (kernarg base 16 plus a field offset of 16) which stays inside the segment.
+; CHECK-LABEL: {{^}}kernarg_i128:
+; CHECK: s_load_b128 s[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], 0x20
+
+; An i128 following a smaller member is padded out to offset 16 rather than
+; packed at offset 8.
+; CHECK-LABEL: {{^}}kernarg_i128_after_i8:
+; CHECK: s_load_b128 s[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], 0x20
+
+; A bare i128 kernel argument is 16-byte aligned in the kernarg segment, so it
+; starts at 16 (not 8) and the argument after it at 32 (not 24).
+; CHECK-LABEL: {{^}}kernarg_i128_scalar:
+; CHECK: s_load_b128 s[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], 0x10
+
+; The kernel metadata is emitted once, after every function, so the per-kernel
+; argument offsets are checked here in order rather than under each label.
+;
+; .kernarg_segment_size is deliberately not checked: it also covers the 256
+; bytes of hidden implicit arguments, which is noise for what this test pins
+; down. The per-argument .offset/.size entries and .kernarg_segment_align are
+; the layout facts that matter.
+; CHECK: .amdgpu_metadata
+
+; CHECK: .name: s
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .size: 32
+; CHECK: .kernarg_segment_align: 16
+; CHECK: .name: kernarg_i128
+;
+; CHECK: .name: s
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .size: 32
+; CHECK: .name: kernarg_i128_after_i8
+;
+; CHECK: .name: a
+; CHECK-NEXT: .offset: 16
+; CHECK-NEXT: .size: 16
+; CHECK: .name: b
+; CHECK-NEXT: .offset: 32
+; CHECK-NEXT: .size: 8
+; CHECK: .name: kernarg_i128_scalar
+
+define amdgpu_kernel void @kernarg_i128(ptr addrspace(1) %out,
+ ptr addrspace(4) byref({ i64, i128 }) align 16 %s) {
+ %pb = getelementptr inbounds i8, ptr addrspace(4) %s, i64 16
+ %b = load i128, ptr addrspace(4) %pb, align 16
+ store i128 %b, ptr addrspace(1) %out, align 16
+ ret void
+}
+
+define amdgpu_kernel void @kernarg_i128_after_i8(ptr addrspace(1) %out,
+ ptr addrspace(4) byref({ i8, i128 }) align 16 %s) {
+ %pb = getelementptr inbounds i8, ptr addrspace(4) %s, i64 16
+ %b = load i128, ptr addrspace(4) %pb, align 16
+ store i128 %b, ptr addrspace(1) %out, align 16
+ ret void
+}
+
+define amdgpu_kernel void @kernarg_i128_scalar(ptr addrspace(1) %out, i128 %a, i64 %b) {
+ %ext = zext i64 %b to i128
+ %sum = add i128 %a, %ext
+ store i128 %sum, ptr addrspace(1) %out, align 16
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
index a082adbf6565e..d9cf400e799d4 100644
--- a/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
+++ b/llvm/unittests/Bitcode/DataLayoutUpgradeTest.cpp
@@ -68,11 +68,18 @@ TEST(DataLayoutUpgradeTest, ValidDataLayoutUpgrade) {
"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");
- // Check that SystemZ adds -S64 if needed.
+ // Check that SystemZ adds -S64 and -i128:64 if needed.
EXPECT_EQ(UpgradeDataLayoutString(
"E-m:e-i1:8:16-i8:8:16-i64:64-f128:64-v128:64-a:8:16-n32:64",
"systemz"),
- "E-S64-m:e-i1:8:16-i8:8:16-i64:64-f128:64-v128:64-a:8:16-n32:64");
+ "E-S64-m:e-i1:8:16-i8:8:16-i64:64-f128:64-v128:64-a:8:16-n32:64-"
+ "i128:64");
+ // but doesn't add -i128:64 if it is already present.
+ EXPECT_EQ(UpgradeDataLayoutString("E-m:e-i1:8:16-i8:8:16-i64:64-i128:64-"
+ "f128:64-v128:64-a:8:16-n32:64-S64",
+ "systemz"),
+ "E-m:e-i1:8:16-i8:8:16-i64:64-i128:64-f128:64-v128:64-a:8:16-n32:"
+ "64-S64");
// Check that RISCV64 upgrades -n64 to -n32:64.
EXPECT_EQ(UpgradeDataLayoutString("e-m:e-p:64:64-i64:64-i128:128-n64-S128",
>From bd5980c74a12f74280e0b6ebfb9dba263a39307f Mon Sep 17 00:00:00 2001
From: Simeon David Schaub <simeon at schaub.rocks>
Date: Thu, 6 Aug 2026 15:48:17 +0000
Subject: [PATCH 2/2] [Tests] Update test expectations for the i128 default
alignment change
Mostly mechanical regeneration with the UTC scripts: wide integer
(i65 and up) allocas, loads and stores in modules without an explicit
i128 data layout entry now get an ABI alignment of 16, and AMDGPU
kernarg offsets after an i128 argument shift to the next 16-byte slot.
Notable non-mechanical updates:
- llubi/loadstore_{be,le}.ll: the alignment-less b256 load would now
imply an alignment of 16 and trap on the 8-aligned alloca, so it gets
an explicit align 8.
- InstCombine/select.ll (test86): the i128 loads are no longer
speculatable across the select (required alignment 16 exceeds the
8-aligned allocas), so the select of pointers is kept, matching the
existing behavior on targets that declare i128:128 explicitly.
- OpenMPIRBuilderTest.CreateTask: the {ptr, i128} shareds struct is now
32 bytes instead of 24.
Assisted-by: Claude Code (claude-fable-5)
---
clang/test/CodeGenHIP/printf_nonhostcall.cpp | 4 +-
.../CostModel/AMDGPU/load-to-trunc.ll | 4 +-
.../LoopAccessAnalysis/depend_diff_types.ll | 2 +-
.../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 +++++----
llvm/test/CodeGen/AMDGPU/kernel-args.ll | 61 ++++----
.../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 | 64 ++++-----
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 +--
.../BoundsChecking/runtimes.ll | 131 +++++++++++-------
.../Instrumentation/BoundsChecking/simple.ll | 30 ++--
.../Instrumentation/DataFlowSanitizer/load.ll | 2 +-
.../DataFlowSanitizer/origin_load.ll | 10 +-
.../test/Instrumentation/Instrumentor/cast.ll | 6 +-
.../Instrumentor/cast_crash.ll | 10 +-
.../Instrumentor/load_store_gpu_ind.ll | 8 +-
.../Instrumentation/Instrumentor/numeric.ll | 42 +++---
.../MemorySanitizer/PowerPC32/kernel-ppcle.ll | 12 +-
.../Instrumentation/MemorySanitizer/byval.ll | 14 +-
.../MemorySanitizer/msan_kernel_basic.ll | 14 +-
.../lifetime-markers-on-inputs-2.ll | 6 +-
.../InferAlignment/irregular-size.ll | 6 +-
llvm/test/Transforms/InstCombine/apint-and.ll | 2 +-
.../InstCombine/load-store-forward.ll | 4 +-
llvm/test/Transforms/InstCombine/or-xor.ll | 2 +-
llvm/test/Transforms/InstCombine/select.ll | 13 +-
llvm/test/Transforms/InstCombine/shift.ll | 16 +--
llvm/test/Transforms/InstCombine/xor.ll | 2 +-
.../LoopIdiom/X86/unordered-atomic-memcpy.ll | 8 +-
.../Transforms/LowerMatrixIntrinsics/unary.ll | 4 +-
llvm/test/Transforms/SCCP/apint-bigint2.ll | 4 +-
llvm/test/Transforms/SCCP/ub-shift.ll | 18 +--
...hreading-max-jump-threading-live-blocks.ll | 8 +-
llvm/test/tools/llubi/loadstore_be.ll | 2 +-
llvm/test/tools/llubi/loadstore_le.ll | 2 +-
.../Frontend/OpenMPIRBuilderTest.cpp | 2 +-
46 files changed, 449 insertions(+), 413 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/Analysis/LoopAccessAnalysis/depend_diff_types.ll b/llvm/test/Analysis/LoopAccessAnalysis/depend_diff_types.ll
index 5d59660a68e90..033c7bb3e7871 100644
--- a/llvm/test/Analysis/LoopAccessAnalysis/depend_diff_types.ll
+++ b/llvm/test/Analysis/LoopAccessAnalysis/depend_diff_types.ll
@@ -373,7 +373,7 @@ define void @different_type_sizes_strided_accesses_store_size_exceeds_depdist(pt
; CHECK-NEXT: Dependences:
; CHECK-NEXT: IndirectUnsafe:
; CHECK-NEXT: store i16 0, ptr %gep.iv, align 2 ->
-; CHECK-NEXT: store i128 1, ptr %gep.10.iv, align 4
+; CHECK-NEXT: store i128 1, ptr %gep.10.iv, align 16
; CHECK-EMPTY:
; CHECK-NEXT: Run-time memory checks:
; CHECK-NEXT: Grouped accesses:
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/kernel-args.ll b/llvm/test/CodeGen/AMDGPU/kernel-args.ll
index 5b95ef59499ac..0671aec5ef55c 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:
@@ -4522,7 +4523,7 @@ define amdgpu_kernel void @i65_arg(ptr addrspace(1) nocapture %out, i65 %in) nou
; EG-NEXT: AND_INT * T1.W, PV.W, literal.x,
; EG-NEXT: 3(4.203895e-45), 0(0.000000e+00)
; EG-NEXT: LSHL T1.W, PV.W, literal.x,
-; EG-NEXT: AND_INT * T2.W, KC0[3].Y, 1,
+; EG-NEXT: AND_INT * T2.W, KC0[3].W, 1,
; EG-NEXT: 3(4.203895e-45), 0(0.000000e+00)
; EG-NEXT: LSHL T1.X, PS, PV.W,
; EG-NEXT: LSHL * T1.W, literal.x, PV.W,
@@ -4533,10 +4534,10 @@ define amdgpu_kernel void @i65_arg(ptr addrspace(1) nocapture %out, i65 %in) nou
; EG-NEXT: ADD_INT * T0.W, KC0[2].Y, literal.y,
; EG-NEXT: 2(2.802597e-45), 4(5.605194e-45)
; EG-NEXT: LSHR T2.X, PV.W, literal.x,
-; EG-NEXT: MOV * T3.X, KC0[3].X,
+; EG-NEXT: MOV * T3.X, KC0[3].Z,
; EG-NEXT: 2(2.802597e-45), 0(0.000000e+00)
; EG-NEXT: LSHR T4.X, KC0[2].Y, literal.x,
-; EG-NEXT: MOV * T5.X, KC0[2].W,
+; EG-NEXT: MOV * T5.X, KC0[3].Y,
; EG-NEXT: 2(2.802597e-45), 0(0.000000e+00)
;
; CM-LABEL: i65_arg:
@@ -4553,7 +4554,7 @@ define amdgpu_kernel void @i65_arg(ptr addrspace(1) nocapture %out, i65 %in) nou
; CM-NEXT: AND_INT * T1.W, PV.W, literal.x,
; CM-NEXT: 3(4.203895e-45), 0(0.000000e+00)
; CM-NEXT: LSHL T0.Z, PV.W, literal.x,
-; CM-NEXT: AND_INT * T1.W, KC0[3].Y, 1,
+; CM-NEXT: AND_INT * T1.W, KC0[3].W, 1,
; CM-NEXT: 3(4.203895e-45), 0(0.000000e+00)
; CM-NEXT: LSHL T1.X, PV.W, PV.Z,
; CM-NEXT: LSHL * T1.W, literal.x, PV.Z,
@@ -4562,12 +4563,12 @@ define amdgpu_kernel void @i65_arg(ptr addrspace(1) nocapture %out, i65 %in) nou
; CM-NEXT: MOV * T1.Z, 0.0,
; CM-NEXT: LSHR * T0.X, KC0[2].Y, literal.x,
; CM-NEXT: 2(2.802597e-45), 0(0.000000e+00)
-; CM-NEXT: MOV T2.X, KC0[2].W,
+; CM-NEXT: MOV T2.X, KC0[3].Y,
; CM-NEXT: ADD_INT * T2.W, KC0[2].Y, literal.x,
; CM-NEXT: 4(5.605194e-45), 0(0.000000e+00)
; CM-NEXT: LSHR * T3.X, PV.W, literal.x,
; CM-NEXT: 2(2.802597e-45), 0(0.000000e+00)
-; CM-NEXT: MOV * T4.X, KC0[3].X,
+; CM-NEXT: MOV * T4.X, KC0[3].Z,
; CM-NEXT: LSHR * T5.X, T0.W, literal.x,
; CM-NEXT: 2(2.802597e-45), 0(0.000000e+00)
entry:
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 97aec3a73ec2f..b7d0e44ba1af4 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
@@ -3770,32 +3770,32 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a,
; EG-NEXT: CF_END
; EG-NEXT: PAD
; EG-NEXT: ALU clause starting at 4:
-; EG-NEXT: MULLO_INT * T0.X, KC0[5].X, KC0[8].X,
-; EG-NEXT: MULHI * T0.Y, KC0[5].X, KC0[8].X,
-; EG-NEXT: MULLO_INT * T0.Z, KC0[8].Y, KC0[4].W,
-; EG-NEXT: MULLO_INT * T0.W, KC0[8].X, KC0[5].Y,
-; EG-NEXT: MULHI * T1.X, KC0[5].X, KC0[7].W,
-; EG-NEXT: MULHI * T1.Y, KC0[4].W, KC0[8].X,
-; EG-NEXT: MULHI * T1.Z, KC0[8].Y, KC0[4].W,
-; EG-NEXT: MULLO_INT * T1.W, KC0[8].Y, KC0[5].X,
-; EG-NEXT: MULHI * T2.X, KC0[7].W, KC0[5].Y,
-; EG-NEXT: MULLO_INT * T2.Y, KC0[5].X, KC0[7].W,
-; EG-NEXT: MULHI * T2.Z, KC0[4].W, KC0[7].W,
+; EG-NEXT: MULLO_INT * T0.X, KC0[5].Z, KC0[8].Z,
+; EG-NEXT: MULHI * T0.Y, KC0[5].Z, KC0[8].Z,
+; EG-NEXT: MULLO_INT * T0.Z, KC0[8].W, KC0[5].Y,
+; EG-NEXT: MULLO_INT * T0.W, KC0[8].Z, KC0[5].W,
+; EG-NEXT: MULHI * T1.X, KC0[5].Z, KC0[8].Y,
+; EG-NEXT: MULHI * T1.Y, KC0[5].Y, KC0[8].Z,
+; EG-NEXT: MULHI * T1.Z, KC0[8].W, KC0[5].Y,
+; EG-NEXT: MULLO_INT * T1.W, KC0[8].W, KC0[5].Z,
+; EG-NEXT: MULHI * T2.X, KC0[8].Y, KC0[5].W,
+; EG-NEXT: MULLO_INT * T2.Y, KC0[5].Z, KC0[8].Y,
+; EG-NEXT: MULHI * T2.Z, KC0[5].Y, KC0[8].Y,
; EG-NEXT: ADD_INT T2.W, T2.Y, PS,
-; EG-NEXT: MULLO_INT * T3.X, KC0[4].W, KC0[8].X,
+; EG-NEXT: MULLO_INT * T3.X, KC0[5].Y, KC0[8].Z,
; EG-NEXT: ADDC_UINT T2.Z, T2.Y, T2.Z,
; EG-NEXT: ADDC_UINT T3.W, PS, PV.W,
-; EG-NEXT: MULLO_INT * T2.Y, KC0[7].W, KC0[5].Z,
+; EG-NEXT: MULLO_INT * T2.Y, KC0[8].Y, KC0[6].X,
; EG-NEXT: ADD_INT T2.X, T2.X, PS,
; EG-NEXT: ADD_INT T2.Y, T1.Z, T1.W,
; EG-NEXT: ADD_INT T1.Z, T1.Y, PV.W,
; EG-NEXT: ADD_INT T1.W, T1.X, PV.Z, BS:VEC_120/SCL_212
-; EG-NEXT: MULLO_INT * T1.X, KC0[8].Z, KC0[4].W,
+; EG-NEXT: MULLO_INT * T1.X, KC0[9].X, KC0[5].Y,
; EG-NEXT: ADD_INT T4.X, PV.W, PV.Z,
; EG-NEXT: ADDC_UINT T1.Y, PV.W, PV.Z,
; EG-NEXT: ADD_INT T1.Z, PV.Y, PS,
; EG-NEXT: ADD_INT T0.W, PV.X, T0.W,
-; EG-NEXT: MULLO_INT * T1.X, KC0[7].W, KC0[5].Y,
+; EG-NEXT: MULLO_INT * T1.X, KC0[8].Y, KC0[5].W,
; EG-NEXT: ADD_INT T2.Y, PV.Z, PV.W,
; EG-NEXT: ADDC_UINT T1.Z, T0.Z, PS,
; EG-NEXT: ADD_INT T0.W, T0.Y, PV.Y,
@@ -3811,7 +3811,7 @@ define amdgpu_kernel void @s_mul_i128(ptr addrspace(1) %out, [8 x i32], i128 %a,
; EG-NEXT: ADD_INT * T0.Y, T3.X, T2.W,
; EG-NEXT: LSHR * T1.X, KC0[2].Y, literal.x,
; EG-NEXT: 2(2.802597e-45), 0(0.000000e+00)
-; EG-NEXT: MULLO_INT * T0.X, KC0[4].W, KC0[7].W,
+; EG-NEXT: MULLO_INT * T0.X, KC0[5].Y, KC0[8].Y,
entry:
%mul = mul i128 %a, %b
store i128 %mul, ptr addrspace(1) %out
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 17dd254a9bf33..e84a682d2e678 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, null 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()
@@ -1044,14 +1044,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
;
@@ -1063,7 +1066,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
@@ -1075,7 +1078,7 @@ define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg %
; GFX1250-NEXT: global_prefetch_b8 v0, null scope:SCOPE_SE
; GFX1250-NEXT: v_nop
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
-; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s12
+; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s16
; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]
; GFX1250-NEXT: s_endpgm
%imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
diff --git a/llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll b/llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll
index 5fd6b75662efb..82abb64c218c2 100644
--- a/llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll
+++ b/llvm/test/CodeGen/AMDGPU/store-weird-sizes.ll
@@ -252,9 +252,9 @@ define amdgpu_kernel void @local_store_i48(ptr addrspace(3) %ptr, i48 %arg) #0 {
define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 {
; HAWAII-LABEL: local_store_i65:
; HAWAII: ; %bb.0:
-; HAWAII-NEXT: s_load_dword s2, s[8:9], 0x4
+; HAWAII-NEXT: s_load_dword s2, s[8:9], 0x6
; HAWAII-NEXT: s_load_dword s3, s[8:9], 0x0
-; HAWAII-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x2
+; HAWAII-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x4
; HAWAII-NEXT: s_mov_b32 m0, -1
; HAWAII-NEXT: s_waitcnt lgkmcnt(0)
; HAWAII-NEXT: s_and_b32 s2, s2, 1
@@ -268,9 +268,9 @@ define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 {
;
; FIJI-LABEL: local_store_i65:
; FIJI: ; %bb.0:
-; FIJI-NEXT: s_load_dword s2, s[8:9], 0x10
+; FIJI-NEXT: s_load_dword s2, s[8:9], 0x18
; FIJI-NEXT: s_load_dword s3, s[8:9], 0x0
-; FIJI-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x8
+; FIJI-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x10
; FIJI-NEXT: s_mov_b32 m0, -1
; FIJI-NEXT: s_waitcnt lgkmcnt(0)
; FIJI-NEXT: s_and_b32 s2, s2, 1
@@ -284,9 +284,9 @@ define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 {
;
; GFX9-LABEL: local_store_i65:
; GFX9: ; %bb.0:
-; GFX9-NEXT: s_load_dword s2, s[8:9], 0x10
+; GFX9-NEXT: s_load_dword s2, s[8:9], 0x18
; GFX9-NEXT: s_load_dword s3, s[8:9], 0x0
-; GFX9-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x8
+; GFX9-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x10
; GFX9-NEXT: s_waitcnt lgkmcnt(0)
; GFX9-NEXT: s_and_b32 s2, s2, 1
; GFX9-NEXT: v_mov_b32_e32 v2, s3
@@ -300,9 +300,9 @@ define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 {
; GFX10-LABEL: local_store_i65:
; GFX10: ; %bb.0:
; GFX10-NEXT: s_clause 0x2
-; GFX10-NEXT: s_load_dword s2, s[8:9], 0x10
+; GFX10-NEXT: s_load_dword s2, s[8:9], 0x18
; GFX10-NEXT: s_load_dword s3, s[8:9], 0x0
-; GFX10-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x8
+; GFX10-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x10
; GFX10-NEXT: s_waitcnt lgkmcnt(0)
; GFX10-NEXT: s_and_b32 s2, s2, 1
; GFX10-NEXT: v_mov_b32_e32 v2, s3
@@ -316,9 +316,9 @@ define amdgpu_kernel void @local_store_i65(ptr addrspace(3) %ptr, i65 %arg) #0 {
; GFX11-LABEL: local_store_i65:
; GFX11: ; %bb.0:
; GFX11-NEXT: s_clause 0x2
-; GFX11-NEXT: s_load_b32 s2, s[4:5], 0x10
+; GFX11-NEXT: s_load_b32 s2, s[4:5], 0x18
; GFX11-NEXT: s_load_b32 s3, s[4:5], 0x0
-; GFX11-NEXT: s_load_b64 s[0:1], s[4:5], 0x8
+; GFX11-NEXT: s_load_b64 s[0:1], s[4:5], 0x10
; GFX11-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-NEXT: s_and_b32 s2, s2, 1
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
diff --git a/llvm/test/Instrumentation/BoundsChecking/runtimes.ll b/llvm/test/Instrumentation/BoundsChecking/runtimes.ll
index 2a756197dfd7e..1a0796a4344d9 100644
--- a/llvm/test/Instrumentation/BoundsChecking/runtimes.ll
+++ b/llvm/test/Instrumentation/BoundsChecking/runtimes.ll
@@ -21,7 +21,7 @@ define void @f1(i64 %x) nounwind {
; TR-LABEL: define void @f1(
; TR-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
; TR-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; TR-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; TR-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; TR-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
; TR-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; TR-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -37,7 +37,7 @@ define void @f1(i64 %x) nounwind {
; RT-LABEL: define void @f1(
; RT-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
; RT-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; RT-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; RT-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; RT-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
; RT-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; RT-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -53,7 +53,7 @@ define void @f1(i64 %x) nounwind {
; TR-NOMERGE-LABEL: define void @f1(
; TR-NOMERGE-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
; TR-NOMERGE-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; TR-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; TR-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; TR-NOMERGE-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
; TR-NOMERGE-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; TR-NOMERGE-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -69,7 +69,7 @@ define void @f1(i64 %x) nounwind {
; RT-NOMERGE-LABEL: define void @f1(
; RT-NOMERGE-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
; RT-NOMERGE-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; RT-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; RT-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; RT-NOMERGE-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
; RT-NOMERGE-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; RT-NOMERGE-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -85,7 +85,7 @@ define void @f1(i64 %x) nounwind {
; RTABORT-NOMERGE-LABEL: define void @f1(
; RTABORT-NOMERGE-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
; RTABORT-NOMERGE-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; RTABORT-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; RTABORT-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; RTABORT-NOMERGE-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
; RTABORT-NOMERGE-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; RTABORT-NOMERGE-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -98,26 +98,10 @@ define void @f1(i64 %x) nounwind {
; RTABORT-NOMERGE-NEXT: call void @__ubsan_handle_local_out_of_bounds_abort() #[[ATTR2:[0-9]+]], !nosanitize [[META0]]
; RTABORT-NOMERGE-NEXT: unreachable, !nosanitize [[META0]]
;
-; MINRT-PRESERVE-NOMERGE-LABEL: define void @f1(
-; MINRT-PRESERVE-NOMERGE-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
-; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
-; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
-; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
-; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
-; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP6:%.*]] = or i1 false, [[TMP5]], !nosanitize [[META0]]
-; MINRT-PRESERVE-NOMERGE-NEXT: br i1 [[TMP6]], label %[[TRAP:.*]], label %[[BB7:.*]]
-; MINRT-PRESERVE-NOMERGE: [[BB7]]:
-; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP8:%.*]] = load i128, ptr [[TMP2]], align 4
-; MINRT-PRESERVE-NOMERGE-NEXT: ret void
-; MINRT-PRESERVE-NOMERGE: [[TRAP]]:
-; MINRT-PRESERVE-NOMERGE-NEXT: call preserve_allcc void @__ubsan_handle_local_out_of_bounds_minimal_preserve() #[[ATTR1:[0-9]+]], !nosanitize [[META0]]
-; MINRT-PRESERVE-NOMERGE-NEXT: br label %[[BB7]], !nosanitize [[META0]]
-;
; MINRT-NOMERGE-LABEL: define void @f1(
; MINRT-NOMERGE-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
; MINRT-NOMERGE-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; MINRT-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; MINRT-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; MINRT-NOMERGE-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
; MINRT-NOMERGE-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; MINRT-NOMERGE-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -133,7 +117,7 @@ define void @f1(i64 %x) nounwind {
; MINRTABORT-NOMERGE-LABEL: define void @f1(
; MINRTABORT-NOMERGE-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
; MINRTABORT-NOMERGE-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; MINRTABORT-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; MINRTABORT-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; MINRTABORT-NOMERGE-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
; MINRTABORT-NOMERGE-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; MINRTABORT-NOMERGE-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -146,34 +130,62 @@ define void @f1(i64 %x) nounwind {
; MINRTABORT-NOMERGE-NEXT: call void @__ubsan_handle_local_out_of_bounds_minimal_abort() #[[ATTR2:[0-9]+]], !nosanitize [[META0]]
; MINRTABORT-NOMERGE-NEXT: unreachable, !nosanitize [[META0]]
;
-; TR-GUARD-COMMON-LABEL: define void @f1(
-; TR-GUARD-COMMON-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
-; TR-GUARD-COMMON-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; TR-GUARD-COMMON-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
-; TR-GUARD-COMMON-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
-; TR-GUARD-COMMON-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
-; TR-GUARD-COMMON-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
-; TR-GUARD-COMMON-NEXT: [[TMP6:%.*]] = or i1 false, [[TMP5]], !nosanitize [[META0]]
-;
-; TR-GUARD-THREE: [[TMP7:%.*]] = call i1 @llvm.allow.ubsan.check(i8 3), !nosanitize [[META0]]
-; TR-GUARD-THIRTEEN: [[TMP7:%.*]] = call i1 @llvm.allow.ubsan.check(i8 13), !nosanitize [[META0]]
-;
-; TR-GUARD-COMMON: [[TMP8:%.*]] = and i1 [[TMP6]], [[TMP7]], !nosanitize [[META0]]
-; TR-GUARD-COMMON-NEXT: br i1 [[TMP8]], label %[[TRAP:.*]], label %[[BB9:.*]]
-; TR-GUARD-COMMON: [[BB9]]:
-; TR-GUARD-COMMON-NEXT: [[TMP10:%.*]] = load i128, ptr [[TMP2]], align 4
-; TR-GUARD-COMMON-NEXT: ret void
-; TR-GUARD-COMMON: [[TRAP]]:
+; MINRT-PRESERVE-NOMERGE-LABEL: define void @f1(
+; MINRT-PRESERVE-NOMERGE-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
+; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
+; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
+; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
+; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
+; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
+; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP6:%.*]] = or i1 false, [[TMP5]], !nosanitize [[META0]]
+; MINRT-PRESERVE-NOMERGE-NEXT: br i1 [[TMP6]], label %[[TRAP:.*]], label %[[BB7:.*]]
+; MINRT-PRESERVE-NOMERGE: [[BB7]]:
+; MINRT-PRESERVE-NOMERGE-NEXT: [[TMP8:%.*]] = load i128, ptr [[TMP2]], align 4
+; MINRT-PRESERVE-NOMERGE-NEXT: ret void
+; MINRT-PRESERVE-NOMERGE: [[TRAP]]:
+; MINRT-PRESERVE-NOMERGE-NEXT: call preserve_allcc void @__ubsan_handle_local_out_of_bounds_minimal_preserve() #[[ATTR1:[0-9]+]], !nosanitize [[META0]]
+; MINRT-PRESERVE-NOMERGE-NEXT: br label %[[BB7]], !nosanitize [[META0]]
;
-; TR-GUARD-THREE: call void @llvm.ubsantrap(i8 3) #[[ATTR3:[0-9]+]], !nosanitize [[META0]]
-; TR-GUARD-THIRTEEN: call void @llvm.ubsantrap(i8 13) #[[ATTR3:[0-9]+]], !nosanitize [[META0]]
+; TR-GUARD-THREE-LABEL: define void @f1(
+; TR-GUARD-THREE-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
+; TR-GUARD-THREE-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
+; TR-GUARD-THREE-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
+; TR-GUARD-THREE-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
+; TR-GUARD-THREE-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
+; TR-GUARD-THREE-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
+; TR-GUARD-THREE-NEXT: [[TMP6:%.*]] = or i1 false, [[TMP5]], !nosanitize [[META0]]
+; TR-GUARD-THREE-NEXT: [[TMP7:%.*]] = call i1 @llvm.allow.ubsan.check(i8 3), !nosanitize [[META0]]
+; TR-GUARD-THREE-NEXT: [[TMP8:%.*]] = and i1 [[TMP6]], [[TMP7]], !nosanitize [[META0]]
+; TR-GUARD-THREE-NEXT: br i1 [[TMP8]], label %[[TRAP:.*]], label %[[BB9:.*]]
+; TR-GUARD-THREE: [[BB9]]:
+; TR-GUARD-THREE-NEXT: [[TMP10:%.*]] = load i128, ptr [[TMP2]], align 4
+; TR-GUARD-THREE-NEXT: ret void
+; TR-GUARD-THREE: [[TRAP]]:
+; TR-GUARD-THREE-NEXT: call void @llvm.ubsantrap(i8 3) #[[ATTR3:[0-9]+]], !nosanitize [[META0]]
+; TR-GUARD-THREE-NEXT: unreachable, !nosanitize [[META0]]
;
-; TR-GUARD-COMMON: unreachable, !nosanitize [[META0]]
+; TR-GUARD-THIRTEEN-LABEL: define void @f1(
+; TR-GUARD-THIRTEEN-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
+; TR-GUARD-THIRTEEN-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
+; TR-GUARD-THIRTEEN-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
+; TR-GUARD-THIRTEEN-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
+; TR-GUARD-THIRTEEN-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
+; TR-GUARD-THIRTEEN-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
+; TR-GUARD-THIRTEEN-NEXT: [[TMP6:%.*]] = or i1 false, [[TMP5]], !nosanitize [[META0]]
+; TR-GUARD-THIRTEEN-NEXT: [[TMP7:%.*]] = call i1 @llvm.allow.ubsan.check(i8 13), !nosanitize [[META0]]
+; TR-GUARD-THIRTEEN-NEXT: [[TMP8:%.*]] = and i1 [[TMP6]], [[TMP7]], !nosanitize [[META0]]
+; TR-GUARD-THIRTEEN-NEXT: br i1 [[TMP8]], label %[[TRAP:.*]], label %[[BB9:.*]]
+; TR-GUARD-THIRTEEN: [[BB9]]:
+; TR-GUARD-THIRTEEN-NEXT: [[TMP10:%.*]] = load i128, ptr [[TMP2]], align 4
+; TR-GUARD-THIRTEEN-NEXT: ret void
+; TR-GUARD-THIRTEEN: [[TRAP]]:
+; TR-GUARD-THIRTEEN-NEXT: call void @llvm.ubsantrap(i8 13) #[[ATTR3:[0-9]+]], !nosanitize [[META0]]
+; TR-GUARD-THIRTEEN-NEXT: unreachable, !nosanitize [[META0]]
;
; RT-GUARD-LABEL: define void @f1(
; RT-GUARD-SAME: i64 [[X:%.*]]) #[[ATTR0:[0-9]+]] {
; RT-GUARD-NEXT: [[TMP1:%.*]] = mul i64 [[X]], 16
-; RT-GUARD-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; RT-GUARD-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; RT-GUARD-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0:![0-9]+]]
; RT-GUARD-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; RT-GUARD-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -193,6 +205,11 @@ define void @f1(i64 %x) nounwind {
ret void
}
+; TR-GUARD: attributes #[[ATTR0]] = { nounwind }
+; TR-GUARD: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) }
+; TR-GUARD: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
+; TR-GUARD: attributes #[[ATTR3]] = { nomerge noreturn nounwind }
+; TR-GUARD: [[META0]] = !{}
;.
; TR: attributes #[[ATTR0]] = { nounwind }
; TR: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
@@ -218,10 +235,18 @@ define void @f1(i64 %x) nounwind {
; MINRTABORT-NOMERGE: attributes #[[ATTR1:[0-9]+]] = { noreturn nounwind }
; MINRTABORT-NOMERGE: attributes #[[ATTR2]] = { nomerge noreturn nounwind }
;.
-; TR-GUARD: attributes #[[ATTR0]] = { nounwind }
-; TR-GUARD: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) }
-; TR-GUARD: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
-; TR-GUARD: attributes #[[ATTR3]] = { nomerge noreturn nounwind }
+; MINRT-PRESERVE-NOMERGE: attributes #[[ATTR0]] = { nounwind }
+; MINRT-PRESERVE-NOMERGE: attributes #[[ATTR1]] = { nomerge nounwind }
+;.
+; TR-GUARD-THREE: attributes #[[ATTR0]] = { nounwind }
+; TR-GUARD-THREE: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) }
+; TR-GUARD-THREE: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
+; TR-GUARD-THREE: attributes #[[ATTR3]] = { nomerge noreturn nounwind }
+;.
+; TR-GUARD-THIRTEEN: attributes #[[ATTR0]] = { nounwind }
+; TR-GUARD-THIRTEEN: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) }
+; TR-GUARD-THIRTEEN: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
+; TR-GUARD-THIRTEEN: attributes #[[ATTR3]] = { nomerge noreturn nounwind }
;.
; RT-GUARD: attributes #[[ATTR0]] = { nounwind }
; RT-GUARD: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) }
@@ -241,7 +266,13 @@ define void @f1(i64 %x) nounwind {
;.
; MINRTABORT-NOMERGE: [[META0]] = !{}
;.
-; TR-GUARD: [[META0]] = !{}
+; MINRT-PRESERVE-NOMERGE: [[META0]] = !{}
+;.
+; TR-GUARD-THREE: [[META0]] = !{}
+;.
+; TR-GUARD-THIRTEEN: [[META0]] = !{}
;.
; RT-GUARD: [[META0]] = !{}
;.
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; TR-GUARD-COMMON: {{.*}}
diff --git a/llvm/test/Instrumentation/BoundsChecking/simple.ll b/llvm/test/Instrumentation/BoundsChecking/simple.ll
index 81592d79d3a83..90a531b8b2e0c 100644
--- a/llvm/test/Instrumentation/BoundsChecking/simple.ll
+++ b/llvm/test/Instrumentation/BoundsChecking/simple.ll
@@ -171,7 +171,7 @@ define void @f5_as2(i32 %x) nounwind {;
define void @f6(i64 %x) nounwind {
; CHECK-LABEL: @f6(
-; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 8
+; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 16
; CHECK-NEXT: [[TMP2:%.*]] = load i128, ptr [[TMP1]], align 4
; CHECK-NEXT: ret void
;
@@ -183,7 +183,7 @@ define void @f6(i64 %x) nounwind {
define void @f7(i64 %x) nounwind {
; CHECK-LABEL: @f7(
; CHECK-NEXT: [[TMP1:%.*]] = mul i64 [[X:%.*]], 16
-; CHECK-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; CHECK-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; CHECK-NEXT: [[TMP3:%.*]] = sub i64 [[TMP1]], 0, !nosanitize [[META0]]
; CHECK-NEXT: [[TMP4:%.*]] = icmp ult i64 [[TMP3]], 16, !nosanitize [[META0]]
; CHECK-NEXT: [[TMP5:%.*]] = or i1 false, [[TMP4]], !nosanitize [[META0]]
@@ -203,8 +203,8 @@ define void @f7(i64 %x) nounwind {
define void @f8() nounwind {
; CHECK-LABEL: @f8(
-; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 8
-; CHECK-NEXT: [[TMP2:%.*]] = alloca i128, align 8
+; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 16
+; CHECK-NEXT: [[TMP2:%.*]] = alloca i128, align 16
; CHECK-NEXT: [[TMP3:%.*]] = select i1 undef, ptr [[TMP1]], ptr [[TMP2]]
; CHECK-NEXT: [[TMP4:%.*]] = load i128, ptr [[TMP3]], align 4
; CHECK-NEXT: ret void
@@ -218,7 +218,7 @@ define void @f8() nounwind {
define void @f9(ptr %arg) nounwind {
; CHECK-LABEL: @f9(
-; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 8
+; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 16
; CHECK-NEXT: [[TMP2:%.*]] = select i1 undef, ptr [[ARG:%.*]], ptr [[TMP1]]
; CHECK-NEXT: [[TMP3:%.*]] = load i128, ptr [[TMP2]], align 4
; CHECK-NEXT: ret void
@@ -232,9 +232,9 @@ define void @f9(ptr %arg) nounwind {
define void @f10(i64 %x, i64 %y) nounwind {
; CHECK-LABEL: @f10(
; CHECK-NEXT: [[TMP1:%.*]] = mul i64 [[X:%.*]], 16
-; CHECK-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 8
+; CHECK-NEXT: [[TMP2:%.*]] = alloca i128, i64 [[X]], align 16
; CHECK-NEXT: [[TMP3:%.*]] = mul i64 [[Y:%.*]], 16
-; CHECK-NEXT: [[TMP4:%.*]] = alloca i128, i64 [[Y]], align 8
+; CHECK-NEXT: [[TMP4:%.*]] = alloca i128, i64 [[Y]], align 16
; CHECK-NEXT: [[TMP5:%.*]] = select i1 undef, i64 [[TMP1]], i64 [[TMP3]]
; CHECK-NEXT: [[TMP6:%.*]] = select i1 undef, ptr [[TMP2]], ptr [[TMP4]]
; CHECK-NEXT: [[TMP7:%.*]] = sub i64 [[TMP5]], 0, !nosanitize [[META0]]
@@ -532,14 +532,14 @@ define void @scalable_alloca2(i64 %y) nounwind {
; CHECK-NEXT: [[DOTIDX:%.*]] = mul i64 [[Y:%.*]], [[TMP6]]
; CHECK-NEXT: [[TMP7:%.*]] = add i64 0, [[DOTIDX]]
; CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds <vscale x 4 x i64>, ptr [[TMP4]], i64 [[Y]]
-; CHECK-NEXT: [[TMP9:%.*]] = call i64 @llvm.vscale.i64(), !nosanitize [[META0]]
-; CHECK-NEXT: [[TMP10:%.*]] = mul nuw i64 [[TMP9]], 32, !nosanitize [[META0]]
-; CHECK-NEXT: [[TMP11:%.*]] = sub i64 [[TMP2]], [[TMP7]], !nosanitize [[META0]]
-; CHECK-NEXT: [[TMP12:%.*]] = icmp ult i64 [[TMP2]], [[TMP7]], !nosanitize [[META0]]
-; CHECK-NEXT: [[TMP13:%.*]] = icmp ult i64 [[TMP11]], [[TMP10]], !nosanitize [[META0]]
-; CHECK-NEXT: [[TMP14:%.*]] = or i1 [[TMP12]], [[TMP13]], !nosanitize [[META0]]
-; CHECK-NEXT: [[TMP15:%.*]] = icmp slt i64 [[TMP7]], 0, !nosanitize [[META0]]
-; CHECK-NEXT: [[TMP16:%.*]] = or i1 [[TMP15]], [[TMP14]], !nosanitize [[META0]]
+; CHECK-NEXT: [[TMP15:%.*]] = call i64 @llvm.vscale.i64(), !nosanitize [[META0]]
+; CHECK-NEXT: [[TMP9:%.*]] = mul nuw i64 [[TMP15]], 32, !nosanitize [[META0]]
+; CHECK-NEXT: [[TMP10:%.*]] = sub i64 [[TMP2]], [[TMP7]], !nosanitize [[META0]]
+; CHECK-NEXT: [[TMP11:%.*]] = icmp ult i64 [[TMP2]], [[TMP7]], !nosanitize [[META0]]
+; CHECK-NEXT: [[TMP12:%.*]] = icmp ult i64 [[TMP10]], [[TMP9]], !nosanitize [[META0]]
+; CHECK-NEXT: [[TMP13:%.*]] = or i1 [[TMP11]], [[TMP12]], !nosanitize [[META0]]
+; CHECK-NEXT: [[TMP14:%.*]] = icmp slt i64 [[TMP7]], 0, !nosanitize [[META0]]
+; CHECK-NEXT: [[TMP16:%.*]] = or i1 [[TMP14]], [[TMP13]], !nosanitize [[META0]]
; CHECK-NEXT: br i1 [[TMP16]], label [[TRAP:%.*]], label [[TMP17:%.*]]
; CHECK: 16:
; CHECK-NEXT: [[TMP18:%.*]] = load <vscale x 4 x i64>, ptr [[TMP8]], align 4
diff --git a/llvm/test/Instrumentation/DataFlowSanitizer/load.ll b/llvm/test/Instrumentation/DataFlowSanitizer/load.ll
index bf8ba909e0be0..0bac48675855b 100644
--- a/llvm/test/Instrumentation/DataFlowSanitizer/load.ll
+++ b/llvm/test/Instrumentation/DataFlowSanitizer/load.ll
@@ -114,7 +114,7 @@ define i128 @load128(ptr %p) {
; CHECK-NEXT: %[[#WIDE_SHADOW:]] = or i64 %[[#WIDE_SHADOW]], %[[#WIDE_SHADOW_SHIFTED]]
; CHECK-NEXT: %[[#SHADOW:]] = trunc i64 %[[#WIDE_SHADOW]] to i8
; COMBINE_LOAD_PTR-NEXT: %[[#SHADOW:]] = or i8 %[[#SHADOW]], %[[#PS]]
- ; CHECK-NEXT: %a = load i128, ptr %p, align 8
+ ; CHECK-NEXT: %a = load i128, ptr %p, align 16
; CHECK-NEXT: store i8 %[[#SHADOW]], ptr @__dfsan_retval_tls, align [[ALIGN]]
; CHECK-NEXT: ret i128 %a
diff --git a/llvm/test/Instrumentation/DataFlowSanitizer/origin_load.ll b/llvm/test/Instrumentation/DataFlowSanitizer/origin_load.ll
index a0c642a3cd0e1..bdda5fc4d275c 100644
--- a/llvm/test/Instrumentation/DataFlowSanitizer/origin_load.ll
+++ b/llvm/test/Instrumentation/DataFlowSanitizer/origin_load.ll
@@ -224,19 +224,19 @@ define i128 @load128(ptr %p) {
; CHECK-NEXT: %[[#SHADOW_PTR:]] = inttoptr i64 %[[#SHADOW_OFFSET]] to ptr
; CHECK-NEXT: %[[#ORIGIN_ADDR:]] = add i64 %[[#SHADOW_OFFSET]], [[#ORIGIN_BASE]]
; CHECK-NEXT: %[[#ORIGIN1_PTR:]] = inttoptr i64 %[[#ORIGIN_ADDR]] to ptr
- ; CHECK-NEXT: %[[#ORIGIN1:]] = load i32, ptr %[[#ORIGIN1_PTR]], align 8
+ ; CHECK-NEXT: %[[#ORIGIN1:]] = load i32, ptr %[[#ORIGIN1_PTR]], align 16
; CHECK-NEXT: %[[#WIDE_SHADOW1:]] = load i64, ptr %[[#SHADOW_PTR]], align 1
; CHECK-NEXT: %[[#WIDE_SHADOW1_LO:]] = shl i64 %[[#WIDE_SHADOW1]], 32
; CHECK-NEXT: %[[#ORIGIN2_PTR:]] = getelementptr i32, ptr %[[#ORIGIN1_PTR]], i64 1
- ; CHECK-NEXT: %[[#ORIGIN2:]] = load i32, ptr %[[#ORIGIN2_PTR]], align 8
+ ; CHECK-NEXT: %[[#ORIGIN2:]] = load i32, ptr %[[#ORIGIN2_PTR]], align 16
; CHECK-NEXT: %[[#WIDE_SHADOW2_PTR:]] = getelementptr i64, ptr %[[#SHADOW_PTR]], i64 1
; CHECK-NEXT: %[[#WIDE_SHADOW2:]] = load i64, ptr %[[#WIDE_SHADOW2_PTR]], align 1
; CHECK-NEXT: %[[#WIDE_SHADOW:]] = or i64 %[[#WIDE_SHADOW1]], %[[#WIDE_SHADOW2]]
; CHECK-NEXT: %[[#ORIGIN3_PTR:]] = getelementptr i32, ptr %[[#ORIGIN2_PTR]], i64 1
- ; CHECK-NEXT: %[[#ORIGIN3:]] = load i32, ptr %[[#ORIGIN3_PTR]], align 8
+ ; CHECK-NEXT: %[[#ORIGIN3:]] = load i32, ptr %[[#ORIGIN3_PTR]], align 16
; CHECK-NEXT: %[[#WIDE_SHADOW2_LO:]] = shl i64 %[[#WIDE_SHADOW2]], 32
; CHECK-NEXT: %[[#ORIGIN4_PTR:]] = getelementptr i32, ptr %[[#ORIGIN3_PTR]], i64 1
- ; CHECK-NEXT: %[[#ORIGIN4:]] = load i32, ptr %[[#ORIGIN4_PTR]], align 8
+ ; CHECK-NEXT: %[[#ORIGIN4:]] = load i32, ptr %[[#ORIGIN4_PTR]], align 16
; CHECK-NEXT: %[[#WIDE_SHADOW_SHIFTED:]] = lshr i64 %[[#WIDE_SHADOW]], 32
; CHECK-NEXT: %[[#WIDE_SHADOW:]] = or i64 %[[#WIDE_SHADOW]], %[[#WIDE_SHADOW_SHIFTED]]
; CHECK-NEXT: %[[#WIDE_SHADOW_SHIFTED:]] = lshr i64 %[[#WIDE_SHADOW]], 16
@@ -255,7 +255,7 @@ define i128 @load128(ptr %p) {
; COMBINE_LOAD_PTR-NEXT: %[[#NZ:]] = icmp ne i8 %[[#PS]], 0
; COMBINE_LOAD_PTR-NEXT: %[[#ORIGIN:]] = select i1 %[[#NZ]], i32 %[[#PO]], i32 %[[#ORIGIN]]
- ; CHECK-NEXT: %a = load i128, ptr %p, align 8
+ ; CHECK-NEXT: %a = load i128, ptr %p, align 16
; CHECK-NEXT: store i8 %[[#SHADOW]], ptr @__dfsan_retval_tls, align [[ALIGN]]
; CHECK-NEXT: store i32 %[[#ORIGIN]], ptr @__dfsan_retval_origin_tls, align 4
diff --git a/llvm/test/Instrumentation/Instrumentor/cast.ll b/llvm/test/Instrumentation/Instrumentor/cast.ll
index 77d9e0452784d..a1660dba489fa 100644
--- a/llvm/test/Instrumentation/Instrumentor/cast.ll
+++ b/llvm/test/Instrumentation/Instrumentor/cast.ll
@@ -275,15 +275,15 @@ entry:
define i128 @test_ext(i32 %p1) {
; CHECK-LABEL: define i128 @test_ext(
; CHECK-SAME: i32 [[P1:%.*]]) {
-; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 8
+; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 16
; CHECK-NEXT: [[TMP2:%.*]] = alloca i64, align 8
; CHECK-NEXT: [[TMP3:%.*]] = zext i32 [[P1]] to i64
; CHECK-NEXT: call void @__instrumentor_pre_cast(i64 [[TMP3]], i32 12, i32 4, i32 12, i32 16, i32 40) #[[ATTR0]]
; CHECK-NEXT: [[I1:%.*]] = zext i32 [[P1]] to i128
; CHECK-NEXT: store i64 [[TMP3]], ptr [[TMP2]], align 4
-; CHECK-NEXT: store i128 [[I1]], ptr [[TMP1]], align 4
+; CHECK-NEXT: store i128 [[I1]], ptr [[TMP1]], align 16
; CHECK-NEXT: call void @__instrumentor_post_cast_ind(ptr [[TMP2]], i32 12, i32 4, ptr [[TMP1]], i32 12, i32 16, i32 40) #[[ATTR0]]
-; CHECK-NEXT: [[TMP4:%.*]] = load i128, ptr [[TMP1]], align 4
+; CHECK-NEXT: [[TMP4:%.*]] = load i128, ptr [[TMP1]], align 16
; CHECK-NEXT: ret i128 [[TMP4]]
;
%i1 = zext i32 %p1 to i128
diff --git a/llvm/test/Instrumentation/Instrumentor/cast_crash.ll b/llvm/test/Instrumentation/Instrumentor/cast_crash.ll
index f9bf05c9c0cca..6067686424c88 100644
--- a/llvm/test/Instrumentation/Instrumentor/cast_crash.ll
+++ b/llvm/test/Instrumentation/Instrumentor/cast_crash.ll
@@ -5,9 +5,9 @@
define i128 @test_ext(i32 %p1) {
; CHECK-LABEL: define i128 @test_ext(
; CHECK-SAME: i32 [[P1:%.*]]) {
-; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 8
+; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 16
; CHECK-NEXT: [[TMP2:%.*]] = alloca i64, align 8
-; CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP1]], ptr @__instrumentor_value_pack, i64 16, i1 false)
+; CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 16 [[TMP1]], ptr @__instrumentor_value_pack, i64 16, i1 false)
; CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw <{ i32, i32, [4 x i8], i32 }>, ptr [[TMP1]], i32 0, i32 3
; CHECK-NEXT: store i32 [[P1]], ptr [[TMP3]], align 4
; CHECK-NEXT: call void @__instrumentor_pre_function(ptr @test_ext, ptr @__instrumentor_.str.2, i32 1, ptr [[TMP1]], i8 0, i32 4) #[[ATTR1:[0-9]+]]
@@ -17,10 +17,10 @@ define i128 @test_ext(i32 %p1) {
; CHECK-NEXT: call void @__instrumentor_pre_cast(i64 [[TMP6]], i32 12, i32 -1, i32 4, i32 12, i32 -1, i32 16, i32 40, i32 3) #[[ATTR1]]
; CHECK-NEXT: [[I1:%.*]] = zext i32 [[TMP5]] to i128
; CHECK-NEXT: store i64 [[TMP6]], ptr [[TMP2]], align 4
-; CHECK-NEXT: store i128 [[I1]], ptr [[TMP1]], align 4
+; CHECK-NEXT: store i128 [[I1]], ptr [[TMP1]], align 16
; CHECK-NEXT: call void @__instrumentor_post_cast_ind(ptr [[TMP2]], i32 12, i32 -1, i32 4, ptr [[TMP1]], i32 12, i32 -1, i32 16, i32 40, i32 -3) #[[ATTR1]]
-; CHECK-NEXT: [[TMP7:%.*]] = load i128, ptr [[TMP1]], align 4
-; CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP1]], ptr @__instrumentor_value_pack, i64 16, i1 false)
+; CHECK-NEXT: [[TMP7:%.*]] = load i128, ptr [[TMP1]], align 16
+; CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 16 [[TMP1]], ptr @__instrumentor_value_pack, i64 16, i1 false)
; CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw <{ i32, i32, [4 x i8], i32 }>, ptr [[TMP1]], i32 0, i32 3
; CHECK-NEXT: store i32 [[P1]], ptr [[TMP8]], align 4
; CHECK-NEXT: call void @__instrumentor_post_function(ptr @test_ext, ptr @__instrumentor_.str.2, i32 1, ptr [[TMP1]], i8 0, i32 -5) #[[ATTR1]]
diff --git a/llvm/test/Instrumentation/Instrumentor/load_store_gpu_ind.ll b/llvm/test/Instrumentation/Instrumentor/load_store_gpu_ind.ll
index dd88398964b83..781c5af7009b6 100644
--- a/llvm/test/Instrumentation/Instrumentor/load_store_gpu_ind.ll
+++ b/llvm/test/Instrumentation/Instrumentor/load_store_gpu_ind.ll
@@ -123,8 +123,8 @@ define noundef i128 @_Z20store_load_long_longPx(ptr captures(none) noundef initi
; CHECK-LABEL: define noundef i128 @_Z20store_load_long_longPx(
; CHECK-SAME: ptr noundef captures(none) initializes((0, 16)) [[A:%.*]]) {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[TMP0:%.*]] = alloca i128, align 8, addrspace(5)
-; CHECK-NEXT: store i128 5, ptr addrspace(5) [[TMP0]], align 8
+; CHECK-NEXT: [[TMP0:%.*]] = alloca i128, align 16, addrspace(5)
+; CHECK-NEXT: store i128 5, ptr addrspace(5) [[TMP0]], align 16
; CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[TMP0]] to ptr
; CHECK-NEXT: [[TMP2:%.*]] = call ptr @__instrumentor_pre_store_ind(ptr [[A]], i32 0, ptr [[TMP1]], i64 16, i64 8, i32 12, i32 0, i8 1, i8 0) #[[ATTR0]]
; CHECK-NEXT: store i128 5, ptr [[TMP2]], align 8
@@ -132,10 +132,10 @@ define noundef i128 @_Z20store_load_long_longPx(ptr captures(none) noundef initi
; CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i8, ptr [[A]], i64 16
; CHECK-NEXT: [[TMP3:%.*]] = call ptr @__instrumentor_pre_load(ptr [[ARRAYIDX]], i32 0, i64 16, i64 8, i32 12, i32 0, i8 1, i8 0) #[[ATTR0]]
; CHECK-NEXT: [[TMP6:%.*]] = load i128, ptr [[TMP3]], align 8
-; CHECK-NEXT: store i128 [[TMP6]], ptr addrspace(5) [[TMP0]], align 8
+; CHECK-NEXT: store i128 [[TMP6]], ptr addrspace(5) [[TMP0]], align 16
; CHECK-NEXT: [[TMP5:%.*]] = addrspacecast ptr addrspace(5) [[TMP0]] to ptr
; CHECK-NEXT: call void @__instrumentor_post_load_ind(ptr [[ARRAYIDX]], i32 0, ptr [[TMP5]], i64 16, i64 8, i32 12, i32 0, i8 1, i8 0) #[[ATTR0]]
-; CHECK-NEXT: [[TMP4:%.*]] = load i128, ptr [[TMP5]], align 8
+; CHECK-NEXT: [[TMP4:%.*]] = load i128, ptr [[TMP5]], align 16
; CHECK-NEXT: ret i128 [[TMP4]]
;
entry:
diff --git a/llvm/test/Instrumentation/Instrumentor/numeric.ll b/llvm/test/Instrumentation/Instrumentor/numeric.ll
index b101b0c860bb0..fcb7b1d107e24 100644
--- a/llvm/test/Instrumentation/Instrumentor/numeric.ll
+++ b/llvm/test/Instrumentation/Instrumentor/numeric.ll
@@ -228,44 +228,44 @@ define i128 @test_i128(i128 %p1, i128 %p2) {
; CHECK-LABEL: define i128 @test_i128(
; CHECK-SAME: i128 [[P1:%.*]], i128 [[P2:%.*]]) {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[TMP0:%.*]] = alloca i128, align 8
-; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 8
-; CHECK-NEXT: [[TMP2:%.*]] = alloca i128, align 8
-; CHECK-NEXT: store i128 [[P1]], ptr [[TMP2]], align 4
-; CHECK-NEXT: store i128 [[P2]], ptr [[TMP1]], align 4
+; CHECK-NEXT: [[TMP0:%.*]] = alloca i128, align 16
+; CHECK-NEXT: [[TMP1:%.*]] = alloca i128, align 16
+; CHECK-NEXT: [[TMP2:%.*]] = alloca i128, align 16
+; CHECK-NEXT: store i128 [[P1]], ptr [[TMP2]], align 16
+; CHECK-NEXT: store i128 [[P2]], ptr [[TMP1]], align 16
; CHECK-NEXT: call void @__instrumentor_pre_numeric_ind(i32 12, i32 16, i32 14, ptr [[TMP2]], ptr [[TMP1]], i64 0, i32 26) #[[ATTR0]]
; CHECK-NEXT: [[A1:%.*]] = add i128 [[P1]], [[P2]]
-; CHECK-NEXT: store i128 [[A1]], ptr [[TMP0]], align 4
+; CHECK-NEXT: store i128 [[A1]], ptr [[TMP0]], align 16
; CHECK-NEXT: call void @__instrumentor_post_numeric_ind(i32 12, i32 16, i32 14, ptr [[TMP2]], ptr [[TMP1]], ptr [[TMP0]], i64 0, i32 -26) #[[ATTR0]]
-; CHECK-NEXT: store i128 [[P1]], ptr [[TMP2]], align 4
-; CHECK-NEXT: store i128 [[A1]], ptr [[TMP0]], align 4
+; CHECK-NEXT: store i128 [[P1]], ptr [[TMP2]], align 16
+; CHECK-NEXT: store i128 [[A1]], ptr [[TMP0]], align 16
; CHECK-NEXT: call void @__instrumentor_pre_numeric_ind(i32 12, i32 16, i32 18, ptr [[TMP2]], ptr [[TMP0]], i64 0, i32 27) #[[ATTR0]]
; CHECK-NEXT: [[A2:%.*]] = mul i128 [[P1]], [[A1]]
-; CHECK-NEXT: store i128 [[A2]], ptr [[TMP1]], align 4
+; CHECK-NEXT: store i128 [[A2]], ptr [[TMP1]], align 16
; CHECK-NEXT: call void @__instrumentor_post_numeric_ind(i32 12, i32 16, i32 18, ptr [[TMP2]], ptr [[TMP0]], ptr [[TMP1]], i64 0, i32 -27) #[[ATTR0]]
-; CHECK-NEXT: store i128 [[A2]], ptr [[TMP2]], align 4
-; CHECK-NEXT: store i128 [[P2]], ptr [[TMP1]], align 4
+; CHECK-NEXT: store i128 [[A2]], ptr [[TMP2]], align 16
+; CHECK-NEXT: store i128 [[P2]], ptr [[TMP1]], align 16
; CHECK-NEXT: call void @__instrumentor_pre_numeric_ind(i32 12, i32 16, i32 24, ptr [[TMP2]], ptr [[TMP1]], i64 0, i32 28) #[[ATTR0]]
; CHECK-NEXT: [[A3:%.*]] = srem i128 [[A2]], [[P2]]
-; CHECK-NEXT: store i128 [[A3]], ptr [[TMP0]], align 4
+; CHECK-NEXT: store i128 [[A3]], ptr [[TMP0]], align 16
; CHECK-NEXT: call void @__instrumentor_post_numeric_ind(i32 12, i32 16, i32 24, ptr [[TMP2]], ptr [[TMP1]], ptr [[TMP0]], i64 0, i32 -28) #[[ATTR0]]
-; CHECK-NEXT: store i128 [[A3]], ptr [[TMP2]], align 4
-; CHECK-NEXT: store i128 [[A1]], ptr [[TMP0]], align 4
+; CHECK-NEXT: store i128 [[A3]], ptr [[TMP2]], align 16
+; CHECK-NEXT: store i128 [[A1]], ptr [[TMP0]], align 16
; CHECK-NEXT: call void @__instrumentor_pre_numeric_ind(i32 12, i32 16, i32 16, ptr [[TMP2]], ptr [[TMP0]], i64 1, i32 29) #[[ATTR0]]
; CHECK-NEXT: [[A4:%.*]] = sub nsw i128 [[A3]], [[A1]]
-; CHECK-NEXT: store i128 [[A4]], ptr [[TMP1]], align 4
+; CHECK-NEXT: store i128 [[A4]], ptr [[TMP1]], align 16
; CHECK-NEXT: call void @__instrumentor_post_numeric_ind(i32 12, i32 16, i32 16, ptr [[TMP2]], ptr [[TMP0]], ptr [[TMP1]], i64 1, i32 -29) #[[ATTR0]]
-; CHECK-NEXT: store i128 [[A4]], ptr [[TMP2]], align 4
-; CHECK-NEXT: store i128 [[A2]], ptr [[TMP1]], align 4
+; CHECK-NEXT: store i128 [[A4]], ptr [[TMP2]], align 16
+; CHECK-NEXT: store i128 [[A2]], ptr [[TMP1]], align 16
; CHECK-NEXT: call void @__instrumentor_pre_numeric_ind(i32 12, i32 16, i32 30, ptr [[TMP2]], ptr [[TMP1]], i64 32, i32 30) #[[ATTR0]]
; CHECK-NEXT: [[A5:%.*]] = or disjoint i128 [[A4]], [[A2]]
-; CHECK-NEXT: store i128 [[A5]], ptr [[TMP0]], align 4
+; CHECK-NEXT: store i128 [[A5]], ptr [[TMP0]], align 16
; CHECK-NEXT: call void @__instrumentor_post_numeric_ind(i32 12, i32 16, i32 30, ptr [[TMP2]], ptr [[TMP1]], ptr [[TMP0]], i64 32, i32 -30) #[[ATTR0]]
-; CHECK-NEXT: store i128 [[A5]], ptr [[TMP2]], align 4
-; CHECK-NEXT: store i128 [[A3]], ptr [[TMP0]], align 4
+; CHECK-NEXT: store i128 [[A5]], ptr [[TMP2]], align 16
+; CHECK-NEXT: store i128 [[A3]], ptr [[TMP0]], align 16
; CHECK-NEXT: call void @__instrumentor_pre_numeric_ind(i32 12, i32 16, i32 26, ptr [[TMP2]], ptr [[TMP0]], i64 0, i32 31) #[[ATTR0]]
; CHECK-NEXT: [[A6:%.*]] = shl i128 [[A5]], [[A3]]
-; CHECK-NEXT: store i128 [[A6]], ptr [[TMP1]], align 4
+; CHECK-NEXT: store i128 [[A6]], ptr [[TMP1]], align 16
; CHECK-NEXT: call void @__instrumentor_post_numeric_ind(i32 12, i32 16, i32 26, ptr [[TMP2]], ptr [[TMP0]], ptr [[TMP1]], i64 0, i32 -31) #[[ATTR0]]
; CHECK-NEXT: ret i128 [[A6]]
;
diff --git a/llvm/test/Instrumentation/MemorySanitizer/PowerPC32/kernel-ppcle.ll b/llvm/test/Instrumentation/MemorySanitizer/PowerPC32/kernel-ppcle.ll
index 8ba033061defe..5722e382d3c22 100644
--- a/llvm/test/Instrumentation/MemorySanitizer/PowerPC32/kernel-ppcle.ll
+++ b/llvm/test/Instrumentation/MemorySanitizer/PowerPC32/kernel-ppcle.ll
@@ -229,12 +229,12 @@ define void @Store16(ptr %p, i128 %x) sanitize_memory {
; CHECK-NEXT: [[TMP15:%.*]] = call { ptr, ptr } @__msan_metadata_ptr_for_store_n(ptr [[P]], i32 16)
; CHECK-NEXT: [[TMP16:%.*]] = extractvalue { ptr, ptr } [[TMP15]], 0
; CHECK-NEXT: [[TMP17:%.*]] = extractvalue { ptr, ptr } [[TMP15]], 1
-; CHECK-NEXT: store i128 [[TMP9]], ptr [[TMP16]], align 8
+; CHECK-NEXT: store i128 [[TMP9]], ptr [[TMP16]], align 16
; CHECK-NEXT: [[_MSCMP3:%.*]] = icmp ne i128 [[TMP9]], 0
; CHECK-NEXT: br i1 [[_MSCMP3]], label %[[BB11:.*]], label %[[BB16:.*]], !prof [[PROF1]]
; CHECK: [[BB11]]:
; CHECK-NEXT: [[TMP19:%.*]] = call i32 @__msan_chain_origin(i32 [[TMP12]])
-; CHECK-NEXT: store i32 [[TMP19]], ptr [[TMP17]], align 8
+; CHECK-NEXT: store i32 [[TMP19]], ptr [[TMP17]], align 16
; CHECK-NEXT: [[TMP22:%.*]] = getelementptr i32, ptr [[TMP17]], i32 1
; CHECK-NEXT: store i32 [[TMP19]], ptr [[TMP22]], align 4
; CHECK-NEXT: [[TMP20:%.*]] = getelementptr i32, ptr [[TMP17]], i32 2
@@ -243,7 +243,7 @@ define void @Store16(ptr %p, i128 %x) sanitize_memory {
; CHECK-NEXT: store i32 [[TMP19]], ptr [[TMP21]], align 4
; CHECK-NEXT: br label %[[BB16]]
; CHECK: [[BB16]]:
-; CHECK-NEXT: store i128 [[X]], ptr [[P]], align 8
+; CHECK-NEXT: store i128 [[X]], ptr [[P]], align 16
; CHECK-NEXT: ret void
;
entry:
@@ -436,12 +436,12 @@ define i128 @Load16(ptr %p) sanitize_memory {
; CHECK-NEXT: call void @__msan_warning(i32 [[TMP3]]) #[[ATTR2]]
; CHECK-NEXT: br label %[[BB5]]
; CHECK: [[BB5]]:
-; CHECK-NEXT: [[TMP9:%.*]] = load i128, ptr [[P]], align 8
+; CHECK-NEXT: [[TMP9:%.*]] = load i128, ptr [[P]], align 16
; CHECK-NEXT: [[TMP10:%.*]] = call { ptr, ptr } @__msan_metadata_ptr_for_load_n(ptr [[P]], i32 16)
; CHECK-NEXT: [[TMP11:%.*]] = extractvalue { ptr, ptr } [[TMP10]], 0
; CHECK-NEXT: [[TMP12:%.*]] = extractvalue { ptr, ptr } [[TMP10]], 1
-; CHECK-NEXT: [[_MSLD:%.*]] = load i128, ptr [[TMP11]], align 8
-; CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[TMP12]], align 8
+; CHECK-NEXT: [[_MSLD:%.*]] = load i128, ptr [[TMP11]], align 16
+; CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[TMP12]], align 16
; CHECK-NEXT: store i128 [[_MSLD]], ptr [[RETVAL_SHADOW]], align 8
; CHECK-NEXT: store i32 [[TMP13]], ptr [[RETVAL_ORIGIN]], align 4
; CHECK-NEXT: ret i128 [[TMP9]]
diff --git a/llvm/test/Instrumentation/MemorySanitizer/byval.ll b/llvm/test/Instrumentation/MemorySanitizer/byval.ll
index 9f6a7cb189547..1ddb3a6ef6844 100644
--- a/llvm/test/Instrumentation/MemorySanitizer/byval.ll
+++ b/llvm/test/Instrumentation/MemorySanitizer/byval.ll
@@ -19,14 +19,14 @@ define i128 @ByValArgument(i32, ptr byval(i128) %p) sanitize_memory {
; CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP3]], ptr align 8 getelementptr (i8, ptr @__msan_param_tls, i64 8), i64 16, i1 false)
; CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[TMP5]], ptr align 4 getelementptr (i8, ptr @__msan_param_origin_tls, i64 8), i64 16, i1 false)
; CHECK-NEXT: call void @llvm.donothing()
-; CHECK-NEXT: [[X:%.*]] = load i128, ptr [[P]], align 8
+; CHECK-NEXT: [[X:%.*]] = load i128, ptr [[P]], align 16
; CHECK-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[P]] to i64
; CHECK-NEXT: [[TMP7:%.*]] = xor i64 [[TMP6]], 87960930222080
; CHECK-NEXT: [[TMP8:%.*]] = inttoptr i64 [[TMP7]] to ptr
; CHECK-NEXT: [[TMP9:%.*]] = add i64 [[TMP7]], 17592186044416
; CHECK-NEXT: [[TMP10:%.*]] = inttoptr i64 [[TMP9]] to ptr
-; CHECK-NEXT: [[_MSLD:%.*]] = load i128, ptr [[TMP8]], align 8
-; CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 8
+; CHECK-NEXT: [[_MSLD:%.*]] = load i128, ptr [[TMP8]], align 16
+; CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 16
; CHECK-NEXT: store i128 [[_MSLD]], ptr @__msan_retval_tls, align 8
; CHECK-NEXT: store i32 [[TMP11]], ptr @__msan_retval_origin_tls, align 4
; CHECK-NEXT: ret i128 [[X]]
@@ -45,9 +45,9 @@ define i128 @ByValArgumentNoSanitize(i32, ptr byval(i128) %p) {
; CHECK-NEXT: [[TMP3:%.*]] = inttoptr i64 [[TMP2]] to ptr
; CHECK-NEXT: [[TMP4:%.*]] = add i64 [[TMP2]], 17592186044416
; CHECK-NEXT: [[TMP5:%.*]] = inttoptr i64 [[TMP4]] to ptr
-; CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 8 [[TMP3]], i8 0, i64 16, i1 false)
+; CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 16 [[TMP3]], i8 0, i64 16, i1 false)
; CHECK-NEXT: call void @llvm.donothing()
-; CHECK-NEXT: [[X:%.*]] = load i128, ptr [[P]], align 8
+; CHECK-NEXT: [[X:%.*]] = load i128, ptr [[P]], align 16
; CHECK-NEXT: store i128 0, ptr @__msan_retval_tls, align 8
; CHECK-NEXT: store i32 0, ptr @__msan_retval_origin_tls, align 4
; CHECK-NEXT: ret i128 [[X]]
@@ -87,7 +87,7 @@ define void @ByValForwardNoSanitize(i32, ptr byval(i128) %p) {
; CHECK-NEXT: [[TMP3:%.*]] = inttoptr i64 [[TMP2]] to ptr
; CHECK-NEXT: [[TMP4:%.*]] = add i64 [[TMP2]], 17592186044416
; CHECK-NEXT: [[TMP5:%.*]] = inttoptr i64 [[TMP4]] to ptr
-; CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 8 [[TMP3]], i8 0, i64 16, i1 false)
+; CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 16 [[TMP3]], i8 0, i64 16, i1 false)
; CHECK-NEXT: call void @llvm.donothing()
; CHECK-NEXT: store i64 0, ptr @__msan_param_tls, align 8
; CHECK-NEXT: call void @Fn(ptr [[P]])
@@ -135,7 +135,7 @@ define void @ByValForwardByValNoSanitize(i32, ptr byval(i128) %p) {
; CHECK-NEXT: [[TMP3:%.*]] = inttoptr i64 [[TMP2]] to ptr
; CHECK-NEXT: [[TMP4:%.*]] = add i64 [[TMP2]], 17592186044416
; CHECK-NEXT: [[TMP5:%.*]] = inttoptr i64 [[TMP4]] to ptr
-; CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 8 [[TMP3]], i8 0, i64 16, i1 false)
+; CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 16 [[TMP3]], i8 0, i64 16, i1 false)
; CHECK-NEXT: call void @llvm.donothing()
; CHECK-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[P]] to i64
; CHECK-NEXT: [[TMP7:%.*]] = xor i64 [[TMP6]], 87960930222080
diff --git a/llvm/test/Instrumentation/MemorySanitizer/msan_kernel_basic.ll b/llvm/test/Instrumentation/MemorySanitizer/msan_kernel_basic.ll
index 5d63367919d1a..e0526d56d9c2d 100644
--- a/llvm/test/Instrumentation/MemorySanitizer/msan_kernel_basic.ll
+++ b/llvm/test/Instrumentation/MemorySanitizer/msan_kernel_basic.ll
@@ -243,7 +243,7 @@ define void @Store16(ptr nocapture %p, i128 %x) nounwind uwtable sanitize_memory
; CHECK-NEXT: [[TMP13:%.*]] = call { ptr, ptr } @__msan_metadata_ptr_for_store_n(ptr [[P]], i64 16)
; CHECK-NEXT: [[TMP14:%.*]] = extractvalue { ptr, ptr } [[TMP13]], 0
; CHECK-NEXT: [[TMP15:%.*]] = extractvalue { ptr, ptr } [[TMP13]], 1
-; CHECK-NEXT: store i128 [[TMP7]], ptr [[TMP14]], align 8
+; CHECK-NEXT: store i128 [[TMP7]], ptr [[TMP14]], align 16
; CHECK-NEXT: [[_MSCMP3:%.*]] = icmp ne i128 [[TMP7]], 0
; CHECK-NEXT: br i1 [[_MSCMP3]], label %[[BB10:.*]], label %[[BB16:.*]], !prof [[PROF1]]
; CHECK: [[BB10]]:
@@ -251,12 +251,12 @@ define void @Store16(ptr nocapture %p, i128 %x) nounwind uwtable sanitize_memory
; CHECK-NEXT: [[TMP18:%.*]] = zext i32 [[TMP17]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = shl i64 [[TMP18]], 32
; CHECK-NEXT: [[TMP20:%.*]] = or i64 [[TMP18]], [[TMP19]]
-; CHECK-NEXT: store i64 [[TMP20]], ptr [[TMP15]], align 8
+; CHECK-NEXT: store i64 [[TMP20]], ptr [[TMP15]], align 16
; CHECK-NEXT: [[TMP21:%.*]] = getelementptr i64, ptr [[TMP15]], i32 1
; CHECK-NEXT: store i64 [[TMP20]], ptr [[TMP21]], align 8
; CHECK-NEXT: br label %[[BB16]]
; CHECK: [[BB16]]:
-; CHECK-NEXT: store i128 [[X]], ptr [[P]], align 8
+; CHECK-NEXT: store i128 [[X]], ptr [[P]], align 16
; CHECK-NEXT: ret void
;
entry:
@@ -441,12 +441,12 @@ define i128 @Load16(ptr nocapture %p) nounwind uwtable sanitize_memory {
; CHECK-NEXT: call void @__msan_warning(i32 [[TMP4]]) #[[ATTR8]]
; CHECK-NEXT: br label %[[BB4]]
; CHECK: [[BB4]]:
-; CHECK-NEXT: [[TMP7:%.*]] = load i128, ptr [[P]], align 8
+; CHECK-NEXT: [[TMP7:%.*]] = load i128, ptr [[P]], align 16
; CHECK-NEXT: [[TMP8:%.*]] = call { ptr, ptr } @__msan_metadata_ptr_for_load_n(ptr [[P]], i64 16)
; CHECK-NEXT: [[TMP9:%.*]] = extractvalue { ptr, ptr } [[TMP8]], 0
; CHECK-NEXT: [[TMP10:%.*]] = extractvalue { ptr, ptr } [[TMP8]], 1
-; CHECK-NEXT: [[_MSLD:%.*]] = load i128, ptr [[TMP9]], align 8
-; CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 8
+; CHECK-NEXT: [[_MSLD:%.*]] = load i128, ptr [[TMP9]], align 16
+; CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 16
; CHECK-NEXT: store i128 [[_MSLD]], ptr [[RETVAL_SHADOW]], align 8
; CHECK-NEXT: store i32 [[TMP11]], ptr [[RETVAL_ORIGIN]], align 4
; CHECK-NEXT: ret i128 [[TMP7]]
@@ -490,7 +490,7 @@ define dso_local i32 @VarArgFn(ptr %fmt, ...) local_unnamed_addr sanitize_memory
; CHECK-NEXT: [[TMP9:%.*]] = alloca i8, i64 [[TMP6]], align 8
; CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP9]], ptr align 8 [[VA_ARG_ORIGIN]], i64 [[TMP8]], i1 false)
; CHECK-NEXT: call void @llvm.donothing()
-; CHECK-NEXT: [[ARGS:%.*]] = alloca [1 x %struct.__va_list_tag], align 16
+; CHECK-NEXT: [[ARGS:%.*]] = alloca [1 x [[STRUCT___VA_LIST_TAG:%.*]]], align 16
; CHECK-NEXT: call void @__msan_poison_alloca(ptr [[ARGS]], i64 24, ptr @[[GLOB0:[0-9]+]])
; CHECK-NEXT: [[TMP10:%.*]] = call { ptr, ptr } @__msan_metadata_ptr_for_store_1(ptr [[ARGS]])
; CHECK-NEXT: [[TMP11:%.*]] = extractvalue { ptr, ptr } [[TMP10]], 0
diff --git a/llvm/test/Transforms/HotColdSplit/lifetime-markers-on-inputs-2.ll b/llvm/test/Transforms/HotColdSplit/lifetime-markers-on-inputs-2.ll
index da7a9b8d7531b..c5da2ff2eb47a 100644
--- a/llvm/test/Transforms/HotColdSplit/lifetime-markers-on-inputs-2.ll
+++ b/llvm/test/Transforms/HotColdSplit/lifetime-markers-on-inputs-2.ll
@@ -37,7 +37,7 @@ declare void @use(ptr)
define void @only_lifetime_start_is_cold(i1 %arg) {
; CHECK-LABEL: @only_lifetime_start_is_cold(
; CHECK-NEXT: entry:
-; CHECK-NEXT: [[LOCAL1:%.*]] = alloca i256, align 8
+; CHECK-NEXT: [[LOCAL1:%.*]] = alloca i256, align 16
; CHECK-NEXT: br i1 [[ARG:%.*]], label [[CODEREPL:%.*]], label [[NO_EXTRACT1:%.*]]
; CHECK: codeRepl:
; CHECK-NEXT: call void @llvm.lifetime.start.p0(ptr [[LOCAL1]])
@@ -95,7 +95,7 @@ exit:
define void @only_lifetime_end_is_cold(i1 %arg) {
; CHECK-LABEL: @only_lifetime_end_is_cold(
; CHECK-NEXT: entry:
-; CHECK-NEXT: [[LOCAL1:%.*]] = alloca i256, align 8
+; CHECK-NEXT: [[LOCAL1:%.*]] = alloca i256, align 16
; CHECK-NEXT: call void @llvm.lifetime.start.p0(ptr [[LOCAL1]])
; CHECK-NEXT: br i1 [[ARG:%.*]], label [[NO_EXTRACT1:%.*]], label [[CODEREPL:%.*]]
; CHECK: no-extract1:
@@ -133,7 +133,7 @@ exit:
define void @do_not_lift_lifetime_end(i1 %arg) {
; CHECK-LABEL: @do_not_lift_lifetime_end(
; CHECK-NEXT: entry:
-; CHECK-NEXT: [[LOCAL1:%.*]] = alloca i256, align 8
+; CHECK-NEXT: [[LOCAL1:%.*]] = alloca i256, align 16
; CHECK-NEXT: call void @llvm.lifetime.start.p0(ptr [[LOCAL1]])
; CHECK-NEXT: br label [[HEADER:%.*]]
; CHECK: header:
diff --git a/llvm/test/Transforms/InferAlignment/irregular-size.ll b/llvm/test/Transforms/InferAlignment/irregular-size.ll
index 9413c8ac5be46..93ff9380ce646 100644
--- a/llvm/test/Transforms/InferAlignment/irregular-size.ll
+++ b/llvm/test/Transforms/InferAlignment/irregular-size.ll
@@ -4,9 +4,9 @@
define void @non_pow2_size(i177 %X) {
; CHECK-LABEL: define void @non_pow2_size
; CHECK-SAME: (i177 [[X:%.*]]) {
-; CHECK-NEXT: [[A:%.*]] = alloca i177, align 8
-; CHECK-NEXT: [[L1:%.*]] = load i177, ptr [[A]], align 8
-; CHECK-NEXT: store i177 [[X]], ptr [[A]], align 8
+; CHECK-NEXT: [[A:%.*]] = alloca i177, align 16
+; CHECK-NEXT: [[L1:%.*]] = load i177, ptr [[A]], align 16
+; CHECK-NEXT: store i177 [[X]], ptr [[A]], align 16
; CHECK-NEXT: ret void
;
%A = alloca i177, align 1
diff --git a/llvm/test/Transforms/InstCombine/apint-and.ll b/llvm/test/Transforms/InstCombine/apint-and.ll
index c5f2087ee02fd..5ef0925c3409a 100644
--- a/llvm/test/Transforms/InstCombine/apint-and.ll
+++ b/llvm/test/Transforms/InstCombine/apint-and.ll
@@ -105,7 +105,7 @@ define i117 @test12(i117 %A, ptr %P) {
; CHECK-LABEL: @test12(
; CHECK-NEXT: [[TMP1:%.*]] = and i117 [[A:%.*]], -4
; CHECK-NEXT: [[C:%.*]] = xor i117 [[TMP1]], 15
-; CHECK-NEXT: store i117 [[C]], ptr [[P:%.*]], align 4
+; CHECK-NEXT: store i117 [[C]], ptr [[P:%.*]], align 16
; CHECK-NEXT: ret i117 3
;
%B = or i117 %A, 3
diff --git a/llvm/test/Transforms/InstCombine/load-store-forward.ll b/llvm/test/Transforms/InstCombine/load-store-forward.ll
index 6a0897ff75036..2e18334416455 100644
--- a/llvm/test/Transforms/InstCombine/load-store-forward.ll
+++ b/llvm/test/Transforms/InstCombine/load-store-forward.ll
@@ -458,7 +458,7 @@ define i32 @load_after_memset_0_clobber(ptr %a) {
define i256 @load_after_memset_0_too_small(ptr %a) {
; CHECK-LABEL: @load_after_memset_0_too_small(
; CHECK-NEXT: call void @llvm.memset.p0.i64(ptr noundef nonnull align 1 dereferenceable(16) [[A:%.*]], i8 0, i64 16, i1 false)
-; CHECK-NEXT: [[V:%.*]] = load i256, ptr [[A]], align 4
+; CHECK-NEXT: [[V:%.*]] = load i256, ptr [[A]], align 16
; CHECK-NEXT: ret i256 [[V]]
;
call void @llvm.memset.p0.i64(ptr %a, i8 0, i64 16, i1 false)
@@ -469,7 +469,7 @@ define i256 @load_after_memset_0_too_small(ptr %a) {
define i129 @load_after_memset_0_too_small_by_one_bit(ptr %a) {
; CHECK-LABEL: @load_after_memset_0_too_small_by_one_bit(
; CHECK-NEXT: call void @llvm.memset.p0.i64(ptr noundef nonnull align 1 dereferenceable(16) [[A:%.*]], i8 0, i64 16, i1 false)
-; CHECK-NEXT: [[V:%.*]] = load i129, ptr [[A]], align 4
+; CHECK-NEXT: [[V:%.*]] = load i129, ptr [[A]], align 16
; CHECK-NEXT: ret i129 [[V]]
;
call void @llvm.memset.p0.i64(ptr %a, i8 0, i64 16, i1 false)
diff --git a/llvm/test/Transforms/InstCombine/or-xor.ll b/llvm/test/Transforms/InstCombine/or-xor.ll
index b05ff15b8b3c8..8920711e2ca09 100644
--- a/llvm/test/Transforms/InstCombine/or-xor.ll
+++ b/llvm/test/Transforms/InstCombine/or-xor.ll
@@ -125,7 +125,7 @@ define i8 @test5_extra_use_not(i8 %x, i8 %y, ptr %dst) {
define i65 @test5_extra_use_xor(i65 %x, i65 %y, ptr %dst) {
; CHECK-LABEL: @test5_extra_use_xor(
; CHECK-NEXT: [[XOR:%.*]] = xor i65 [[X:%.*]], [[Y:%.*]]
-; CHECK-NEXT: store i65 [[XOR]], ptr [[DST:%.*]], align 4
+; CHECK-NEXT: store i65 [[XOR]], ptr [[DST:%.*]], align 16
; CHECK-NEXT: [[TMP1:%.*]] = and i65 [[X]], [[Y]]
; CHECK-NEXT: [[Z:%.*]] = xor i65 [[TMP1]], -1
; CHECK-NEXT: ret i65 [[Z]]
diff --git a/llvm/test/Transforms/InstCombine/select.ll b/llvm/test/Transforms/InstCombine/select.ll
index a1e84fad9a827..aaedf9d0e43fa 100644
--- a/llvm/test/Transforms/InstCombine/select.ll
+++ b/llvm/test/Transforms/InstCombine/select.ll
@@ -1348,11 +1348,11 @@ define ptr @test85(i1 %flag) {
; CHECK-LABEL: define ptr @test85(
; CHECK-SAME: i1 [[FLAG:%.*]]) {
; CHECK-NEXT: [[X:%.*]] = alloca [2 x ptr], align 8
-; CHECK-NEXT: [[Y:%.*]] = alloca i128, align 8
+; CHECK-NEXT: [[Y:%.*]] = alloca i128, align 16
; CHECK-NEXT: call void @scribble_on_i128(ptr nonnull [[X]])
; CHECK-NEXT: call void @scribble_on_i128(ptr nonnull [[Y]])
-; CHECK-NEXT: [[T:%.*]] = load i128, ptr [[X]], align 4
-; CHECK-NEXT: store i128 [[T]], ptr [[Y]], align 4
+; CHECK-NEXT: [[T:%.*]] = load i128, ptr [[X]], align 16
+; CHECK-NEXT: store i128 [[T]], ptr [[Y]], align 16
; CHECK-NEXT: [[X_VAL:%.*]] = load ptr, ptr [[X]], align 8
; CHECK-NEXT: [[Y_VAL:%.*]] = load ptr, ptr [[Y]], align 8
; CHECK-NEXT: [[V:%.*]] = select i1 [[FLAG]], ptr [[X_VAL]], ptr [[Y_VAL]]
@@ -1376,14 +1376,13 @@ define i128 @test86(i1 %flag) {
; CHECK-LABEL: define i128 @test86(
; CHECK-SAME: i1 [[FLAG:%.*]]) {
; CHECK-NEXT: [[X:%.*]] = alloca [2 x ptr], align 8
-; CHECK-NEXT: [[Y:%.*]] = alloca i128, align 8
+; CHECK-NEXT: [[Y:%.*]] = alloca i128, align 16
; CHECK-NEXT: call void @scribble_on_i128(ptr nonnull [[X]])
; CHECK-NEXT: call void @scribble_on_i128(ptr nonnull [[Y]])
; CHECK-NEXT: [[T:%.*]] = load ptr, ptr [[X]], align 8
; CHECK-NEXT: store ptr [[T]], ptr [[Y]], align 8
-; CHECK-NEXT: [[X_VAL:%.*]] = load i128, ptr [[X]], align 4
-; CHECK-NEXT: [[Y_VAL:%.*]] = load i128, ptr [[Y]], align 4
-; CHECK-NEXT: [[V:%.*]] = select i1 [[FLAG]], i128 [[X_VAL]], i128 [[Y_VAL]]
+; CHECK-NEXT: [[P:%.*]] = select i1 [[FLAG]], ptr [[X]], ptr [[Y]]
+; CHECK-NEXT: [[V:%.*]] = load i128, ptr [[P]], align 16
; CHECK-NEXT: ret i128 [[V]]
;
%x = alloca [2 x ptr]
diff --git a/llvm/test/Transforms/InstCombine/shift.ll b/llvm/test/Transforms/InstCombine/shift.ll
index d63ba48830068..701ee07b36d84 100644
--- a/llvm/test/Transforms/InstCombine/shift.ll
+++ b/llvm/test/Transforms/InstCombine/shift.ll
@@ -1740,19 +1740,19 @@ define i177 @lshr_out_of_range2(i177 %Y, ptr %A2, ptr %ptr) {
; https://bugs.chromium.org/p/oss-fuzz/issues/detail?id=5032
define void @ashr_out_of_range(ptr %A) {
; CHECK-LABEL: @ashr_out_of_range(
-; CHECK-NEXT: [[L:%.*]] = load i177, ptr [[A:%.*]], align 4
+; CHECK-NEXT: [[L:%.*]] = load i177, ptr [[A:%.*]], align 16
; CHECK-NEXT: [[TMP1:%.*]] = icmp eq i177 [[L]], -1
; CHECK-NEXT: [[TMP2:%.*]] = select i1 [[TMP1]], i64 -1, i64 -2
-; CHECK-NEXT: [[G11:%.*]] = getelementptr [24 x i8], ptr [[A]], i64 [[TMP2]]
-; CHECK-NEXT: [[L7:%.*]] = load i177, ptr [[G11]], align 4
+; CHECK-NEXT: [[G11:%.*]] = getelementptr [32 x i8], ptr [[A]], i64 [[TMP2]]
+; CHECK-NEXT: [[L7:%.*]] = load i177, ptr [[G11]], align 16
; CHECK-NEXT: [[L7_FROZEN:%.*]] = freeze i177 [[L7]]
; CHECK-NEXT: [[C171:%.*]] = icmp slt i177 [[L7_FROZEN]], 0
; CHECK-NEXT: [[C17:%.*]] = and i1 [[TMP1]], [[C171]]
; CHECK-NEXT: [[TMP3:%.*]] = sext i1 [[C17]] to i64
-; CHECK-NEXT: [[G62:%.*]] = getelementptr [24 x i8], ptr [[G11]], i64 [[TMP3]]
+; CHECK-NEXT: [[G62:%.*]] = getelementptr [32 x i8], ptr [[G11]], i64 [[TMP3]]
; CHECK-NEXT: [[TMP4:%.*]] = icmp eq i177 [[L7_FROZEN]], -1
; CHECK-NEXT: [[B28:%.*]] = select i1 [[TMP4]], i177 0, i177 [[L7_FROZEN]]
-; CHECK-NEXT: store i177 [[B28]], ptr [[G62]], align 4
+; CHECK-NEXT: store i177 [[B28]], ptr [[G62]], align 16
; CHECK-NEXT: ret void
;
%L = load i177, ptr %A
@@ -1780,10 +1780,10 @@ define void @ashr_out_of_range_1(ptr %A) {
; CHECK-NEXT: [[TMP1:%.*]] = icmp eq i177 [[L_FROZEN]], -1
; CHECK-NEXT: [[TMP2:%.*]] = trunc i177 [[L_FROZEN]] to i64
; CHECK-NEXT: [[TMP3:%.*]] = select i1 [[TMP1]], i64 0, i64 [[TMP2]]
-; CHECK-NEXT: [[TMP4:%.*]] = getelementptr [24 x i8], ptr [[A]], i64 [[TMP3]]
-; CHECK-NEXT: [[G11:%.*]] = getelementptr i8, ptr [[TMP4]], i64 -24
+; CHECK-NEXT: [[TMP4:%.*]] = getelementptr [32 x i8], ptr [[A]], i64 [[TMP3]]
+; CHECK-NEXT: [[G11:%.*]] = getelementptr i8, ptr [[TMP4]], i64 -32
; CHECK-NEXT: [[TMP5:%.*]] = sext i1 [[TMP1]] to i64
-; CHECK-NEXT: [[G62:%.*]] = getelementptr [24 x i8], ptr [[G11]], i64 [[TMP5]]
+; CHECK-NEXT: [[G62:%.*]] = getelementptr [32 x i8], ptr [[G11]], i64 [[TMP5]]
; CHECK-NEXT: [[TMP6:%.*]] = icmp eq i177 [[L_FROZEN]], -1
; CHECK-NEXT: [[B28:%.*]] = select i1 [[TMP6]], i177 0, i177 [[L_FROZEN]]
; CHECK-NEXT: store i177 [[B28]], ptr [[G62]], align 4
diff --git a/llvm/test/Transforms/InstCombine/xor.ll b/llvm/test/Transforms/InstCombine/xor.ll
index 3abaf74285cc0..b2b1c4787f657 100644
--- a/llvm/test/Transforms/InstCombine/xor.ll
+++ b/llvm/test/Transforms/InstCombine/xor.ll
@@ -1304,7 +1304,7 @@ define i67 @xor_orn_commute3_1use(i67 %pa, i67 %pb, ptr %s) {
; CHECK-NEXT: [[B:%.*]] = udiv i67 42, [[PB:%.*]]
; CHECK-NEXT: [[NOTA:%.*]] = xor i67 [[A]], -1
; CHECK-NEXT: [[L:%.*]] = or i67 [[B]], [[NOTA]]
-; CHECK-NEXT: store i67 [[L]], ptr [[S:%.*]], align 4
+; CHECK-NEXT: store i67 [[L]], ptr [[S:%.*]], align 16
; CHECK-NEXT: [[Z:%.*]] = xor i67 [[A]], [[L]]
; CHECK-NEXT: ret i67 [[Z]]
;
diff --git a/llvm/test/Transforms/LoopIdiom/X86/unordered-atomic-memcpy.ll b/llvm/test/Transforms/LoopIdiom/X86/unordered-atomic-memcpy.ll
index df12769b2527f..35b56ce18b95b 100644
--- a/llvm/test/Transforms/LoopIdiom/X86/unordered-atomic-memcpy.ll
+++ b/llvm/test/Transforms/LoopIdiom/X86/unordered-atomic-memcpy.ll
@@ -544,8 +544,8 @@ for.end: ; preds = %for.body, %entry
define void @test9(i64 %Size) nounwind ssp {
; CHECK-LABEL: @test9(
; CHECK-NEXT: bb.nph:
-; CHECK-NEXT: [[BASE:%.*]] = alloca i128, i32 10000, align 8
-; CHECK-NEXT: [[DEST:%.*]] = alloca i128, i32 10000, align 8
+; CHECK-NEXT: [[BASE:%.*]] = alloca i128, i32 10000, align 16
+; CHECK-NEXT: [[DEST:%.*]] = alloca i128, i32 10000, align 16
; CHECK-NEXT: [[TMP0:%.*]] = shl nuw i64 [[SIZE:%.*]], 4
; CHECK-NEXT: call void @llvm.memcpy.element.unordered.atomic.p0.p0.i64(ptr align 16 [[DEST]], ptr align 16 [[BASE]], i64 [[TMP0]], i32 16)
; CHECK-NEXT: br label [[FOR_BODY:%.*]]
@@ -583,8 +583,8 @@ for.end: ; preds = %for.body, %entry
define void @test10(i64 %Size) nounwind ssp {
; CHECK-LABEL: @test10(
; CHECK-NEXT: bb.nph:
-; CHECK-NEXT: [[BASE:%.*]] = alloca i256, i32 10000, align 8
-; CHECK-NEXT: [[DEST:%.*]] = alloca i256, i32 10000, align 8
+; CHECK-NEXT: [[BASE:%.*]] = alloca i256, i32 10000, align 16
+; CHECK-NEXT: [[DEST:%.*]] = alloca i256, i32 10000, align 16
; CHECK-NEXT: br label [[FOR_BODY:%.*]]
; CHECK: for.body:
; CHECK-NEXT: [[INDVAR:%.*]] = phi i64 [ 0, [[BB_NPH:%.*]] ], [ [[INDVAR_NEXT:%.*]], [[FOR_BODY]] ]
diff --git a/llvm/test/Transforms/LowerMatrixIntrinsics/unary.ll b/llvm/test/Transforms/LowerMatrixIntrinsics/unary.ll
index 5d7777245ebed..697420a2e999a 100644
--- a/llvm/test/Transforms/LowerMatrixIntrinsics/unary.ll
+++ b/llvm/test/Transforms/LowerMatrixIntrinsics/unary.ll
@@ -218,7 +218,7 @@ define void @bitcast_2x2_v4f64_to_v8i32(ptr %in, ptr %out) {
define void @bitcast_2x2_i256_to_v4i64(ptr %in, ptr %out) {
; CHECK-LABEL: @bitcast_2x2_i256_to_v4i64(
-; CHECK-NEXT: [[INV:%.*]] = load i256, ptr [[IN:%.*]], align 4
+; CHECK-NEXT: [[INV:%.*]] = load i256, ptr [[IN:%.*]], align 16
; CHECK-NEXT: [[OP:%.*]] = bitcast i256 [[INV]] to <4 x double>
; CHECK-NEXT: [[SPLIT:%.*]] = shufflevector <4 x double> [[OP]], <4 x double> poison, <2 x i32> <i32 0, i32 1>
; CHECK-NEXT: [[SPLIT1:%.*]] = shufflevector <4 x double> [[OP]], <4 x double> poison, <2 x i32> <i32 2, i32 3>
@@ -240,7 +240,7 @@ define void @bitcast_2x2_4i64_to_i256(ptr %in, ptr %out) {
; CHECK-NEXT: [[COL_LOAD1:%.*]] = load <2 x double>, ptr [[VEC_GEP]], align 8
; CHECK-NEXT: [[TMP1:%.*]] = shufflevector <2 x double> [[COL_LOAD]], <2 x double> [[COL_LOAD1]], <4 x i32> <i32 0, i32 1, i32 2, i32 3>
; CHECK-NEXT: [[OP:%.*]] = bitcast <4 x double> [[TMP1]] to i256
-; CHECK-NEXT: store i256 [[OP]], ptr [[OUT:%.*]], align 4
+; CHECK-NEXT: store i256 [[OP]], ptr [[OUT:%.*]], align 16
; CHECK-NEXT: ret void
;
%inv = call <4 x double> @llvm.matrix.column.major.load(ptr %in, i64 2, i1 false, i32 2, i32 2)
diff --git a/llvm/test/Transforms/SCCP/apint-bigint2.ll b/llvm/test/Transforms/SCCP/apint-bigint2.ll
index c5e7643cdd5ab..ea47a94f6d350 100644
--- a/llvm/test/Transforms/SCCP/apint-bigint2.ll
+++ b/llvm/test/Transforms/SCCP/apint-bigint2.ll
@@ -24,7 +24,7 @@ define i101 @large_aggregate() {
; CHECK-NEXT: [[D:%.*]] = and i101 undef, 1
; CHECK-NEXT: [[DD:%.*]] = or i101 [[D]], 1
; CHECK-NEXT: [[G:%.*]] = getelementptr i101, ptr getelementptr inbounds nuw (i8, ptr @Y, i64 80), i101 [[DD]]
-; CHECK-NEXT: [[L3:%.*]] = load i101, ptr [[G]], align 4
+; CHECK-NEXT: [[L3:%.*]] = load i101, ptr [[G]], align 16
; CHECK-NEXT: ret i101 [[L3]]
;
%B = load i101, ptr undef
@@ -41,7 +41,7 @@ define i101 @large_aggregate_2() {
; CHECK-NEXT: [[D:%.*]] = and i101 undef, 1
; CHECK-NEXT: [[DD:%.*]] = or i101 [[D]], 1
; CHECK-NEXT: [[G:%.*]] = getelementptr i101, ptr getelementptr inbounds nuw (i8, ptr @Y, i64 80), i101 [[DD]]
-; CHECK-NEXT: [[L3:%.*]] = load i101, ptr [[G]], align 4
+; CHECK-NEXT: [[L3:%.*]] = load i101, ptr [[G]], align 16
; CHECK-NEXT: ret i101 [[L3]]
;
%D = and i101 undef, 1
diff --git a/llvm/test/Transforms/SCCP/ub-shift.ll b/llvm/test/Transforms/SCCP/ub-shift.ll
index ae2739d0445d0..4810c16b5685c 100644
--- a/llvm/test/Transforms/SCCP/ub-shift.ll
+++ b/llvm/test/Transforms/SCCP/ub-shift.ll
@@ -23,10 +23,10 @@ define void @shift_undef_64(ptr %p) {
define void @shift_undef_65(ptr %p) {
; CHECK-LABEL: @shift_undef_65(
-; CHECK-NEXT: store i65 0, ptr [[P:%.*]], align 4
-; CHECK-NEXT: store i65 0, ptr [[P]], align 4
+; CHECK-NEXT: store i65 0, ptr [[P:%.*]], align 16
+; CHECK-NEXT: store i65 0, ptr [[P]], align 16
; CHECK-NEXT: [[R3:%.*]] = shl nuw nsw i65 1, -18446744073709551615
-; CHECK-NEXT: store i65 [[R3]], ptr [[P]], align 4
+; CHECK-NEXT: store i65 [[R3]], ptr [[P]], align 16
; CHECK-NEXT: ret void
;
%r1 = lshr i65 2, 18446744073709551617
@@ -43,10 +43,10 @@ define void @shift_undef_65(ptr %p) {
define void @shift_undef_256(ptr %p) {
; CHECK-LABEL: @shift_undef_256(
-; CHECK-NEXT: store i256 0, ptr [[P:%.*]], align 4
-; CHECK-NEXT: store i256 0, ptr [[P]], align 4
+; CHECK-NEXT: store i256 0, ptr [[P:%.*]], align 16
+; CHECK-NEXT: store i256 0, ptr [[P]], align 16
; CHECK-NEXT: [[R3:%.*]] = shl nuw nsw i256 1, 18446744073709551619
-; CHECK-NEXT: store i256 [[R3]], ptr [[P]], align 4
+; CHECK-NEXT: store i256 [[R3]], ptr [[P]], align 16
; CHECK-NEXT: ret void
;
%r1 = lshr i256 2, 18446744073709551617
@@ -63,10 +63,10 @@ define void @shift_undef_256(ptr %p) {
define void @shift_undef_511(ptr %p) {
; CHECK-LABEL: @shift_undef_511(
-; CHECK-NEXT: store i511 0, ptr [[P:%.*]], align 4
-; CHECK-NEXT: store i511 -1, ptr [[P]], align 4
+; CHECK-NEXT: store i511 0, ptr [[P:%.*]], align 16
+; CHECK-NEXT: store i511 -1, ptr [[P]], align 16
; CHECK-NEXT: [[R3:%.*]] = shl nuw nsw i511 -3, 1208925819614629174706180
-; CHECK-NEXT: store i511 [[R3]], ptr [[P]], align 4
+; CHECK-NEXT: store i511 [[R3]], ptr [[P]], align 16
; CHECK-NEXT: ret void
;
%r1 = lshr i511 -1, 1208925819614629174706276 ; 2^80 + 100
diff --git a/llvm/test/Transforms/SimplifyCFG/jump-threading-max-jump-threading-live-blocks.ll b/llvm/test/Transforms/SimplifyCFG/jump-threading-max-jump-threading-live-blocks.ll
index 686869348019d..92bc3b54015e7 100644
--- a/llvm/test/Transforms/SimplifyCFG/jump-threading-max-jump-threading-live-blocks.ll
+++ b/llvm/test/Transforms/SimplifyCFG/jump-threading-max-jump-threading-live-blocks.ll
@@ -21,11 +21,11 @@ define void @testB(ptr %ptrA, ptr %ptrB, i64 %a, i64 %b, i64 %c) {
; CHECK_LIMIT_3-NEXT: br i1 [[COND2]], label %[[IFB_ARM1:.*]], label %[[IFB_ARM2:.*]]
; CHECK_LIMIT_3: [[IFB_ARM1]]:
; CHECK_LIMIT_3-NEXT: [[PTR_ARM1:%.*]] = getelementptr i64, ptr [[PTRB]], i64 8
-; CHECK_LIMIT_3-NEXT: store i128 0, ptr [[PTR_ARM1]], align 4
+; CHECK_LIMIT_3-NEXT: store i128 0, ptr [[PTR_ARM1]], align 16
; CHECK_LIMIT_3-NEXT: br label %[[IFB_JOIN:.*]]
; CHECK_LIMIT_3: [[IFB_ARM2]]:
; CHECK_LIMIT_3-NEXT: [[PTR_ARM2:%.*]] = getelementptr i64, ptr [[PTRB]], i64 16
-; CHECK_LIMIT_3-NEXT: store i128 0, ptr [[PTR_ARM2]], align 4
+; CHECK_LIMIT_3-NEXT: store i128 0, ptr [[PTR_ARM2]], align 16
; CHECK_LIMIT_3-NEXT: br label %[[IFB_JOIN]]
; CHECK_LIMIT_3: [[IFB_JOIN]]:
; CHECK_LIMIT_3-NEXT: [[PTRC:%.*]] = phi ptr [ [[PTR_ARM1]], %[[IFB_ARM1]] ], [ [[PTR_ARM2]], %[[IFB_ARM2]] ]
@@ -45,11 +45,11 @@ define void @testB(ptr %ptrA, ptr %ptrB, i64 %a, i64 %b, i64 %c) {
; CHECK_LIMIT_4-NEXT: br i1 [[COND2]], label %[[IFB_ARM1:.*]], label %[[IFB_ARM2:.*]]
; CHECK_LIMIT_4: [[IFB_ARM1]]:
; CHECK_LIMIT_4-NEXT: [[PTR_ARM1:%.*]] = getelementptr i64, ptr [[PTRB]], i64 8
-; CHECK_LIMIT_4-NEXT: store i128 0, ptr [[PTR_ARM1]], align 4
+; CHECK_LIMIT_4-NEXT: store i128 0, ptr [[PTR_ARM1]], align 16
; CHECK_LIMIT_4-NEXT: br label %[[IFB_JOIN:.*]]
; CHECK_LIMIT_4: [[IFB_ARM2]]:
; CHECK_LIMIT_4-NEXT: [[PTR_ARM2:%.*]] = getelementptr i64, ptr [[PTRB]], i64 16
-; CHECK_LIMIT_4-NEXT: store i128 0, ptr [[PTR_ARM2]], align 4
+; CHECK_LIMIT_4-NEXT: store i128 0, ptr [[PTR_ARM2]], align 16
; CHECK_LIMIT_4-NEXT: br label %[[IFB_JOIN]]
; CHECK_LIMIT_4: [[IFB_JOIN]]:
; CHECK_LIMIT_4-NEXT: [[PTRC:%.*]] = phi ptr [ [[PTR_ARM1]], %[[IFB_ARM1]] ], [ [[PTR_ARM2]], %[[IFB_ARM2]] ]
diff --git a/llvm/test/tools/llubi/loadstore_be.ll b/llvm/test/tools/llubi/loadstore_be.ll
index fe2733dc8149c..7dbdfbea68726 100644
--- a/llvm/test/tools/llubi/loadstore_be.ll
+++ b/llvm/test/tools/llubi/loadstore_be.ll
@@ -165,7 +165,7 @@ define void @main() {
store ptr %alloc_ptr_array, ptr %ptr_array_2
%ptr_array_2_middle = getelementptr i8, ptr %alloc_ptr_array, i64 20
store i8 0, ptr %ptr_array_2_middle
- %mixed_alloc = load b256, ptr %alloc_ptr_array
+ %mixed_alloc = load b256, ptr %alloc_ptr_array, align 8
ret void
}
diff --git a/llvm/test/tools/llubi/loadstore_le.ll b/llvm/test/tools/llubi/loadstore_le.ll
index 9c434461380cc..65ba3c97b778e 100644
--- a/llvm/test/tools/llubi/loadstore_le.ll
+++ b/llvm/test/tools/llubi/loadstore_le.ll
@@ -166,7 +166,7 @@ define void @main() {
store ptr %alloc_ptr_array, ptr %ptr_array_2
%ptr_array_2_middle = getelementptr i8, ptr %alloc_ptr_array, i64 20
store i8 0, ptr %ptr_array_2_middle
- %mixed_alloc = load b256, ptr %alloc_ptr_array
+ %mixed_alloc = load b256, ptr %alloc_ptr_array, align 8
ret void
}
diff --git a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
index 084dcb0a5847f..55e82b8a2aa29 100644
--- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
@@ -7421,7 +7421,7 @@ TEST_F(OpenMPIRBuilderTest, CreateTask) {
ConstantInt *SharedsSize =
dyn_cast<ConstantInt>(TaskAllocCall->getOperand(4));
EXPECT_EQ(SharedsSize->getSExtValue(),
- 24); // 64-bit pointer + 128-bit integer
+ 32); // 64-bit pointer + padding + 128-bit integer
// Verify Wrapper function
Function *OutlinedFn =
More information about the cfe-commits
mailing list