[llvm] [SelectionDAG] Preserve cache hint metadata during legalization (PR #225273)
Yonah Goldberg via llvm-commits
llvm-commits at lists.llvm.org
Fri Oct 2 11:59:03 PDT 2026
https://github.com/YonahGoldberg updated https://github.com/llvm/llvm-project/pull/225273
>From 3fbf514cd541666ae6fff245e9a63332ff301367 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Tue, 22 Sep 2026 02:23:40 +0000
Subject: [PATCH 1/7] [SelectionDAG] Add baseline cache hint legalization tests
---
.../CodeGen/NVPTX/cache-hint-transforms.ll | 214 +++++++++++++++++-
.../NVPTX/cache-hint-unaligned-store-aa.ll | 117 ++++++++++
.../RISCV/mem-cache-hint-legalize-types.ll | 134 +++++++++++
3 files changed, 463 insertions(+), 2 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
create mode 100644 llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll b/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
index afa2342d025c7..6f6d805974ce6 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
@@ -1,4 +1,4 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --filter "^\s*(?:mov\.b64|ld(?:\.[A-Za-z0-9_:]+)*\.global|st(?:\.[A-Za-z0-9_:]+)*\.global|atom(?:\.[A-Za-z0-9_:]+)*\.global|ld\.param\.b32)" --version 6
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --filter "^\s*(?:mov\.b64|ld(?:\.[A-Za-z0-9_:]+)*\.global|st(?:\.[A-Za-z0-9_:]+)*\.global|atom(?:\.[A-Za-z0-9_:]+)*\.global|ld\.param\.b32)" --no-generate-body-for-unused-prefixes --version 6
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | FileCheck %s --check-prefixes=CHECK,O2
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 -O0 | FileCheck %s --check-prefixes=CHECK,O0
; RUN: %if ptxas-sm_80 && ptxas-isa-7.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | %ptxas-verify -arch=sm_80 %}
@@ -164,6 +164,216 @@ define void @test_legalize_scalarize_store_v3i64(ptr addrspace(1) %p, <3 x i64>
ret void
}
+define void @legalize_unaligned_f64_store(ptr addrspace(1) %p, double %v) {
+; O2-LABEL: legalize_unaligned_f64_store(
+; O2: st.global.b8 [%rd1], %rd2;
+; O2: st.global.b8 [%rd1+7], %rd3;
+; O2: st.global.b8 [%rd1+6], %rd4;
+; O2: st.global.b8 [%rd1+5], %rd5;
+; O2: st.global.b8 [%rd1+4], %rd6;
+; O2: st.global.b8 [%rd1+3], %rd7;
+; O2: st.global.b8 [%rd1+2], %rd8;
+; O2: st.global.b8 [%rd1+1], %rd9;
+;
+; O0-LABEL: legalize_unaligned_f64_store(
+; O0: st.global.b8 [%rd1+7], %rd3;
+; O0: st.global.b8 [%rd1+6], %rd4;
+; O0: st.global.b8 [%rd1+5], %rd5;
+; O0: st.global.b8 [%rd1+4], %rd6;
+; O0: st.global.b8 [%rd1+3], %rd7;
+; O0: st.global.b8 [%rd1+2], %rd8;
+; O0: st.global.b8 [%rd1], %rd2;
+; O0: st.global.b8 [%rd1+1], %rd9;
+ store double %v, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
+ ret void
+}
+
+define void @legalize_unaligned_v3i64_store(ptr addrspace(1) %p, <3 x i64> %v) {
+; O2-LABEL: legalize_unaligned_v3i64_store(
+; O2: st.global.b8 [%rd1+16], %rd2;
+; O2: st.global.b8 [%rd1+8], %rd4;
+; O2: st.global.b8 [%rd1], %rd3;
+; O2: st.global.b8 [%rd1+23], %rd5;
+; O2: st.global.b8 [%rd1+22], %rd6;
+; O2: st.global.b8 [%rd1+21], %rd7;
+; O2: st.global.b8 [%rd1+20], %rd8;
+; O2: st.global.b8 [%rd1+19], %rd9;
+; O2: st.global.b8 [%rd1+18], %rd10;
+; O2: st.global.b8 [%rd1+17], %rd11;
+; O2: st.global.b8 [%rd1+15], %rd12;
+; O2: st.global.b8 [%rd1+14], %rd13;
+; O2: st.global.b8 [%rd1+13], %rd14;
+; O2: st.global.b8 [%rd1+12], %rd15;
+; O2: st.global.b8 [%rd1+11], %rd16;
+; O2: st.global.b8 [%rd1+10], %rd17;
+; O2: st.global.b8 [%rd1+9], %rd18;
+; O2: st.global.b8 [%rd1+7], %rd19;
+; O2: st.global.b8 [%rd1+6], %rd20;
+; O2: st.global.b8 [%rd1+5], %rd21;
+; O2: st.global.b8 [%rd1+4], %rd22;
+; O2: st.global.b8 [%rd1+3], %rd23;
+; O2: st.global.b8 [%rd1+2], %rd24;
+; O2: st.global.b8 [%rd1+1], %rd25;
+;
+; O0-LABEL: legalize_unaligned_v3i64_store(
+; O0: st.global.b8 [%rd1+23], %rd5;
+; O0: st.global.b8 [%rd1+22], %rd6;
+; O0: st.global.b8 [%rd1+21], %rd7;
+; O0: st.global.b8 [%rd1+20], %rd8;
+; O0: st.global.b8 [%rd1+19], %rd9;
+; O0: st.global.b8 [%rd1+18], %rd10;
+; O0: st.global.b8 [%rd1+16], %rd4;
+; O0: st.global.b8 [%rd1+17], %rd11;
+; O0: st.global.b8 [%rd1+15], %rd12;
+; O0: st.global.b8 [%rd1+14], %rd13;
+; O0: st.global.b8 [%rd1+13], %rd14;
+; O0: st.global.b8 [%rd1+12], %rd15;
+; O0: st.global.b8 [%rd1+11], %rd16;
+; O0: st.global.b8 [%rd1+10], %rd17;
+; O0: st.global.b8 [%rd1+8], %rd3;
+; O0: st.global.b8 [%rd1+9], %rd18;
+; O0: st.global.b8 [%rd1+7], %rd19;
+; O0: st.global.b8 [%rd1+6], %rd20;
+; O0: st.global.b8 [%rd1+5], %rd21;
+; O0: st.global.b8 [%rd1+4], %rd22;
+; O0: st.global.b8 [%rd1+3], %rd23;
+; O0: st.global.b8 [%rd1+2], %rd24;
+; O0: st.global.b8 [%rd1], %rd2;
+; O0: st.global.b8 [%rd1+1], %rd25;
+ store <3 x i64> %v, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
+ ret void
+}
+
+define double @legalize_unaligned_f64_load(ptr addrspace(1) %p) {
+; CHECK-LABEL: legalize_unaligned_f64_load(
+; CHECK: ld.global.b8 %rd2, [%rd1];
+; CHECK: ld.global.b8 %rd3, [%rd1+1];
+; CHECK: ld.global.b8 %rd6, [%rd1+2];
+; CHECK: ld.global.b8 %rd8, [%rd1+3];
+; CHECK: ld.global.b8 %rd12, [%rd1+4];
+; CHECK: ld.global.b8 %rd13, [%rd1+5];
+; CHECK: ld.global.b8 %rd16, [%rd1+6];
+; CHECK: ld.global.b8 %rd18, [%rd1+7];
+ %v = load double, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
+ ret double %v
+}
+
+define <3 x i64> @legalize_unaligned_v3i64_load(ptr addrspace(1) %p) {
+; CHECK-LABEL: legalize_unaligned_v3i64_load(
+; CHECK: ld.global.b8 %rd2, [%rd1+16];
+; CHECK: ld.global.b8 %rd3, [%rd1+17];
+; CHECK: ld.global.b8 %rd6, [%rd1+18];
+; CHECK: ld.global.b8 %rd8, [%rd1+19];
+; CHECK: ld.global.b8 %rd12, [%rd1+20];
+; CHECK: ld.global.b8 %rd13, [%rd1+21];
+; CHECK: ld.global.b8 %rd16, [%rd1+22];
+; CHECK: ld.global.b8 %rd18, [%rd1+23];
+; CHECK: ld.global.b8 %rd24, [%rd1+8];
+; CHECK: ld.global.b8 %rd25, [%rd1+9];
+; CHECK: ld.global.b8 %rd28, [%rd1+10];
+; CHECK: ld.global.b8 %rd30, [%rd1+11];
+; CHECK: ld.global.b8 %rd34, [%rd1+12];
+; CHECK: ld.global.b8 %rd35, [%rd1+13];
+; CHECK: ld.global.b8 %rd38, [%rd1+14];
+; CHECK: ld.global.b8 %rd40, [%rd1+15];
+; CHECK: ld.global.b8 %rd46, [%rd1];
+; CHECK: ld.global.b8 %rd47, [%rd1+1];
+; CHECK: ld.global.b8 %rd50, [%rd1+2];
+; CHECK: ld.global.b8 %rd52, [%rd1+3];
+; CHECK: ld.global.b8 %rd56, [%rd1+4];
+; CHECK: ld.global.b8 %rd57, [%rd1+5];
+; CHECK: ld.global.b8 %rd60, [%rd1+6];
+; CHECK: ld.global.b8 %rd62, [%rd1+7];
+ %v = load <3 x i64>, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
+ ret <3 x i64> %v
+}
+
+define i32 @legalize_i24_load(ptr addrspace(1) %p) {
+; CHECK-LABEL: legalize_i24_load(
+; CHECK: ld.global.b8 %r1, [%rd1];
+; CHECK: ld.global.b8 %r2, [%rd1+1];
+; CHECK: ld.global.b8 %r5, [%rd1+2];
+ %v = load i24, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
+ %ext = zext i24 %v to i32
+ ret i32 %ext
+}
+
+define void @legalize_i24_store(ptr addrspace(1) %p, i32 %v) {
+; CHECK-LABEL: legalize_i24_store(
+; CHECK: ld.param.b32 %r1, [legalize_i24_store_param_1];
+; CHECK: st.global.b8 [%rd1], %r1;
+; CHECK: st.global.b8 [%rd1+1], %r2;
+; CHECK: st.global.b8 [%rd1+2], %r3;
+ %trunc = trunc i32 %v to i24
+ store i24 %trunc, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
+ ret void
+}
+
+define <1 x i64> @legalize_scalarize_load_v1i64(ptr addrspace(1) %p) {
+; CHECK-LABEL: legalize_scalarize_load_v1i64(
+; CHECK: ld.global.b64 %rd2, [%rd1];
+ %v = load <1 x i64>, ptr addrspace(1) %p, align 8, !mem.cache_hint !13
+ ret <1 x i64> %v
+}
+
+define void @legalize_scalarize_store_v1i64(ptr addrspace(1) %p, <1 x i64> %v) {
+; CHECK-LABEL: legalize_scalarize_store_v1i64(
+; CHECK: st.global.b64 [%rd1], %rd2;
+ store <1 x i64> %v, ptr addrspace(1) %p, align 8, !mem.cache_hint !14
+ ret void
+}
+
+define <3 x i1> @legalize_scalarize_load_v3i1(ptr addrspace(1) %p) {
+; CHECK-LABEL: legalize_scalarize_load_v3i1(
+; CHECK: ld.global.b8 %rs1, [%rd1];
+ %v = load <3 x i1>, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
+ ret <3 x i1> %v
+}
+
+define void @legalize_scalarize_store_v3i1(ptr addrspace(1) %p, <3 x i1> %v) {
+; O2-LABEL: legalize_scalarize_store_v3i1(
+; O2: st.global.b8 [%rd1], %rs10;
+;
+; O0-LABEL: legalize_scalarize_store_v3i1(
+; O0: st.global.b8 [%rd1], %rs11;
+ store <3 x i1> %v, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
+ ret void
+}
+
+define <3 x i32> @legalize_widen_extload_v3i8(ptr addrspace(1) %p) {
+; CHECK-LABEL: legalize_widen_extload_v3i8(
+; CHECK: ld.global.b8 %rs1, [%rd1];
+; CHECK: ld.global.b8 %rs2, [%rd1+1];
+; CHECK: ld.global.b8 %r4, [%rd1+2];
+ %v = load <3 x i8>, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
+ %ext = zext <3 x i8> %v to <3 x i32>
+ ret <3 x i32> %ext
+}
+
+define void @legalize_widen_truncstore_v3i8(ptr addrspace(1) %p, <3 x i32> %v) {
+; O2-LABEL: legalize_widen_truncstore_v3i8(
+; O2: st.global.b8 [%rd1+2], %rs1;
+; O2: st.global.b8 [%rd1], %rs2;
+; O2: st.global.b8 [%rd1+1], %rs3;
+;
+; O0-LABEL: legalize_widen_truncstore_v3i8(
+; O0: ld.param.b32 %r3, [legalize_widen_truncstore_v3i8_param_1+8];
+; O0: st.global.b8 [%rd1], %rs1;
+; O0: st.global.b8 [%rd1+1], %rs2;
+; O0: st.global.b8 [%rd1+2], %r3;
+ %trunc = trunc <3 x i32> %v to <3 x i8>
+ store <3 x i8> %trunc, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
+ ret void
+}
+
+define i32 @scalarize_extracted_vector_load(ptr addrspace(1) %p, i64 %idx) {
+; CHECK-LABEL: scalarize_extracted_vector_load(
+; CHECK: ld.global.b32 %r1, [%rd5];
+ %v = load <4 x i32>, ptr addrspace(1) %p, align 16, !mem.cache_hint !13
+ %elt = extractelement <4 x i32> %v, i64 %idx
+ ret i32 %elt
+}
+
;-----------------------------------------------------------------------------
; Metadata definitions
;-----------------------------------------------------------------------------
@@ -197,4 +407,4 @@ define void @test_legalize_scalarize_store_v3i64(ptr addrspace(1) %p, <3 x i64>
!13 = !{i32 0, !28}
!28 = !{!"nvvm.l1_eviction", !"last"}
!14 = !{i32 1, !29}
-!29 = !{!"nvvm.l2_eviction", !"first"}
+!29 = !{!"nvvm.l1_eviction", !"first"}
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll b/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
new file mode 100644
index 0000000000000..0b354dfb13787
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
@@ -0,0 +1,117 @@
+; NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 \
+; RUN: -stop-after=finalize-isel %s -o - | FileCheck %s
+
+define void @unaligned_f64_store(ptr addrspace(1) %p, double %v) {
+ ; CHECK-LABEL: name: unaligned_f64_store
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: [[LD_i64_:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_f64_store_param_0>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
+ ; CHECK-NEXT: [[LD_i64_1:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_f64_store_param_1>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
+ ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 0, 0, $noreg :: (store (s8) into %ir.p, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 56
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 7, 0, $noreg :: (store (s8) into %ir.p + 7, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri1:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 48
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 6, 0, $noreg :: (store (s8) into %ir.p + 6, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri2:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 40
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 5, 0, $noreg :: (store (s8) into %ir.p + 5, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri3:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 32
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 4, 0, $noreg :: (store (s8) into %ir.p + 4, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri4:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 24
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 3, 0, $noreg :: (store (s8) into %ir.p + 3, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri5:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 16
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 2, 0, $noreg :: (store (s8) into %ir.p + 2, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri6:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 8
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 1, 0, $noreg :: (store (s8) into %ir.p + 1, addrspace 1)
+ ; CHECK-NEXT: Return
+ store double %v, ptr addrspace(1) %p, align 1, !alias.scope !2, !noalias !4, !mem.cache_hint !5
+ ret void
+}
+
+define void @unaligned_v3i64_store(ptr addrspace(1) %p, <3 x i64> %v) {
+ ; CHECK-LABEL: name: unaligned_v3i64_store
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: [[LD_i64_:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_v3i64_store_param_0>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
+ ; CHECK-NEXT: [[LD_i64_1:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_v3i64_store_param_1>, 16, 0, $noreg :: (dereferenceable invariant load (s64), align 16, addrspace 102)
+ ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 16, 0, $noreg :: (store (s8) into %ir.p + 16, addrspace 1)
+ ; CHECK-NEXT: [[LDV_i64_v2_:%[0-9]+]]:b64, [[LDV_i64_v2_1:%[0-9]+]]:b64 = LDV_i64_v2 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_v3i64_store_param_1>, 0, 0, $noreg :: (dereferenceable invariant load (s128), align 32, addrspace 102)
+ ; CHECK-NEXT: ST_i64 [[LDV_i64_v2_1]], 0, 0, 1, 8, [[LD_i64_]], 8, 0, $noreg :: (store (s8) into %ir.p + 8, addrspace 1)
+ ; CHECK-NEXT: ST_i64 [[LDV_i64_v2_]], 0, 0, 1, 8, [[LD_i64_]], 0, 0, $noreg :: (store (s8) into %ir.p, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 56
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 23, 0, $noreg :: (store (s8) into %ir.p + 23, !alias.scope !0, !noalias !3, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri1:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 48
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 22, 0, $noreg :: (store (s8) into %ir.p + 22, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri2:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 40
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 21, 0, $noreg :: (store (s8) into %ir.p + 21, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri3:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 32
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 20, 0, $noreg :: (store (s8) into %ir.p + 20, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri4:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 24
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 19, 0, $noreg :: (store (s8) into %ir.p + 19, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri5:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 16
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 18, 0, $noreg :: (store (s8) into %ir.p + 18, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri6:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 8
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 17, 0, $noreg :: (store (s8) into %ir.p + 17, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri7:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 56
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri7]], 0, 0, 1, 8, [[LD_i64_]], 15, 0, $noreg :: (store (s8) into %ir.p + 15, !alias.scope !0, !noalias !3, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri8:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 48
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri8]], 0, 0, 1, 8, [[LD_i64_]], 14, 0, $noreg :: (store (s8) into %ir.p + 14, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri9:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 40
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri9]], 0, 0, 1, 8, [[LD_i64_]], 13, 0, $noreg :: (store (s8) into %ir.p + 13, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri10:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 32
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri10]], 0, 0, 1, 8, [[LD_i64_]], 12, 0, $noreg :: (store (s8) into %ir.p + 12, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri11:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 24
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri11]], 0, 0, 1, 8, [[LD_i64_]], 11, 0, $noreg :: (store (s8) into %ir.p + 11, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri12:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 16
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri12]], 0, 0, 1, 8, [[LD_i64_]], 10, 0, $noreg :: (store (s8) into %ir.p + 10, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri13:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 8
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri13]], 0, 0, 1, 8, [[LD_i64_]], 9, 0, $noreg :: (store (s8) into %ir.p + 9, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri14:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 56
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri14]], 0, 0, 1, 8, [[LD_i64_]], 7, 0, $noreg :: (store (s8) into %ir.p + 7, !alias.scope !0, !noalias !3, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri15:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 48
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri15]], 0, 0, 1, 8, [[LD_i64_]], 6, 0, $noreg :: (store (s8) into %ir.p + 6, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri16:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 40
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri16]], 0, 0, 1, 8, [[LD_i64_]], 5, 0, $noreg :: (store (s8) into %ir.p + 5, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri17:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 32
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri17]], 0, 0, 1, 8, [[LD_i64_]], 4, 0, $noreg :: (store (s8) into %ir.p + 4, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri18:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 24
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri18]], 0, 0, 1, 8, [[LD_i64_]], 3, 0, $noreg :: (store (s8) into %ir.p + 3, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri19:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 16
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri19]], 0, 0, 1, 8, [[LD_i64_]], 2, 0, $noreg :: (store (s8) into %ir.p + 2, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri20:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 8
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri20]], 0, 0, 1, 8, [[LD_i64_]], 1, 0, $noreg :: (store (s8) into %ir.p + 1, addrspace 1)
+ ; CHECK-NEXT: Return
+ store <3 x i64> %v, ptr addrspace(1) %p, align 1, !alias.scope !2, !noalias !4, !mem.cache_hint !5
+ ret void
+}
+
+define void @unaligned_i64_store(ptr addrspace(1) %p, i64 %v) {
+ ; CHECK-LABEL: name: unaligned_i64_store
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: [[LD_i64_:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_i64_store_param_0>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
+ ; CHECK-NEXT: [[LD_i64_1:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_i64_store_param_1>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
+ ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 0, 0, $noreg :: (store (s8) into %ir.p, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 56
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 7, 0, $noreg :: (store (s8) into %ir.p + 7, !alias.scope !0, !noalias !3, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri1:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 48
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 6, 0, $noreg :: (store (s8) into %ir.p + 6, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri2:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 40
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 5, 0, $noreg :: (store (s8) into %ir.p + 5, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri3:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 32
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 4, 0, $noreg :: (store (s8) into %ir.p + 4, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri4:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 24
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 3, 0, $noreg :: (store (s8) into %ir.p + 3, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri5:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 16
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 2, 0, $noreg :: (store (s8) into %ir.p + 2, addrspace 1)
+ ; CHECK-NEXT: [[SRL64_ri6:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 8
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 1, 0, $noreg :: (store (s8) into %ir.p + 1, addrspace 1)
+ ; CHECK-NEXT: Return
+ store i64 %v, ptr addrspace(1) %p, align 1, !alias.scope !2, !noalias !4, !mem.cache_hint !5
+ ret void
+}
+
+!0 = distinct !{!0, !"domain"}
+!1 = distinct !{!1, !0, !"scope"}
+!2 = !{!1}
+!3 = distinct !{!3, !0, !"other scope"}
+!4 = !{!3}
+!5 = !{i32 1, !6}
+!6 = !{!"nvvm.l1_eviction", !"first"}
diff --git a/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll b/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
new file mode 100644
index 0000000000000..7a2d74d0c9321
--- /dev/null
+++ b/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
@@ -0,0 +1,134 @@
+; NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=riscv32 -stop-after=finalize-isel %s -o - | FileCheck %s
+
+define fp128 @soften_f128_load(ptr %p) {
+ ; CHECK-LABEL: name: soften_f128_load
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: liveins: $x10, $x11
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x11
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:gpr = COPY [[COPY1]]
+ ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 16)
+ ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 16)
+ ; CHECK-NEXT: [[LW2:%[0-9]+]]:gpr = LW [[COPY]], 8 :: (load (s32) from %ir.p + 8, align 8, basealign 16)
+ ; CHECK-NEXT: [[LW3:%[0-9]+]]:gpr = LW [[COPY]], 12 :: (load (s32) from %ir.p + 12, basealign 16)
+ ; CHECK-NEXT: SW killed [[LW3]], [[COPY2]], 12 :: (store (s32) into unknown-address + 12, basealign 16)
+ ; CHECK-NEXT: SW killed [[LW2]], [[COPY2]], 8 :: (store (s32) into unknown-address + 8, align 8, basealign 16)
+ ; CHECK-NEXT: SW killed [[LW1]], [[COPY2]], 4 :: (store (s32) into unknown-address + 4, basealign 16)
+ ; CHECK-NEXT: SW killed [[LW]], [[COPY2]], 0 :: (store (s32), align 16)
+ ; CHECK-NEXT: PseudoRET
+ %v = load fp128, ptr %p, align 16, !mem.cache_hint !0
+ ret fp128 %v
+}
+
+define double @soften_f16_extload(ptr %p) {
+ ; CHECK-LABEL: name: soften_f16_extload
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: liveins: $x10
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x10
+ ; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def dead $x2, implicit $x2
+ ; CHECK-NEXT: [[LH:%[0-9]+]]:gpr = LH [[COPY]], 0 :: (load (s16) from %ir.p)
+ ; CHECK-NEXT: $x10 = COPY [[LH]]
+ ; CHECK-NEXT: PseudoCALL target-flags(riscv-call) &__extendhfsf2, csr_ilp32_lp64, implicit-def dead $x1, implicit $x10, implicit-def $x2, implicit-def $x10
+ ; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def dead $x2, implicit $x2
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
+ ; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def dead $x2, implicit $x2
+ ; CHECK-NEXT: $x10 = COPY [[COPY1]]
+ ; CHECK-NEXT: PseudoCALL target-flags(riscv-call) &__extendsfdf2, csr_ilp32_lp64, implicit-def dead $x1, implicit $x10, implicit-def $x2, implicit-def $x10, implicit-def $x11
+ ; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def dead $x2, implicit $x2
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:gpr = COPY $x10
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:gpr = COPY $x11
+ ; CHECK-NEXT: $x10 = COPY [[COPY2]]
+ ; CHECK-NEXT: $x11 = COPY [[COPY3]]
+ ; CHECK-NEXT: PseudoRET implicit $x10, implicit $x11
+ %v = load half, ptr %p, align 2, !mem.cache_hint !0
+ %ext = fpext half %v to double
+ ret double %ext
+}
+
+define i128 @expand_i128_load(ptr %p) {
+ ; CHECK-LABEL: name: expand_i128_load
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: liveins: $x10, $x11
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x11
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:gpr = COPY [[COPY1]]
+ ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 16)
+ ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 16)
+ ; CHECK-NEXT: [[LW2:%[0-9]+]]:gpr = LW [[COPY]], 8 :: (load (s32) from %ir.p + 8, align 8, basealign 16)
+ ; CHECK-NEXT: [[LW3:%[0-9]+]]:gpr = LW [[COPY]], 12 :: (load (s32) from %ir.p + 12, basealign 16)
+ ; CHECK-NEXT: SW killed [[LW3]], [[COPY2]], 12 :: (store (s32) into unknown-address + 12, basealign 8)
+ ; CHECK-NEXT: SW killed [[LW2]], [[COPY2]], 8 :: (store (s32) into unknown-address + 8, align 8)
+ ; CHECK-NEXT: SW killed [[LW1]], [[COPY2]], 4 :: (store (s32) into unknown-address + 4, basealign 8)
+ ; CHECK-NEXT: SW killed [[LW]], [[COPY2]], 0 :: (store (s32), align 8)
+ ; CHECK-NEXT: PseudoRET
+ %v = load i128, ptr %p, align 16, !range !2, !mem.cache_hint !0
+ ret i128 %v
+}
+
+define void @expand_i128_store(ptr %p, i128 %v) {
+ ; CHECK-LABEL: name: expand_i128_store
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: liveins: $x10, $x11
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x11
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
+ ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32))
+ ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32))
+ ; CHECK-NEXT: [[LW2:%[0-9]+]]:gpr = LW [[COPY]], 8 :: (load (s32))
+ ; CHECK-NEXT: [[LW3:%[0-9]+]]:gpr = LW [[COPY]], 12 :: (load (s32))
+ ; CHECK-NEXT: SW killed [[LW3]], [[COPY1]], 12 :: (store (s32) into %ir.p + 12, basealign 16)
+ ; CHECK-NEXT: SW killed [[LW2]], [[COPY1]], 8 :: (store (s32) into %ir.p + 8, align 8, basealign 16)
+ ; CHECK-NEXT: SW killed [[LW1]], [[COPY1]], 4 :: (store (s32) into %ir.p + 4, basealign 16)
+ ; CHECK-NEXT: SW killed [[LW]], [[COPY1]], 0 :: (store (s32) into %ir.p, align 16)
+ ; CHECK-NEXT: PseudoRET
+ store i128 %v, ptr %p, align 16, !mem.cache_hint !1
+ ret void
+}
+
+define i128 @expand_i128_extload(ptr %p) {
+ ; CHECK-LABEL: name: expand_i128_extload
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: liveins: $x10, $x11
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x11
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:gpr = COPY [[COPY1]]
+ ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 8)
+ ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 8)
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:gpr = COPY $x0
+ ; CHECK-NEXT: SW [[COPY3]], [[COPY2]], 12 :: (store (s32) into unknown-address + 12, basealign 8)
+ ; CHECK-NEXT: SW [[COPY3]], [[COPY2]], 8 :: (store (s32) into unknown-address + 8, align 8)
+ ; CHECK-NEXT: SW killed [[LW1]], [[COPY2]], 4 :: (store (s32) into unknown-address + 4, basealign 8)
+ ; CHECK-NEXT: SW killed [[LW]], [[COPY2]], 0 :: (store (s32), align 8)
+ ; CHECK-NEXT: PseudoRET
+ %v = load i64, ptr %p, align 8, !range !3, !mem.cache_hint !0
+ %ext = zext i64 %v to i128
+ ret i128 %ext
+}
+
+define void @expand_i128_truncstore(ptr %p, i128 %v) {
+ ; CHECK-LABEL: name: expand_i128_truncstore
+ ; CHECK: bb.0 (%ir-block.0):
+ ; CHECK-NEXT: liveins: $x10, $x11
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x11
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
+ ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32))
+ ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32))
+ ; CHECK-NEXT: SW killed [[LW1]], [[COPY1]], 4 :: (store (s32) into %ir.p + 4, basealign 8)
+ ; CHECK-NEXT: SW killed [[LW]], [[COPY1]], 0 :: (store (s32) into %ir.p, align 8)
+ ; CHECK-NEXT: PseudoRET
+ %trunc = trunc i128 %v to i64
+ store i64 %trunc, ptr %p, align 8, !mem.cache_hint !1
+ ret void
+}
+
+!0 = !{i32 0, !4}
+!1 = !{i32 1, !4}
+!2 = !{i128 0, i128 1000}
+!3 = !{i64 0, i64 1000}
+!4 = !{!"test.cache_hint", !"value"}
>From 112b704b685ea3227fd0b6a20fa21aac94813521 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Tue, 22 Sep 2026 03:10:52 +0000
Subject: [PATCH 2/7] [SelectionDAG] Preserve cache hints through legalization
---
llvm/include/llvm/CodeGen/SelectionDAGNodes.h | 7 +
llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp | 33 +-
.../SelectionDAG/LegalizeFloatTypes.cpp | 21 +-
.../SelectionDAG/LegalizeIntegerTypes.cpp | 24 +-
.../SelectionDAG/LegalizeTypesGeneric.cpp | 19 +-
.../SelectionDAG/LegalizeVectorTypes.cpp | 54 ++--
.../CodeGen/SelectionDAG/TargetLowering.cpp | 80 +++--
.../CodeGen/NVPTX/cache-hint-transforms.ll | 303 +++++++++---------
.../NVPTX/cache-hint-unaligned-store-aa.ll | 80 ++---
.../RISCV/mem-cache-hint-legalize-types.ll | 34 +-
10 files changed, 340 insertions(+), 315 deletions(-)
diff --git a/llvm/include/llvm/CodeGen/SelectionDAGNodes.h b/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
index 667c98f086827..a3d2f2c882642 100644
--- a/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
+++ b/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
@@ -1493,6 +1493,13 @@ class MemSDNode : public SDNode {
return getMemOperand()->getMemCacheHint();
}
+ /// Returns LLVM IR metadata carried by this memory access, except for range
+ /// metadata. Range metadata describes a loaded value and cannot be blindly
+ /// transferred when an access is split or its type changes.
+ MMOMetadata getNonRangeMMOMetadata() const {
+ return MMOMetadata(getAAInfo(), /*Ranges=*/nullptr, getMemCacheHint());
+ }
+
/// Returns the synchronization scope ID for this memory operation.
SyncScope::ID getSyncScopeID() const {
return getMemOperand()->getSyncScopeID();
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
index bd9c47806a9ca..d3de3538cb3b0 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
@@ -496,7 +496,7 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
SDLoc dl(Node);
MachineMemOperand::Flags MMOFlags = ST->getMemOperand()->getFlags();
- AAMDNodes AAInfo = ST->getAAInfo();
+ MMOMetadata Metadata = ST->getNonRangeMMOMetadata();
if (!ST->isTruncatingStore()) {
LLVM_DEBUG(dbgs() << "Legalizing store operation\n");
@@ -536,7 +536,7 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
"Can only promote stores to same size type");
Value = DAG.getNode(ISD::BITCAST, dl, NVT, Value);
SDValue Result = DAG.getStore(Chain, dl, Value, Ptr, ST->getPointerInfo(),
- ST->getBaseAlign(), MMOFlags, AAInfo);
+ ST->getBaseAlign(), MMOFlags, Metadata);
ReplaceNode(SDValue(Node, 0), Result);
break;
}
@@ -559,7 +559,7 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
Value = DAG.getZeroExtendInReg(Value, dl, StVT);
SDValue Result =
DAG.getTruncStore(Chain, dl, Value, Ptr, ST->getPointerInfo(), NVT,
- ST->getBaseAlign(), MMOFlags, AAInfo);
+ ST->getBaseAlign(), MMOFlags, Metadata);
ReplaceNode(SDValue(Node, 0), Result);
} else if (!StVT.isVector() && !isPowerOf2_64(StWidth.getFixedValue())) {
// If not storing a power-of-2 number of bits, expand as two stores.
@@ -582,7 +582,7 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
// TRUNCSTORE:i24 X -> TRUNCSTORE:i16 X, TRUNCSTORE at +2:i8 (srl X, 16)
// Store the bottom RoundWidth bits.
Lo = DAG.getTruncStore(Chain, dl, Value, Ptr, ST->getPointerInfo(),
- RoundVT, ST->getBaseAlign(), MMOFlags, AAInfo);
+ RoundVT, ST->getBaseAlign(), MMOFlags, Metadata);
// Store the remaining ExtraWidth bits.
IncrementSize = RoundWidth / 8;
@@ -593,7 +593,7 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
DAG.getShiftAmountConstant(RoundWidth, Value.getValueType(), dl));
Hi = DAG.getTruncStore(Chain, dl, Hi, Ptr,
ST->getPointerInfo().getWithOffset(IncrementSize),
- ExtraVT, ST->getBaseAlign(), MMOFlags, AAInfo);
+ ExtraVT, ST->getBaseAlign(), MMOFlags, Metadata);
} else {
// Big endian - avoid unaligned stores.
// TRUNCSTORE:i24 X -> TRUNCSTORE:i16 (srl X, 8), TRUNCSTORE at +2:i8 X
@@ -602,7 +602,7 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
ISD::SRL, dl, Value.getValueType(), Value,
DAG.getShiftAmountConstant(ExtraWidth, Value.getValueType(), dl));
Hi = DAG.getTruncStore(Chain, dl, Hi, Ptr, ST->getPointerInfo(), RoundVT,
- ST->getBaseAlign(), MMOFlags, AAInfo);
+ ST->getBaseAlign(), MMOFlags, Metadata);
// Store the remaining ExtraWidth bits.
IncrementSize = RoundWidth / 8;
@@ -611,7 +611,7 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
Ptr.getValueType()));
Lo = DAG.getTruncStore(Chain, dl, Value, Ptr,
ST->getPointerInfo().getWithOffset(IncrementSize),
- ExtraVT, ST->getBaseAlign(), MMOFlags, AAInfo);
+ ExtraVT, ST->getBaseAlign(), MMOFlags, Metadata);
}
// The order of the stores doesn't matter.
@@ -649,15 +649,16 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
if (TLI.isTypeLegal(StVT)) {
Value = DAG.getNode(ISD::TRUNCATE, dl, StVT, Value);
Result = DAG.getStore(Chain, dl, Value, Ptr, ST->getPointerInfo(),
- ST->getBaseAlign(), MMOFlags, AAInfo);
+ ST->getBaseAlign(), MMOFlags, Metadata);
} else {
// The in-memory type isn't legal. Truncate to the type it would promote
// to, and then do a truncstore.
Value = DAG.getNode(ISD::TRUNCATE, dl,
TLI.getTypeToTransformTo(*DAG.getContext(), StVT),
Value);
- Result = DAG.getTruncStore(Chain, dl, Value, Ptr, ST->getPointerInfo(),
- StVT, ST->getBaseAlign(), MMOFlags, AAInfo);
+ Result =
+ DAG.getTruncStore(Chain, dl, Value, Ptr, ST->getPointerInfo(), StVT,
+ ST->getBaseAlign(), MMOFlags, Metadata);
}
ReplaceNode(SDValue(Node, 0), Result);
@@ -736,7 +737,7 @@ void SelectionDAGLegalize::LegalizeLoadOps(SDNode *Node) {
EVT SrcVT = LD->getMemoryVT();
TypeSize SrcWidth = SrcVT.getSizeInBits();
MachineMemOperand::Flags MMOFlags = LD->getMemOperand()->getFlags();
- AAMDNodes AAInfo = LD->getAAInfo();
+ MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
if (SrcWidth != SrcVT.getStoreSizeInBits() &&
// Some targets pretend to have an i1 loading operation, and actually
@@ -764,7 +765,7 @@ void SelectionDAGLegalize::LegalizeLoadOps(SDNode *Node) {
SDValue Result = DAG.getExtLoad(NewExtType, dl, Node->getValueType(0),
Chain, Ptr, LD->getPointerInfo(), NVT,
- LD->getBaseAlign(), MMOFlags, AAInfo);
+ LD->getBaseAlign(), MMOFlags, Metadata);
Ch = Result.getValue(1); // The chain.
@@ -804,7 +805,7 @@ void SelectionDAGLegalize::LegalizeLoadOps(SDNode *Node) {
// Load the bottom RoundWidth bits.
Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, Node->getValueType(0), Chain, Ptr,
LD->getPointerInfo(), RoundVT, LD->getBaseAlign(),
- MMOFlags, AAInfo);
+ MMOFlags, Metadata);
// Load the remaining ExtraWidth bits.
IncrementSize = RoundWidth / 8;
@@ -812,7 +813,7 @@ void SelectionDAGLegalize::LegalizeLoadOps(SDNode *Node) {
DAG.getMemBasePlusOffset(Ptr, TypeSize::getFixed(IncrementSize), dl);
Hi = DAG.getExtLoad(ExtType, dl, Node->getValueType(0), Chain, Ptr,
LD->getPointerInfo().getWithOffset(IncrementSize),
- ExtraVT, LD->getBaseAlign(), MMOFlags, AAInfo);
+ ExtraVT, LD->getBaseAlign(), MMOFlags, Metadata);
// Build a factor node to remember that this load is independent of
// the other one.
@@ -832,7 +833,7 @@ void SelectionDAGLegalize::LegalizeLoadOps(SDNode *Node) {
// Load the top RoundWidth bits.
Hi = DAG.getExtLoad(ExtType, dl, Node->getValueType(0), Chain, Ptr,
LD->getPointerInfo(), RoundVT, LD->getBaseAlign(),
- MMOFlags, AAInfo);
+ MMOFlags, Metadata);
// Load the remaining ExtraWidth bits.
IncrementSize = RoundWidth / 8;
@@ -840,7 +841,7 @@ void SelectionDAGLegalize::LegalizeLoadOps(SDNode *Node) {
DAG.getMemBasePlusOffset(Ptr, TypeSize::getFixed(IncrementSize), dl);
Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, Node->getValueType(0), Chain, Ptr,
LD->getPointerInfo().getWithOffset(IncrementSize),
- ExtraVT, LD->getBaseAlign(), MMOFlags, AAInfo);
+ ExtraVT, LD->getBaseAlign(), MMOFlags, Metadata);
// Build a factor node to remember that this load is independent of
// the other one.
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
index b7efd80d93fd4..1b9b68705195a 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
@@ -920,7 +920,7 @@ SDValue DAGTypeLegalizer::SoftenFloatRes_LOAD(SDNode *N) {
NewL = DAG.getLoad(L->getAddressingMode(), ISD::EXTLOAD, NVT, dl,
L->getChain(), L->getBasePtr(), L->getOffset(),
L->getPointerInfo(), MemVT, L->getBaseAlign(), MMOFlags,
- L->getAAInfo());
+ L->getNonRangeMMOMetadata());
// Legalized the chain result - switch anything that used the old chain to
// use the new one.
ReplaceValueWith(SDValue(N, 1), NewL.getValue(1));
@@ -928,10 +928,11 @@ SDValue DAGTypeLegalizer::SoftenFloatRes_LOAD(SDNode *N) {
}
// Do a non-extending load followed by FP_EXTEND.
- NewL = DAG.getLoad(L->getAddressingMode(), ISD::NON_EXTLOAD, L->getMemoryVT(),
- dl, L->getChain(), L->getBasePtr(), L->getOffset(),
- L->getPointerInfo(), L->getMemoryVT(), L->getBaseAlign(),
- MMOFlags, L->getAAInfo());
+ NewL = DAG.getLoad(
+ L->getAddressingMode(), ISD::NON_EXTLOAD, L->getMemoryVT(), dl,
+ L->getChain(), L->getBasePtr(), L->getOffset(), L->getPointerInfo(),
+ L->getMemoryVT(), L->getBaseAlign(), MMOFlags,
+ L->getNonRangeMMOMetadata());
// Legalized the chain result - switch anything that used the old chain to
// use the new one.
ReplaceValueWith(SDValue(N, 1), NewL.getValue(1));
@@ -2818,11 +2819,11 @@ SDValue DAGTypeLegalizer::SoftPromoteHalfRes_LOAD(SDNode *N) {
// Load the value as an integer value with the same number of bits.
assert(L->getExtensionType() == ISD::NON_EXTLOAD && "Unexpected extension!");
- SDValue NewL =
- DAG.getLoad(L->getAddressingMode(), L->getExtensionType(), MVT::i16,
- SDLoc(N), L->getChain(), L->getBasePtr(), L->getOffset(),
- L->getPointerInfo(), MVT::i16, L->getBaseAlign(),
- L->getMemOperand()->getFlags(), L->getAAInfo());
+ SDValue NewL = DAG.getLoad(
+ L->getAddressingMode(), L->getExtensionType(), MVT::i16, SDLoc(N),
+ L->getChain(), L->getBasePtr(), L->getOffset(), L->getPointerInfo(),
+ MVT::i16, L->getBaseAlign(), L->getMemOperand()->getFlags(),
+ L->getNonRangeMMOMetadata());
// Legalize the chain result by replacing uses of the old value chain with the
// new one
ReplaceValueWith(SDValue(N, 1), NewL.getValue(1));
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
index 55d89ed05545b..60e5788262992 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
@@ -4428,7 +4428,7 @@ void DAGTypeLegalizer::ExpandIntRes_LOAD(LoadSDNode *N,
SDValue Ptr = N->getBasePtr();
ISD::LoadExtType ExtType = N->getExtensionType();
MachineMemOperand::Flags MMOFlags = N->getMemOperand()->getFlags();
- AAMDNodes AAInfo = N->getAAInfo();
+ MMOMetadata Metadata = N->getNonRangeMMOMetadata();
SDLoc dl(N);
assert(NVT.isByteSized() && "Expanded type not byte sized!");
@@ -4437,7 +4437,7 @@ void DAGTypeLegalizer::ExpandIntRes_LOAD(LoadSDNode *N,
EVT MemVT = N->getMemoryVT();
Lo = DAG.getExtLoad(ExtType, dl, NVT, Ch, Ptr, N->getPointerInfo(), MemVT,
- N->getBaseAlign(), MMOFlags, AAInfo);
+ N->getBaseAlign(), MMOFlags, Metadata);
// Remember the chain.
Ch = Lo.getValue(1);
@@ -4459,7 +4459,7 @@ void DAGTypeLegalizer::ExpandIntRes_LOAD(LoadSDNode *N,
} else if (DAG.getDataLayout().isLittleEndian()) {
// Little-endian - low bits are at low addresses.
Lo = DAG.getLoad(NVT, dl, Ch, Ptr, N->getPointerInfo(), N->getBaseAlign(),
- MMOFlags, AAInfo);
+ MMOFlags, Metadata);
unsigned ExcessBits =
N->getMemoryVT().getSizeInBits() - NVT.getSizeInBits();
@@ -4470,7 +4470,7 @@ void DAGTypeLegalizer::ExpandIntRes_LOAD(LoadSDNode *N,
Ptr = DAG.getMemBasePlusOffset(Ptr, TypeSize::getFixed(IncrementSize), dl);
Hi = DAG.getExtLoad(ExtType, dl, NVT, Ch, Ptr,
N->getPointerInfo().getWithOffset(IncrementSize), NEVT,
- N->getBaseAlign(), MMOFlags, AAInfo);
+ N->getBaseAlign(), MMOFlags, Metadata);
// Build a factor node to remember that this load is independent of the
// other one.
@@ -4488,7 +4488,7 @@ void DAGTypeLegalizer::ExpandIntRes_LOAD(LoadSDNode *N,
Hi = DAG.getExtLoad(ExtType, dl, NVT, Ch, Ptr, N->getPointerInfo(),
EVT::getIntegerVT(*DAG.getContext(),
MemVT.getSizeInBits() - ExcessBits),
- N->getBaseAlign(), MMOFlags, AAInfo);
+ N->getBaseAlign(), MMOFlags, Metadata);
// Increment the pointer to the other half.
Ptr = DAG.getMemBasePlusOffset(Ptr, TypeSize::getFixed(IncrementSize), dl);
@@ -4496,7 +4496,7 @@ void DAGTypeLegalizer::ExpandIntRes_LOAD(LoadSDNode *N,
Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, NVT, Ch, Ptr,
N->getPointerInfo().getWithOffset(IncrementSize),
EVT::getIntegerVT(*DAG.getContext(), ExcessBits),
- N->getBaseAlign(), MMOFlags, AAInfo);
+ N->getBaseAlign(), MMOFlags, Metadata);
// Build a factor node to remember that this load is independent of the
// other one.
@@ -5991,7 +5991,7 @@ SDValue DAGTypeLegalizer::ExpandIntOp_STORE(StoreSDNode *N, unsigned OpNo) {
SDValue Ch = N->getChain();
SDValue Ptr = N->getBasePtr();
MachineMemOperand::Flags MMOFlags = N->getMemOperand()->getFlags();
- AAMDNodes AAInfo = N->getAAInfo();
+ MMOMetadata Metadata = N->getNonRangeMMOMetadata();
SDLoc dl(N);
SDValue Lo, Hi;
@@ -6001,7 +6001,7 @@ SDValue DAGTypeLegalizer::ExpandIntOp_STORE(StoreSDNode *N, unsigned OpNo) {
GetExpandedInteger(N->getValue(), Lo, Hi);
return DAG.getTruncStore(Ch, dl, Lo, Ptr, N->getPointerInfo(),
N->getMemoryVT(), N->getBaseAlign(), MMOFlags,
- AAInfo);
+ Metadata);
}
if (DAG.getDataLayout().isLittleEndian()) {
@@ -6009,7 +6009,7 @@ SDValue DAGTypeLegalizer::ExpandIntOp_STORE(StoreSDNode *N, unsigned OpNo) {
GetExpandedInteger(N->getValue(), Lo, Hi);
Lo = DAG.getStore(Ch, dl, Lo, Ptr, N->getPointerInfo(), N->getBaseAlign(),
- MMOFlags, AAInfo);
+ MMOFlags, Metadata);
unsigned ExcessBits =
N->getMemoryVT().getSizeInBits() - NVT.getSizeInBits();
@@ -6020,7 +6020,7 @@ SDValue DAGTypeLegalizer::ExpandIntOp_STORE(StoreSDNode *N, unsigned OpNo) {
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
Hi = DAG.getTruncStore(Ch, dl, Hi, Ptr,
N->getPointerInfo().getWithOffset(IncrementSize),
- NEVT, N->getBaseAlign(), MMOFlags, AAInfo);
+ NEVT, N->getBaseAlign(), MMOFlags, Metadata);
return DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Lo, Hi);
}
@@ -6048,7 +6048,7 @@ SDValue DAGTypeLegalizer::ExpandIntOp_STORE(StoreSDNode *N, unsigned OpNo) {
// Store both the high bits and maybe some of the low bits.
Hi = DAG.getTruncStore(Ch, dl, Hi, Ptr, N->getPointerInfo(), HiVT,
- N->getBaseAlign(), MMOFlags, AAInfo);
+ N->getBaseAlign(), MMOFlags, Metadata);
// Increment the pointer to the other half.
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
@@ -6056,7 +6056,7 @@ SDValue DAGTypeLegalizer::ExpandIntOp_STORE(StoreSDNode *N, unsigned OpNo) {
Lo = DAG.getTruncStore(Ch, dl, Lo, Ptr,
N->getPointerInfo().getWithOffset(IncrementSize),
EVT::getIntegerVT(*DAG.getContext(), ExcessBits),
- N->getBaseAlign(), MMOFlags, AAInfo);
+ N->getBaseAlign(), MMOFlags, Metadata);
return DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Lo, Hi);
}
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeTypesGeneric.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeTypesGeneric.cpp
index 8c252c3491540..2e3a6495cbe05 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeTypesGeneric.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeTypesGeneric.cpp
@@ -254,19 +254,20 @@ void DAGTypeLegalizer::ExpandRes_NormalLoad(SDNode *N, SDValue &Lo,
EVT NVT = TLI.getTypeToTransformTo(*DAG.getContext(), ValueVT);
SDValue Chain = LD->getChain();
SDValue Ptr = LD->getBasePtr();
- AAMDNodes AAInfo = LD->getAAInfo();
+ MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
assert(NVT.isByteSized() && "Expanded type not byte sized!");
- Lo = DAG.getLoad(NVT, dl, Chain, Ptr, LD->getPointerInfo(),
- LD->getBaseAlign(), LD->getMemOperand()->getFlags(), AAInfo);
+ Lo =
+ DAG.getLoad(NVT, dl, Chain, Ptr, LD->getPointerInfo(), LD->getBaseAlign(),
+ LD->getMemOperand()->getFlags(), Metadata);
// Increment the pointer to the other half.
unsigned IncrementSize = NVT.getSizeInBits() / 8;
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
- Hi = DAG.getLoad(NVT, dl, Chain, Ptr,
- LD->getPointerInfo().getWithOffset(IncrementSize),
- LD->getBaseAlign(), LD->getMemOperand()->getFlags(), AAInfo);
+ Hi = DAG.getLoad(
+ NVT, dl, Chain, Ptr, LD->getPointerInfo().getWithOffset(IncrementSize),
+ LD->getBaseAlign(), LD->getMemOperand()->getFlags(), Metadata);
// Build a factor node to remember that this load is independent of the
// other one.
@@ -482,7 +483,7 @@ SDValue DAGTypeLegalizer::ExpandOp_NormalStore(SDNode *N, unsigned OpNo) {
EVT NVT = TLI.getTypeToTransformTo(*DAG.getContext(), ValueVT);
SDValue Chain = St->getChain();
SDValue Ptr = St->getBasePtr();
- AAMDNodes AAInfo = St->getAAInfo();
+ MMOMetadata Metadata = St->getNonRangeMMOMetadata();
assert(NVT.isByteSized() && "Expanded type not byte sized!");
unsigned IncrementSize = NVT.getSizeInBits() / 8;
@@ -495,12 +496,12 @@ SDValue DAGTypeLegalizer::ExpandOp_NormalStore(SDNode *N, unsigned OpNo) {
Lo =
DAG.getStore(Chain, dl, Lo, Ptr, St->getPointerInfo(), St->getBaseAlign(),
- St->getMemOperand()->getFlags(), AAInfo);
+ St->getMemOperand()->getFlags(), Metadata);
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
Hi = DAG.getStore(
Chain, dl, Hi, Ptr, St->getPointerInfo().getWithOffset(IncrementSize),
- St->getBaseAlign(), St->getMemOperand()->getFlags(), AAInfo);
+ St->getBaseAlign(), St->getMemOperand()->getFlags(), Metadata);
return DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Lo, Hi);
}
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
index dc3b589087421..8b0dccdac71fc 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
@@ -578,7 +578,8 @@ SDValue DAGTypeLegalizer::ScalarizeVecRes_LOAD(LoadSDNode *N) {
N->getValueType(0).getVectorElementType(), SDLoc(N), N->getChain(),
N->getBasePtr(), DAG.getPOISON(N->getBasePtr().getValueType()),
N->getPointerInfo(), N->getMemoryVT().getVectorElementType(),
- N->getBaseAlign(), N->getMemOperand()->getFlags(), N->getAAInfo());
+ N->getBaseAlign(), N->getMemOperand()->getFlags(),
+ N->getNonRangeMMOMetadata());
// Legalize the chain result - switch anything that used the old chain to
// use the new one.
@@ -1204,11 +1205,13 @@ SDValue DAGTypeLegalizer::ScalarizeVecOp_STORE(StoreSDNode *N, unsigned OpNo){
N->getChain(), dl, GetScalarizedVector(N->getOperand(1)),
N->getBasePtr(), N->getPointerInfo(),
N->getMemoryVT().getVectorElementType(), N->getBaseAlign(),
- N->getMemOperand()->getFlags(), N->getAAInfo());
+ N->getMemOperand()->getFlags(),
+ N->getNonRangeMMOMetadata());
- return DAG.getStore(N->getChain(), dl, GetScalarizedVector(N->getOperand(1)),
- N->getBasePtr(), N->getPointerInfo(), N->getBaseAlign(),
- N->getMemOperand()->getFlags(), N->getAAInfo());
+ return DAG.getStore(
+ N->getChain(), dl, GetScalarizedVector(N->getOperand(1)), N->getBasePtr(),
+ N->getPointerInfo(), N->getBaseAlign(), N->getMemOperand()->getFlags(),
+ N->getNonRangeMMOMetadata());
}
/// If the value to store is a vector that needs to be scalarized, it must be
@@ -2430,7 +2433,7 @@ void DAGTypeLegalizer::SplitVecRes_LOAD(LoadSDNode *LD, SDValue &Lo,
SDValue Offset = DAG.getPOISON(Ptr.getValueType());
EVT MemoryVT = LD->getMemoryVT();
MachineMemOperand::Flags MMOFlags = LD->getMemOperand()->getFlags();
- AAMDNodes AAInfo = LD->getAAInfo();
+ MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
EVT LoMemVT, HiMemVT;
std::tie(LoMemVT, HiMemVT) = DAG.GetSplitDestVTs(MemoryVT);
@@ -2445,13 +2448,13 @@ void DAGTypeLegalizer::SplitVecRes_LOAD(LoadSDNode *LD, SDValue &Lo,
Lo = DAG.getLoad(ISD::UNINDEXED, ExtType, LoVT, dl, Ch, Ptr, Offset,
LD->getPointerInfo(), LoMemVT, LD->getBaseAlign(), MMOFlags,
- AAInfo);
+ Metadata);
MachinePointerInfo MPI;
IncrementPointer(LD, LoMemVT, MPI, Ptr);
Hi = DAG.getLoad(ISD::UNINDEXED, ExtType, HiVT, dl, Ch, Ptr, Offset, MPI,
- HiMemVT, LD->getBaseAlign(), MMOFlags, AAInfo);
+ HiMemVT, LD->getBaseAlign(), MMOFlags, Metadata);
// Build a factor node to remember that this load is independent of the
// other one.
@@ -4744,7 +4747,7 @@ SDValue DAGTypeLegalizer::SplitVecOp_STORE(StoreSDNode *N, unsigned OpNo) {
EVT MemoryVT = N->getMemoryVT();
Align Alignment = N->getBaseAlign();
MachineMemOperand::Flags MMOFlags = N->getMemOperand()->getFlags();
- AAMDNodes AAInfo = N->getAAInfo();
+ MMOMetadata Metadata = N->getNonRangeMMOMetadata();
SDValue Lo, Hi;
GetSplitVector(N->getOperand(1), Lo, Hi);
@@ -4757,19 +4760,19 @@ SDValue DAGTypeLegalizer::SplitVecOp_STORE(StoreSDNode *N, unsigned OpNo) {
if (isTruncating)
Lo = DAG.getTruncStore(Ch, DL, Lo, Ptr, N->getPointerInfo(), LoMemVT,
- Alignment, MMOFlags, AAInfo);
+ Alignment, MMOFlags, Metadata);
else
Lo = DAG.getStore(Ch, DL, Lo, Ptr, N->getPointerInfo(), Alignment, MMOFlags,
- AAInfo);
+ Metadata);
MachinePointerInfo MPI;
IncrementPointer(N, LoMemVT, MPI, Ptr);
if (isTruncating)
- Hi = DAG.getTruncStore(Ch, DL, Hi, Ptr, MPI,
- HiMemVT, Alignment, MMOFlags, AAInfo);
+ Hi = DAG.getTruncStore(Ch, DL, Hi, Ptr, MPI, HiMemVT, Alignment, MMOFlags,
+ Metadata);
else
- Hi = DAG.getStore(Ch, DL, Hi, Ptr, MPI, Alignment, MMOFlags, AAInfo);
+ Hi = DAG.getStore(Ch, DL, Hi, Ptr, MPI, Alignment, MMOFlags, Metadata);
return DAG.getNode(ISD::TokenFactor, DL, MVT::Other, Lo, Hi);
}
@@ -8987,7 +8990,7 @@ SDValue DAGTypeLegalizer::GenWidenVectorLoads(SmallVectorImpl<SDValue> &LdChain,
SDValue Chain = LD->getChain();
SDValue BasePtr = LD->getBasePtr();
MachineMemOperand::Flags MMOFlags = LD->getMemOperand()->getFlags();
- AAMDNodes AAInfo = LD->getAAInfo();
+ MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
TypeSize LdWidth = LdVT.getSizeInBits();
TypeSize WidenWidth = WidenVT.getSizeInBits();
@@ -9029,7 +9032,7 @@ SDValue DAGTypeLegalizer::GenWidenVectorLoads(SmallVectorImpl<SDValue> &LdChain,
}
SDValue LdOp = DAG.getLoad(*FirstVT, dl, Chain, BasePtr, LD->getPointerInfo(),
- LD->getBaseAlign(), MMOFlags, AAInfo);
+ LD->getBaseAlign(), MMOFlags, Metadata);
LdChain.push_back(LdOp.getValue(1));
// Check if we can load the element with one instruction.
@@ -9052,8 +9055,8 @@ SDValue DAGTypeLegalizer::GenWidenVectorLoads(SmallVectorImpl<SDValue> &LdChain,
Align NewAlign = ScaledOffset == 0
? LD->getBaseAlign()
: commonAlignment(LD->getAlign(), ScaledOffset);
- SDValue L =
- DAG.getLoad(MemVT, dl, Chain, BasePtr, MPI, NewAlign, MMOFlags, AAInfo);
+ SDValue L = DAG.getLoad(MemVT, dl, Chain, BasePtr, MPI, NewAlign, MMOFlags,
+ Metadata);
LdOps.push_back(L);
LdChain.push_back(L.getValue(1));
@@ -9144,7 +9147,7 @@ DAGTypeLegalizer::GenWidenVectorExtLoads(SmallVectorImpl<SDValue> &LdChain,
SDValue Chain = LD->getChain();
SDValue BasePtr = LD->getBasePtr();
MachineMemOperand::Flags MMOFlags = LD->getMemOperand()->getFlags();
- AAMDNodes AAInfo = LD->getAAInfo();
+ MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
if (LdVT.isScalableVector())
return SDValue();
@@ -9159,7 +9162,7 @@ DAGTypeLegalizer::GenWidenVectorExtLoads(SmallVectorImpl<SDValue> &LdChain,
unsigned Increment = LdEltVT.getSizeInBits() / 8;
Ops[0] =
DAG.getExtLoad(ExtType, dl, EltVT, Chain, BasePtr, LD->getPointerInfo(),
- LdEltVT, LD->getBaseAlign(), MMOFlags, AAInfo);
+ LdEltVT, LD->getBaseAlign(), MMOFlags, Metadata);
LdChain.push_back(Ops[0].getValue(1));
unsigned i = 0, Offset = Increment;
for (i=1; i < NumElts; ++i, Offset += Increment) {
@@ -9167,7 +9170,7 @@ DAGTypeLegalizer::GenWidenVectorExtLoads(SmallVectorImpl<SDValue> &LdChain,
DAG.getObjectPtrOffset(dl, BasePtr, TypeSize::getFixed(Offset));
Ops[i] = DAG.getExtLoad(ExtType, dl, EltVT, Chain, NewBasePtr,
LD->getPointerInfo().getWithOffset(Offset), LdEltVT,
- LD->getBaseAlign(), MMOFlags, AAInfo);
+ LD->getBaseAlign(), MMOFlags, Metadata);
LdChain.push_back(Ops[i].getValue(1));
}
@@ -9187,7 +9190,7 @@ bool DAGTypeLegalizer::GenWidenVectorStores(SmallVectorImpl<SDValue> &StChain,
SDValue Chain = ST->getChain();
SDValue BasePtr = ST->getBasePtr();
MachineMemOperand::Flags MMOFlags = ST->getMemOperand()->getFlags();
- AAMDNodes AAInfo = ST->getAAInfo();
+ MMOMetadata Metadata = ST->getNonRangeMMOMetadata();
SDValue ValOp = GetWidenedVector(ST->getValue());
SDLoc dl(ST);
@@ -9239,7 +9242,7 @@ bool DAGTypeLegalizer::GenWidenVectorStores(SmallVectorImpl<SDValue> &StChain,
: commonAlignment(ST->getAlign(), ScaledOffset);
SDValue EOp = DAG.getExtractSubvector(dl, NewVT, ValOp, Idx);
SDValue PartStore = DAG.getStore(Chain, dl, EOp, BasePtr, MPI, NewAlign,
- MMOFlags, AAInfo);
+ MMOFlags, Metadata);
StChain.push_back(PartStore);
Idx += NumVTElts;
@@ -9255,8 +9258,9 @@ bool DAGTypeLegalizer::GenWidenVectorStores(SmallVectorImpl<SDValue> &StChain,
Idx = Idx * ValEltWidth / NewVTWidth.getFixedValue();
do {
SDValue EOp = DAG.getExtractVectorElt(dl, NewVT, VecOp, Idx++);
- SDValue PartStore = DAG.getStore(Chain, dl, EOp, BasePtr, MPI,
- ST->getBaseAlign(), MMOFlags, AAInfo);
+ SDValue PartStore =
+ DAG.getStore(Chain, dl, EOp, BasePtr, MPI, ST->getBaseAlign(),
+ MMOFlags, Metadata);
StChain.push_back(PartStore);
IncrementPointer(cast<StoreSDNode>(PartStore), NewVT, MPI, BasePtr);
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index 6e559d91952f2..67e14373bfc03 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -11680,10 +11680,10 @@ TargetLowering::scalarizeVectorLoad(LoadSDNode *LD,
// Load the whole vector and avoid masking off the top bits as it makes
// the codegen worse.
- SDValue Load =
- DAG.getExtLoad(ISD::EXTLOAD, SL, LoadVT, Chain, BasePTR,
- LD->getPointerInfo(), SrcIntVT, LD->getBaseAlign(),
- LD->getMemOperand()->getFlags(), LD->getAAInfo());
+ SDValue Load = DAG.getExtLoad(
+ ISD::EXTLOAD, SL, LoadVT, Chain, BasePTR, LD->getPointerInfo(),
+ SrcIntVT, LD->getBaseAlign(), LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
SmallVector<SDValue, 8> Vals;
for (unsigned Idx = 0; Idx < NumElem; ++Idx) {
@@ -11718,7 +11718,8 @@ TargetLowering::scalarizeVectorLoad(LoadSDNode *LD,
SDValue ScalarLoad = DAG.getExtLoad(
ExtType, SL, DstEltVT, Chain, BasePTR,
LD->getPointerInfo().getWithOffset(Idx * Stride), SrcEltVT,
- LD->getBaseAlign(), LD->getMemOperand()->getFlags(), LD->getAAInfo());
+ LD->getBaseAlign(), LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
BasePTR = DAG.getObjectPtrOffset(SL, BasePTR, TypeSize::getFixed(Stride));
@@ -11778,9 +11779,10 @@ SDValue TargetLowering::scalarizeVectorStore(StoreSDNode *ST,
CurrVal = DAG.getNode(ISD::OR, SL, IntVT, CurrVal, ShiftedElt);
}
- return DAG.getStore(Chain, SL, CurrVal, BasePtr, ST->getPointerInfo(),
- ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
- ST->getAAInfo());
+ return DAG.getStore(
+ Chain, SL, CurrVal, BasePtr, ST->getPointerInfo(), ST->getBaseAlign(),
+ ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata());
}
// Store Stride in bytes
@@ -11799,7 +11801,7 @@ SDValue TargetLowering::scalarizeVectorStore(StoreSDNode *ST,
SDValue Store = DAG.getTruncStore(
Chain, SL, Elt, Ptr, ST->getPointerInfo().getWithOffset(Idx * Stride),
MemSclVT, ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
- ST->getAAInfo());
+ ST->getNonRangeMMOMetadata());
Stores.push_back(Store);
}
@@ -11864,7 +11866,8 @@ TargetLowering::expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const {
// Load one integer register's worth from the original location.
SDValue Load = DAG.getLoad(
RegVT, dl, Chain, Ptr, LD->getPointerInfo().getWithOffset(Offset),
- LD->getBaseAlign(), LD->getMemOperand()->getFlags(), LD->getAAInfo());
+ LD->getBaseAlign(), LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
// Follow the load with a store to the stack slot. Remember the store.
Stores.push_back(DAG.getStore(
Load.getValue(1), dl, Load, StackPtr,
@@ -11882,7 +11885,8 @@ TargetLowering::expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const {
SDValue Load = DAG.getExtLoad(
ISD::EXTLOAD, dl, RegVT, Chain, Ptr,
LD->getPointerInfo().getWithOffset(Offset), MemVT, LD->getBaseAlign(),
- LD->getMemOperand()->getFlags(), LD->getAAInfo());
+ LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
// Follow the load with a store to the stack slot. Remember the store.
// On big-endian machines this requires a truncating store to ensure
// that the bits end up in the right place.
@@ -11923,25 +11927,29 @@ TargetLowering::expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const {
// Load the value in two parts
SDValue Lo, Hi;
if (DAG.getDataLayout().isLittleEndian()) {
- Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, VT, Chain, Ptr, LD->getPointerInfo(),
- NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
- LD->getAAInfo());
+ Lo = DAG.getExtLoad(
+ ISD::ZEXTLOAD, dl, VT, Chain, Ptr, LD->getPointerInfo(), NewLoadedVT,
+ Alignment, LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
- Hi = DAG.getExtLoad(HiExtType, dl, VT, Chain, Ptr,
- LD->getPointerInfo().getWithOffset(IncrementSize),
- NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
- LD->getAAInfo());
+ Hi = DAG.getExtLoad(
+ HiExtType, dl, VT, Chain, Ptr,
+ LD->getPointerInfo().getWithOffset(IncrementSize), NewLoadedVT,
+ Alignment, LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
} else {
- Hi = DAG.getExtLoad(HiExtType, dl, VT, Chain, Ptr, LD->getPointerInfo(),
- NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
- LD->getAAInfo());
+ Hi = DAG.getExtLoad(
+ HiExtType, dl, VT, Chain, Ptr, LD->getPointerInfo(), NewLoadedVT,
+ Alignment, LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
- Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, VT, Chain, Ptr,
- LD->getPointerInfo().getWithOffset(IncrementSize),
- NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
- LD->getAAInfo());
+ Lo = DAG.getExtLoad(
+ ISD::ZEXTLOAD, dl, VT, Chain, Ptr,
+ LD->getPointerInfo().getWithOffset(IncrementSize), NewLoadedVT,
+ Alignment, LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
}
// aggregate the two parts
@@ -11981,8 +11989,10 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
// same size, then a (misaligned) int store.
// FIXME: Does not handle truncating floating point stores!
SDValue Result = DAG.getNode(ISD::BITCAST, dl, intVT, Val);
- Result = DAG.getStore(Chain, dl, Result, Ptr, ST->getPointerInfo(),
- Alignment, ST->getMemOperand()->getFlags());
+ Result = DAG.getStore(
+ Chain, dl, Result, Ptr, ST->getPointerInfo(), Alignment,
+ ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata());
return Result;
}
// Do a (aligned) store to a stack slot, then copy from the stack slot
@@ -12021,7 +12031,8 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
Stores.push_back(DAG.getStore(Load.getValue(1), dl, Load, Ptr,
ST->getPointerInfo().getWithOffset(Offset),
ST->getBaseAlign(),
- ST->getMemOperand()->getFlags()));
+ ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata()));
// Increment the pointers.
Offset += RegBytes;
StackPtr = DAG.getObjectPtrOffset(dl, StackPtr, StackPtrIncrement);
@@ -12042,7 +12053,8 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
Stores.push_back(DAG.getTruncStore(
Load.getValue(1), dl, Load, Ptr,
ST->getPointerInfo().getWithOffset(Offset), LoadMemVT,
- ST->getBaseAlign(), ST->getMemOperand()->getFlags(), ST->getAAInfo()));
+ ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata()));
// The order of the stores doesn't matter - say it with a TokenFactor.
SDValue Result = DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Stores);
return Result;
@@ -12074,13 +12086,15 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
Store1 = DAG.getTruncStore(Chain, dl,
DAG.getDataLayout().isLittleEndian() ? Lo : Hi,
Ptr, ST->getPointerInfo(), NewStoredVT, Alignment,
- ST->getMemOperand()->getFlags());
+ ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
Store2 = DAG.getTruncStore(
Chain, dl, DAG.getDataLayout().isLittleEndian() ? Hi : Lo, Ptr,
ST->getPointerInfo().getWithOffset(IncrementSize), NewStoredVT, Alignment,
- ST->getMemOperand()->getFlags(), ST->getAAInfo());
+ ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata());
SDValue Result =
DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Store1, Store2);
@@ -14303,13 +14317,13 @@ SDValue TargetLowering::scalarizeExtractedVectorLoad(EVT ResultVT,
Load = DAG.getExtLoad(ExtType, DL, ResultVT, OriginalLoad->getChain(),
NewPtr, MPI, VecEltVT, Alignment,
OriginalLoad->getMemOperand()->getFlags(),
- OriginalLoad->getAAInfo());
+ OriginalLoad->getNonRangeMMOMetadata());
DAG.makeEquivalentMemoryOrdering(OriginalLoad, Load);
} else {
// The result type is narrower or the same width as the vector element
Load = DAG.getLoad(VecEltVT, DL, OriginalLoad->getChain(), NewPtr, MPI,
Alignment, OriginalLoad->getMemOperand()->getFlags(),
- OriginalLoad->getAAInfo());
+ OriginalLoad->getNonRangeMMOMetadata());
DAG.makeEquivalentMemoryOrdering(OriginalLoad, Load);
if (ResultVT.bitsLT(VecEltVT))
Load = DAG.getNode(ISD::TRUNCATE, DL, ResultVT, Load);
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll b/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
index 6f6d805974ce6..ac86b770dd79f 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
@@ -86,16 +86,11 @@ define i32 @test_load_l2_normal(ptr addrspace(1) %p) {
ret i32 %v
}
-;-----------------------------------------------------------------------------
-; TODO: Preserve cache hints across DAGCombiner-created memory rewrites.
-; This documents the current store-of-concat-trunc behavior: copied MMOs for
-; the split stores do not retain !mem.cache_hint metadata yet.
-;-----------------------------------------------------------------------------
-
define void @test_dagcombine_store_concat_trunc_v8i32(ptr addrspace(1) %p, <4 x i64> %a, <4 x i64> %b) {
; CHECK-LABEL: test_dagcombine_store_concat_trunc_v8i32(
-; CHECK: st.global.v4.b32 [%rd1+16], {%r8, %r7, %r6, %r5};
-; CHECK: st.global.v4.b32 [%rd1], {%r4, %r3, %r2, %r1};
+; CHECK: mov.b64 %rd10, 12345;
+; CHECK: st.global.L2::cache_hint.v4.b32 [%rd1+16], {%r8, %r7, %r6, %r5}, %rd10;
+; CHECK: st.global.L2::cache_hint.v4.b32 [%rd1], {%r4, %r3, %r2, %r1}, %rd10;
%ta = trunc <4 x i64> %a to <4 x i32>
%tb = trunc <4 x i64> %b to <4 x i32>
%c = shufflevector <4 x i32> %ta, <4 x i32> %tb, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7>
@@ -104,195 +99,197 @@ define void @test_dagcombine_store_concat_trunc_v8i32(ptr addrspace(1) %p, <4 x
}
;-----------------------------------------------------------------------------
-; TODO: Preserve cache hints across one-to-N DAG memory rewrites.
-; These tests document the current split/scalarized behavior: newly-created
-; memory ops do not retain !mem.cache_hint metadata yet.
+; Preserve cache hints across one-to-N DAG memory rewrites.
;-----------------------------------------------------------------------------
define <16 x i32> @test_legalize_split_load_v16i32(ptr addrspace(1) %p) {
; O2-LABEL: test_legalize_split_load_v16i32(
-; O2: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
-; O2: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];
-; O2: ld.global.v4.b32 {%r9, %r10, %r11, %r12}, [%rd1+32];
-; O2: ld.global.v4.b32 {%r13, %r14, %r15, %r16}, [%rd1+48];
+; O2: mov.b64 %rd2, 12345;
+; O2: ld.global.L2::cache_hint.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1], %rd2;
+; O2: ld.global.L2::cache_hint.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16], %rd2;
+; O2: ld.global.L2::cache_hint.v4.b32 {%r9, %r10, %r11, %r12}, [%rd1+32], %rd2;
+; O2: ld.global.L2::cache_hint.v4.b32 {%r13, %r14, %r15, %r16}, [%rd1+48], %rd2;
;
; O0-LABEL: test_legalize_split_load_v16i32(
-; O0: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+48];
-; O0: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+32];
-; O0: ld.global.v4.b32 {%r9, %r10, %r11, %r12}, [%rd1+16];
-; O0: ld.global.v4.b32 {%r13, %r14, %r15, %r16}, [%rd1];
+; O0: mov.b64 %rd2, 12345;
+; O0: ld.global.L2::cache_hint.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+48], %rd2;
+; O0: ld.global.L2::cache_hint.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+32], %rd2;
+; O0: ld.global.L2::cache_hint.v4.b32 {%r9, %r10, %r11, %r12}, [%rd1+16], %rd2;
+; O0: ld.global.L2::cache_hint.v4.b32 {%r13, %r14, %r15, %r16}, [%rd1], %rd2;
%v = load <16 x i32>, ptr addrspace(1) %p, align 16, !mem.cache_hint !11
ret <16 x i32> %v
}
define void @test_legalize_split_store_v16i32(ptr addrspace(1) %p, <16 x i32> %v) {
; O2-LABEL: test_legalize_split_store_v16i32(
-; O2: st.global.v4.b32 [%rd1+48], {%r1, %r2, %r3, %r4};
-; O2: st.global.v4.b32 [%rd1+32], {%r5, %r6, %r7, %r8};
-; O2: st.global.v4.b32 [%rd1+16], {%r9, %r10, %r11, %r12};
-; O2: st.global.v4.b32 [%rd1], {%r13, %r14, %r15, %r16};
+; O2: mov.b64 %rd2, 12345;
+; O2: st.global.L2::cache_hint.v4.b32 [%rd1+48], {%r1, %r2, %r3, %r4}, %rd2;
+; O2: st.global.L2::cache_hint.v4.b32 [%rd1+32], {%r5, %r6, %r7, %r8}, %rd2;
+; O2: st.global.L2::cache_hint.v4.b32 [%rd1+16], {%r9, %r10, %r11, %r12}, %rd2;
+; O2: st.global.L2::cache_hint.v4.b32 [%rd1], {%r13, %r14, %r15, %r16}, %rd2;
;
; O0-LABEL: test_legalize_split_store_v16i32(
-; O0: st.global.v4.b32 [%rd1+48], {%r13, %r14, %r15, %r16};
-; O0: st.global.v4.b32 [%rd1+32], {%r9, %r10, %r11, %r12};
-; O0: st.global.v4.b32 [%rd1+16], {%r5, %r6, %r7, %r8};
-; O0: st.global.v4.b32 [%rd1], {%r1, %r2, %r3, %r4};
+; O0: mov.b64 %rd2, 12345;
+; O0: st.global.L2::cache_hint.v4.b32 [%rd1+48], {%r13, %r14, %r15, %r16}, %rd2;
+; O0: st.global.L2::cache_hint.v4.b32 [%rd1+32], {%r9, %r10, %r11, %r12}, %rd2;
+; O0: st.global.L2::cache_hint.v4.b32 [%rd1+16], {%r5, %r6, %r7, %r8}, %rd2;
+; O0: st.global.L2::cache_hint.v4.b32 [%rd1], {%r1, %r2, %r3, %r4}, %rd2;
store <16 x i32> %v, ptr addrspace(1) %p, align 16, !mem.cache_hint !12
ret void
}
define <3 x i64> @test_legalize_scalarize_load_v3i64(ptr addrspace(1) %p) {
; CHECK-LABEL: test_legalize_scalarize_load_v3i64(
-; CHECK: ld.global.b64 %rd2, [%rd1+16];
-; CHECK: ld.global.b64 %rd3, [%rd1+8];
-; CHECK: ld.global.b64 %rd4, [%rd1];
+; CHECK: ld.global.L1::evict_last.b64 %rd2, [%rd1+16];
+; CHECK: ld.global.L1::evict_last.b64 %rd3, [%rd1+8];
+; CHECK: ld.global.L1::evict_last.b64 %rd4, [%rd1];
%v = load <3 x i64>, ptr addrspace(1) %p, align 8, !mem.cache_hint !13
ret <3 x i64> %v
}
define void @test_legalize_scalarize_store_v3i64(ptr addrspace(1) %p, <3 x i64> %v) {
; O2-LABEL: test_legalize_scalarize_store_v3i64(
-; O2: st.global.b64 [%rd1+16], %rd2;
-; O2: st.global.b64 [%rd1+8], %rd4;
-; O2: st.global.b64 [%rd1], %rd3;
+; O2: st.global.L1::evict_first.b64 [%rd1+16], %rd2;
+; O2: st.global.L1::evict_first.b64 [%rd1+8], %rd4;
+; O2: st.global.L1::evict_first.b64 [%rd1], %rd3;
;
; O0-LABEL: test_legalize_scalarize_store_v3i64(
-; O0: st.global.b64 [%rd1+16], %rd4;
-; O0: st.global.b64 [%rd1+8], %rd3;
-; O0: st.global.b64 [%rd1], %rd2;
+; O0: st.global.L1::evict_first.b64 [%rd1+16], %rd4;
+; O0: st.global.L1::evict_first.b64 [%rd1+8], %rd3;
+; O0: st.global.L1::evict_first.b64 [%rd1], %rd2;
store <3 x i64> %v, ptr addrspace(1) %p, align 8, !mem.cache_hint !14
ret void
}
define void @legalize_unaligned_f64_store(ptr addrspace(1) %p, double %v) {
; O2-LABEL: legalize_unaligned_f64_store(
-; O2: st.global.b8 [%rd1], %rd2;
-; O2: st.global.b8 [%rd1+7], %rd3;
-; O2: st.global.b8 [%rd1+6], %rd4;
-; O2: st.global.b8 [%rd1+5], %rd5;
-; O2: st.global.b8 [%rd1+4], %rd6;
-; O2: st.global.b8 [%rd1+3], %rd7;
-; O2: st.global.b8 [%rd1+2], %rd8;
-; O2: st.global.b8 [%rd1+1], %rd9;
+; O2: st.global.L1::evict_first.b8 [%rd1], %rd2;
+; O2: st.global.L1::evict_first.b8 [%rd1+7], %rd3;
+; O2: st.global.L1::evict_first.b8 [%rd1+6], %rd4;
+; O2: st.global.L1::evict_first.b8 [%rd1+5], %rd5;
+; O2: st.global.L1::evict_first.b8 [%rd1+4], %rd6;
+; O2: st.global.L1::evict_first.b8 [%rd1+3], %rd7;
+; O2: st.global.L1::evict_first.b8 [%rd1+2], %rd8;
+; O2: st.global.L1::evict_first.b8 [%rd1+1], %rd9;
;
; O0-LABEL: legalize_unaligned_f64_store(
-; O0: st.global.b8 [%rd1+7], %rd3;
-; O0: st.global.b8 [%rd1+6], %rd4;
-; O0: st.global.b8 [%rd1+5], %rd5;
-; O0: st.global.b8 [%rd1+4], %rd6;
-; O0: st.global.b8 [%rd1+3], %rd7;
-; O0: st.global.b8 [%rd1+2], %rd8;
-; O0: st.global.b8 [%rd1], %rd2;
-; O0: st.global.b8 [%rd1+1], %rd9;
+; O0: st.global.L1::evict_first.b8 [%rd1+7], %rd3;
+; O0: st.global.L1::evict_first.b8 [%rd1+6], %rd4;
+; O0: st.global.L1::evict_first.b8 [%rd1+5], %rd5;
+; O0: st.global.L1::evict_first.b8 [%rd1+4], %rd6;
+; O0: st.global.L1::evict_first.b8 [%rd1+3], %rd7;
+; O0: st.global.L1::evict_first.b8 [%rd1+2], %rd8;
+; O0: st.global.L1::evict_first.b8 [%rd1], %rd2;
+; O0: st.global.L1::evict_first.b8 [%rd1+1], %rd9;
store double %v, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
ret void
}
define void @legalize_unaligned_v3i64_store(ptr addrspace(1) %p, <3 x i64> %v) {
; O2-LABEL: legalize_unaligned_v3i64_store(
-; O2: st.global.b8 [%rd1+16], %rd2;
-; O2: st.global.b8 [%rd1+8], %rd4;
-; O2: st.global.b8 [%rd1], %rd3;
-; O2: st.global.b8 [%rd1+23], %rd5;
-; O2: st.global.b8 [%rd1+22], %rd6;
-; O2: st.global.b8 [%rd1+21], %rd7;
-; O2: st.global.b8 [%rd1+20], %rd8;
-; O2: st.global.b8 [%rd1+19], %rd9;
-; O2: st.global.b8 [%rd1+18], %rd10;
-; O2: st.global.b8 [%rd1+17], %rd11;
-; O2: st.global.b8 [%rd1+15], %rd12;
-; O2: st.global.b8 [%rd1+14], %rd13;
-; O2: st.global.b8 [%rd1+13], %rd14;
-; O2: st.global.b8 [%rd1+12], %rd15;
-; O2: st.global.b8 [%rd1+11], %rd16;
-; O2: st.global.b8 [%rd1+10], %rd17;
-; O2: st.global.b8 [%rd1+9], %rd18;
-; O2: st.global.b8 [%rd1+7], %rd19;
-; O2: st.global.b8 [%rd1+6], %rd20;
-; O2: st.global.b8 [%rd1+5], %rd21;
-; O2: st.global.b8 [%rd1+4], %rd22;
-; O2: st.global.b8 [%rd1+3], %rd23;
-; O2: st.global.b8 [%rd1+2], %rd24;
-; O2: st.global.b8 [%rd1+1], %rd25;
+; O2: st.global.L1::evict_first.b8 [%rd1+16], %rd2;
+; O2: st.global.L1::evict_first.b8 [%rd1+8], %rd4;
+; O2: st.global.L1::evict_first.b8 [%rd1], %rd3;
+; O2: st.global.L1::evict_first.b8 [%rd1+23], %rd5;
+; O2: st.global.L1::evict_first.b8 [%rd1+22], %rd6;
+; O2: st.global.L1::evict_first.b8 [%rd1+21], %rd7;
+; O2: st.global.L1::evict_first.b8 [%rd1+20], %rd8;
+; O2: st.global.L1::evict_first.b8 [%rd1+19], %rd9;
+; O2: st.global.L1::evict_first.b8 [%rd1+18], %rd10;
+; O2: st.global.L1::evict_first.b8 [%rd1+17], %rd11;
+; O2: st.global.L1::evict_first.b8 [%rd1+15], %rd12;
+; O2: st.global.L1::evict_first.b8 [%rd1+14], %rd13;
+; O2: st.global.L1::evict_first.b8 [%rd1+13], %rd14;
+; O2: st.global.L1::evict_first.b8 [%rd1+12], %rd15;
+; O2: st.global.L1::evict_first.b8 [%rd1+11], %rd16;
+; O2: st.global.L1::evict_first.b8 [%rd1+10], %rd17;
+; O2: st.global.L1::evict_first.b8 [%rd1+9], %rd18;
+; O2: st.global.L1::evict_first.b8 [%rd1+7], %rd19;
+; O2: st.global.L1::evict_first.b8 [%rd1+6], %rd20;
+; O2: st.global.L1::evict_first.b8 [%rd1+5], %rd21;
+; O2: st.global.L1::evict_first.b8 [%rd1+4], %rd22;
+; O2: st.global.L1::evict_first.b8 [%rd1+3], %rd23;
+; O2: st.global.L1::evict_first.b8 [%rd1+2], %rd24;
+; O2: st.global.L1::evict_first.b8 [%rd1+1], %rd25;
;
; O0-LABEL: legalize_unaligned_v3i64_store(
-; O0: st.global.b8 [%rd1+23], %rd5;
-; O0: st.global.b8 [%rd1+22], %rd6;
-; O0: st.global.b8 [%rd1+21], %rd7;
-; O0: st.global.b8 [%rd1+20], %rd8;
-; O0: st.global.b8 [%rd1+19], %rd9;
-; O0: st.global.b8 [%rd1+18], %rd10;
-; O0: st.global.b8 [%rd1+16], %rd4;
-; O0: st.global.b8 [%rd1+17], %rd11;
-; O0: st.global.b8 [%rd1+15], %rd12;
-; O0: st.global.b8 [%rd1+14], %rd13;
-; O0: st.global.b8 [%rd1+13], %rd14;
-; O0: st.global.b8 [%rd1+12], %rd15;
-; O0: st.global.b8 [%rd1+11], %rd16;
-; O0: st.global.b8 [%rd1+10], %rd17;
-; O0: st.global.b8 [%rd1+8], %rd3;
-; O0: st.global.b8 [%rd1+9], %rd18;
-; O0: st.global.b8 [%rd1+7], %rd19;
-; O0: st.global.b8 [%rd1+6], %rd20;
-; O0: st.global.b8 [%rd1+5], %rd21;
-; O0: st.global.b8 [%rd1+4], %rd22;
-; O0: st.global.b8 [%rd1+3], %rd23;
-; O0: st.global.b8 [%rd1+2], %rd24;
-; O0: st.global.b8 [%rd1], %rd2;
-; O0: st.global.b8 [%rd1+1], %rd25;
+; O0: st.global.L1::evict_first.b8 [%rd1+23], %rd5;
+; O0: st.global.L1::evict_first.b8 [%rd1+22], %rd6;
+; O0: st.global.L1::evict_first.b8 [%rd1+21], %rd7;
+; O0: st.global.L1::evict_first.b8 [%rd1+20], %rd8;
+; O0: st.global.L1::evict_first.b8 [%rd1+19], %rd9;
+; O0: st.global.L1::evict_first.b8 [%rd1+18], %rd10;
+; O0: st.global.L1::evict_first.b8 [%rd1+16], %rd4;
+; O0: st.global.L1::evict_first.b8 [%rd1+17], %rd11;
+; O0: st.global.L1::evict_first.b8 [%rd1+15], %rd12;
+; O0: st.global.L1::evict_first.b8 [%rd1+14], %rd13;
+; O0: st.global.L1::evict_first.b8 [%rd1+13], %rd14;
+; O0: st.global.L1::evict_first.b8 [%rd1+12], %rd15;
+; O0: st.global.L1::evict_first.b8 [%rd1+11], %rd16;
+; O0: st.global.L1::evict_first.b8 [%rd1+10], %rd17;
+; O0: st.global.L1::evict_first.b8 [%rd1+8], %rd3;
+; O0: st.global.L1::evict_first.b8 [%rd1+9], %rd18;
+; O0: st.global.L1::evict_first.b8 [%rd1+7], %rd19;
+; O0: st.global.L1::evict_first.b8 [%rd1+6], %rd20;
+; O0: st.global.L1::evict_first.b8 [%rd1+5], %rd21;
+; O0: st.global.L1::evict_first.b8 [%rd1+4], %rd22;
+; O0: st.global.L1::evict_first.b8 [%rd1+3], %rd23;
+; O0: st.global.L1::evict_first.b8 [%rd1+2], %rd24;
+; O0: st.global.L1::evict_first.b8 [%rd1], %rd2;
+; O0: st.global.L1::evict_first.b8 [%rd1+1], %rd25;
store <3 x i64> %v, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
ret void
}
define double @legalize_unaligned_f64_load(ptr addrspace(1) %p) {
; CHECK-LABEL: legalize_unaligned_f64_load(
-; CHECK: ld.global.b8 %rd2, [%rd1];
-; CHECK: ld.global.b8 %rd3, [%rd1+1];
-; CHECK: ld.global.b8 %rd6, [%rd1+2];
-; CHECK: ld.global.b8 %rd8, [%rd1+3];
-; CHECK: ld.global.b8 %rd12, [%rd1+4];
-; CHECK: ld.global.b8 %rd13, [%rd1+5];
-; CHECK: ld.global.b8 %rd16, [%rd1+6];
-; CHECK: ld.global.b8 %rd18, [%rd1+7];
+; CHECK: ld.global.L1::evict_last.b8 %rd2, [%rd1];
+; CHECK: ld.global.L1::evict_last.b8 %rd3, [%rd1+1];
+; CHECK: ld.global.L1::evict_last.b8 %rd6, [%rd1+2];
+; CHECK: ld.global.L1::evict_last.b8 %rd8, [%rd1+3];
+; CHECK: ld.global.L1::evict_last.b8 %rd12, [%rd1+4];
+; CHECK: ld.global.L1::evict_last.b8 %rd13, [%rd1+5];
+; CHECK: ld.global.L1::evict_last.b8 %rd16, [%rd1+6];
+; CHECK: ld.global.L1::evict_last.b8 %rd18, [%rd1+7];
%v = load double, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
ret double %v
}
define <3 x i64> @legalize_unaligned_v3i64_load(ptr addrspace(1) %p) {
; CHECK-LABEL: legalize_unaligned_v3i64_load(
-; CHECK: ld.global.b8 %rd2, [%rd1+16];
-; CHECK: ld.global.b8 %rd3, [%rd1+17];
-; CHECK: ld.global.b8 %rd6, [%rd1+18];
-; CHECK: ld.global.b8 %rd8, [%rd1+19];
-; CHECK: ld.global.b8 %rd12, [%rd1+20];
-; CHECK: ld.global.b8 %rd13, [%rd1+21];
-; CHECK: ld.global.b8 %rd16, [%rd1+22];
-; CHECK: ld.global.b8 %rd18, [%rd1+23];
-; CHECK: ld.global.b8 %rd24, [%rd1+8];
-; CHECK: ld.global.b8 %rd25, [%rd1+9];
-; CHECK: ld.global.b8 %rd28, [%rd1+10];
-; CHECK: ld.global.b8 %rd30, [%rd1+11];
-; CHECK: ld.global.b8 %rd34, [%rd1+12];
-; CHECK: ld.global.b8 %rd35, [%rd1+13];
-; CHECK: ld.global.b8 %rd38, [%rd1+14];
-; CHECK: ld.global.b8 %rd40, [%rd1+15];
-; CHECK: ld.global.b8 %rd46, [%rd1];
-; CHECK: ld.global.b8 %rd47, [%rd1+1];
-; CHECK: ld.global.b8 %rd50, [%rd1+2];
-; CHECK: ld.global.b8 %rd52, [%rd1+3];
-; CHECK: ld.global.b8 %rd56, [%rd1+4];
-; CHECK: ld.global.b8 %rd57, [%rd1+5];
-; CHECK: ld.global.b8 %rd60, [%rd1+6];
-; CHECK: ld.global.b8 %rd62, [%rd1+7];
+; CHECK: ld.global.L1::evict_last.b8 %rd2, [%rd1+16];
+; CHECK: ld.global.L1::evict_last.b8 %rd3, [%rd1+17];
+; CHECK: ld.global.L1::evict_last.b8 %rd6, [%rd1+18];
+; CHECK: ld.global.L1::evict_last.b8 %rd8, [%rd1+19];
+; CHECK: ld.global.L1::evict_last.b8 %rd12, [%rd1+20];
+; CHECK: ld.global.L1::evict_last.b8 %rd13, [%rd1+21];
+; CHECK: ld.global.L1::evict_last.b8 %rd16, [%rd1+22];
+; CHECK: ld.global.L1::evict_last.b8 %rd18, [%rd1+23];
+; CHECK: ld.global.L1::evict_last.b8 %rd24, [%rd1+8];
+; CHECK: ld.global.L1::evict_last.b8 %rd25, [%rd1+9];
+; CHECK: ld.global.L1::evict_last.b8 %rd28, [%rd1+10];
+; CHECK: ld.global.L1::evict_last.b8 %rd30, [%rd1+11];
+; CHECK: ld.global.L1::evict_last.b8 %rd34, [%rd1+12];
+; CHECK: ld.global.L1::evict_last.b8 %rd35, [%rd1+13];
+; CHECK: ld.global.L1::evict_last.b8 %rd38, [%rd1+14];
+; CHECK: ld.global.L1::evict_last.b8 %rd40, [%rd1+15];
+; CHECK: ld.global.L1::evict_last.b8 %rd46, [%rd1];
+; CHECK: ld.global.L1::evict_last.b8 %rd47, [%rd1+1];
+; CHECK: ld.global.L1::evict_last.b8 %rd50, [%rd1+2];
+; CHECK: ld.global.L1::evict_last.b8 %rd52, [%rd1+3];
+; CHECK: ld.global.L1::evict_last.b8 %rd56, [%rd1+4];
+; CHECK: ld.global.L1::evict_last.b8 %rd57, [%rd1+5];
+; CHECK: ld.global.L1::evict_last.b8 %rd60, [%rd1+6];
+; CHECK: ld.global.L1::evict_last.b8 %rd62, [%rd1+7];
%v = load <3 x i64>, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
ret <3 x i64> %v
}
define i32 @legalize_i24_load(ptr addrspace(1) %p) {
; CHECK-LABEL: legalize_i24_load(
-; CHECK: ld.global.b8 %r1, [%rd1];
-; CHECK: ld.global.b8 %r2, [%rd1+1];
-; CHECK: ld.global.b8 %r5, [%rd1+2];
+; CHECK: ld.global.L1::evict_last.b8 %r1, [%rd1];
+; CHECK: ld.global.L1::evict_last.b8 %r2, [%rd1+1];
+; CHECK: ld.global.L1::evict_last.b8 %r5, [%rd1+2];
%v = load i24, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
%ext = zext i24 %v to i32
ret i32 %ext
@@ -301,9 +298,9 @@ define i32 @legalize_i24_load(ptr addrspace(1) %p) {
define void @legalize_i24_store(ptr addrspace(1) %p, i32 %v) {
; CHECK-LABEL: legalize_i24_store(
; CHECK: ld.param.b32 %r1, [legalize_i24_store_param_1];
-; CHECK: st.global.b8 [%rd1], %r1;
-; CHECK: st.global.b8 [%rd1+1], %r2;
-; CHECK: st.global.b8 [%rd1+2], %r3;
+; CHECK: st.global.L1::evict_first.b8 [%rd1], %r1;
+; CHECK: st.global.L1::evict_first.b8 [%rd1+1], %r2;
+; CHECK: st.global.L1::evict_first.b8 [%rd1+2], %r3;
%trunc = trunc i32 %v to i24
store i24 %trunc, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
ret void
@@ -311,40 +308,40 @@ define void @legalize_i24_store(ptr addrspace(1) %p, i32 %v) {
define <1 x i64> @legalize_scalarize_load_v1i64(ptr addrspace(1) %p) {
; CHECK-LABEL: legalize_scalarize_load_v1i64(
-; CHECK: ld.global.b64 %rd2, [%rd1];
+; CHECK: ld.global.L1::evict_last.b64 %rd2, [%rd1];
%v = load <1 x i64>, ptr addrspace(1) %p, align 8, !mem.cache_hint !13
ret <1 x i64> %v
}
define void @legalize_scalarize_store_v1i64(ptr addrspace(1) %p, <1 x i64> %v) {
; CHECK-LABEL: legalize_scalarize_store_v1i64(
-; CHECK: st.global.b64 [%rd1], %rd2;
+; CHECK: st.global.L1::evict_first.b64 [%rd1], %rd2;
store <1 x i64> %v, ptr addrspace(1) %p, align 8, !mem.cache_hint !14
ret void
}
define <3 x i1> @legalize_scalarize_load_v3i1(ptr addrspace(1) %p) {
; CHECK-LABEL: legalize_scalarize_load_v3i1(
-; CHECK: ld.global.b8 %rs1, [%rd1];
+; CHECK: ld.global.L1::evict_last.b8 %rs1, [%rd1];
%v = load <3 x i1>, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
ret <3 x i1> %v
}
define void @legalize_scalarize_store_v3i1(ptr addrspace(1) %p, <3 x i1> %v) {
; O2-LABEL: legalize_scalarize_store_v3i1(
-; O2: st.global.b8 [%rd1], %rs10;
+; O2: st.global.L1::evict_first.b8 [%rd1], %rs10;
;
; O0-LABEL: legalize_scalarize_store_v3i1(
-; O0: st.global.b8 [%rd1], %rs11;
+; O0: st.global.L1::evict_first.b8 [%rd1], %rs11;
store <3 x i1> %v, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
ret void
}
define <3 x i32> @legalize_widen_extload_v3i8(ptr addrspace(1) %p) {
; CHECK-LABEL: legalize_widen_extload_v3i8(
-; CHECK: ld.global.b8 %rs1, [%rd1];
-; CHECK: ld.global.b8 %rs2, [%rd1+1];
-; CHECK: ld.global.b8 %r4, [%rd1+2];
+; CHECK: ld.global.L1::evict_last.b8 %rs1, [%rd1];
+; CHECK: ld.global.L1::evict_last.b8 %rs2, [%rd1+1];
+; CHECK: ld.global.L1::evict_last.b8 %r4, [%rd1+2];
%v = load <3 x i8>, ptr addrspace(1) %p, align 1, !mem.cache_hint !13
%ext = zext <3 x i8> %v to <3 x i32>
ret <3 x i32> %ext
@@ -352,15 +349,15 @@ define <3 x i32> @legalize_widen_extload_v3i8(ptr addrspace(1) %p) {
define void @legalize_widen_truncstore_v3i8(ptr addrspace(1) %p, <3 x i32> %v) {
; O2-LABEL: legalize_widen_truncstore_v3i8(
-; O2: st.global.b8 [%rd1+2], %rs1;
-; O2: st.global.b8 [%rd1], %rs2;
-; O2: st.global.b8 [%rd1+1], %rs3;
+; O2: st.global.L1::evict_first.b8 [%rd1+2], %rs1;
+; O2: st.global.L1::evict_first.b8 [%rd1], %rs2;
+; O2: st.global.L1::evict_first.b8 [%rd1+1], %rs3;
;
; O0-LABEL: legalize_widen_truncstore_v3i8(
; O0: ld.param.b32 %r3, [legalize_widen_truncstore_v3i8_param_1+8];
-; O0: st.global.b8 [%rd1], %rs1;
-; O0: st.global.b8 [%rd1+1], %rs2;
-; O0: st.global.b8 [%rd1+2], %r3;
+; O0: st.global.L1::evict_first.b8 [%rd1], %rs1;
+; O0: st.global.L1::evict_first.b8 [%rd1+1], %rs2;
+; O0: st.global.L1::evict_first.b8 [%rd1+2], %r3;
%trunc = trunc <3 x i32> %v to <3 x i8>
store <3 x i8> %trunc, ptr addrspace(1) %p, align 1, !mem.cache_hint !14
ret void
@@ -368,7 +365,7 @@ define void @legalize_widen_truncstore_v3i8(ptr addrspace(1) %p, <3 x i32> %v) {
define i32 @scalarize_extracted_vector_load(ptr addrspace(1) %p, i64 %idx) {
; CHECK-LABEL: scalarize_extracted_vector_load(
-; CHECK: ld.global.b32 %r1, [%rd5];
+; CHECK: ld.global.L1::evict_last.b32 %r1, [%rd5];
%v = load <4 x i32>, ptr addrspace(1) %p, align 16, !mem.cache_hint !13
%elt = extractelement <4 x i32> %v, i64 %idx
ret i32 %elt
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll b/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
index 0b354dfb13787..cc92dbc6b8d58 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
@@ -7,21 +7,21 @@ define void @unaligned_f64_store(ptr addrspace(1) %p, double %v) {
; CHECK: bb.0 (%ir-block.0):
; CHECK-NEXT: [[LD_i64_:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_f64_store_param_0>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
; CHECK-NEXT: [[LD_i64_1:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_f64_store_param_1>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
- ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 0, 0, $noreg :: (store (s8) into %ir.p, addrspace 1)
+ ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 0, 2, $noreg :: (store (s8) into %ir.p, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 56
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 7, 0, $noreg :: (store (s8) into %ir.p + 7, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 7, 2, $noreg :: (store (s8) into %ir.p + 7, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri1:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 48
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 6, 0, $noreg :: (store (s8) into %ir.p + 6, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 6, 2, $noreg :: (store (s8) into %ir.p + 6, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri2:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 40
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 5, 0, $noreg :: (store (s8) into %ir.p + 5, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 5, 2, $noreg :: (store (s8) into %ir.p + 5, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri3:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 32
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 4, 0, $noreg :: (store (s8) into %ir.p + 4, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 4, 2, $noreg :: (store (s8) into %ir.p + 4, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri4:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 24
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 3, 0, $noreg :: (store (s8) into %ir.p + 3, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 3, 2, $noreg :: (store (s8) into %ir.p + 3, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri5:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 16
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 2, 0, $noreg :: (store (s8) into %ir.p + 2, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 2, 2, $noreg :: (store (s8) into %ir.p + 2, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri6:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 8
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 1, 0, $noreg :: (store (s8) into %ir.p + 1, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 1, 2, $noreg :: (store (s8) into %ir.p + 1, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: Return
store double %v, ptr addrspace(1) %p, align 1, !alias.scope !2, !noalias !4, !mem.cache_hint !5
ret void
@@ -32,52 +32,52 @@ define void @unaligned_v3i64_store(ptr addrspace(1) %p, <3 x i64> %v) {
; CHECK: bb.0 (%ir-block.0):
; CHECK-NEXT: [[LD_i64_:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_v3i64_store_param_0>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
; CHECK-NEXT: [[LD_i64_1:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_v3i64_store_param_1>, 16, 0, $noreg :: (dereferenceable invariant load (s64), align 16, addrspace 102)
- ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 16, 0, $noreg :: (store (s8) into %ir.p + 16, addrspace 1)
+ ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 16, 2, $noreg :: (store (s8) into %ir.p + 16, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[LDV_i64_v2_:%[0-9]+]]:b64, [[LDV_i64_v2_1:%[0-9]+]]:b64 = LDV_i64_v2 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_v3i64_store_param_1>, 0, 0, $noreg :: (dereferenceable invariant load (s128), align 32, addrspace 102)
- ; CHECK-NEXT: ST_i64 [[LDV_i64_v2_1]], 0, 0, 1, 8, [[LD_i64_]], 8, 0, $noreg :: (store (s8) into %ir.p + 8, addrspace 1)
- ; CHECK-NEXT: ST_i64 [[LDV_i64_v2_]], 0, 0, 1, 8, [[LD_i64_]], 0, 0, $noreg :: (store (s8) into %ir.p, addrspace 1)
+ ; CHECK-NEXT: ST_i64 [[LDV_i64_v2_1]], 0, 0, 1, 8, [[LD_i64_]], 8, 2, $noreg :: (store (s8) into %ir.p + 8, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
+ ; CHECK-NEXT: ST_i64 [[LDV_i64_v2_]], 0, 0, 1, 8, [[LD_i64_]], 0, 2, $noreg :: (store (s8) into %ir.p, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 56
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 23, 0, $noreg :: (store (s8) into %ir.p + 23, !alias.scope !0, !noalias !3, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 23, 2, $noreg :: (store (s8) into %ir.p + 23, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri1:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 48
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 22, 0, $noreg :: (store (s8) into %ir.p + 22, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 22, 2, $noreg :: (store (s8) into %ir.p + 22, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri2:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 40
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 21, 0, $noreg :: (store (s8) into %ir.p + 21, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 21, 2, $noreg :: (store (s8) into %ir.p + 21, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri3:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 32
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 20, 0, $noreg :: (store (s8) into %ir.p + 20, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 20, 2, $noreg :: (store (s8) into %ir.p + 20, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri4:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 24
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 19, 0, $noreg :: (store (s8) into %ir.p + 19, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 19, 2, $noreg :: (store (s8) into %ir.p + 19, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri5:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 16
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 18, 0, $noreg :: (store (s8) into %ir.p + 18, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 18, 2, $noreg :: (store (s8) into %ir.p + 18, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri6:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 8
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 17, 0, $noreg :: (store (s8) into %ir.p + 17, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 17, 2, $noreg :: (store (s8) into %ir.p + 17, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri7:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 56
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri7]], 0, 0, 1, 8, [[LD_i64_]], 15, 0, $noreg :: (store (s8) into %ir.p + 15, !alias.scope !0, !noalias !3, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri7]], 0, 0, 1, 8, [[LD_i64_]], 15, 2, $noreg :: (store (s8) into %ir.p + 15, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri8:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 48
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri8]], 0, 0, 1, 8, [[LD_i64_]], 14, 0, $noreg :: (store (s8) into %ir.p + 14, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri8]], 0, 0, 1, 8, [[LD_i64_]], 14, 2, $noreg :: (store (s8) into %ir.p + 14, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri9:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 40
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri9]], 0, 0, 1, 8, [[LD_i64_]], 13, 0, $noreg :: (store (s8) into %ir.p + 13, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri9]], 0, 0, 1, 8, [[LD_i64_]], 13, 2, $noreg :: (store (s8) into %ir.p + 13, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri10:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 32
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri10]], 0, 0, 1, 8, [[LD_i64_]], 12, 0, $noreg :: (store (s8) into %ir.p + 12, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri10]], 0, 0, 1, 8, [[LD_i64_]], 12, 2, $noreg :: (store (s8) into %ir.p + 12, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri11:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 24
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri11]], 0, 0, 1, 8, [[LD_i64_]], 11, 0, $noreg :: (store (s8) into %ir.p + 11, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri11]], 0, 0, 1, 8, [[LD_i64_]], 11, 2, $noreg :: (store (s8) into %ir.p + 11, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri12:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 16
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri12]], 0, 0, 1, 8, [[LD_i64_]], 10, 0, $noreg :: (store (s8) into %ir.p + 10, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri12]], 0, 0, 1, 8, [[LD_i64_]], 10, 2, $noreg :: (store (s8) into %ir.p + 10, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri13:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_1]], 8
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri13]], 0, 0, 1, 8, [[LD_i64_]], 9, 0, $noreg :: (store (s8) into %ir.p + 9, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri13]], 0, 0, 1, 8, [[LD_i64_]], 9, 2, $noreg :: (store (s8) into %ir.p + 9, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri14:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 56
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri14]], 0, 0, 1, 8, [[LD_i64_]], 7, 0, $noreg :: (store (s8) into %ir.p + 7, !alias.scope !0, !noalias !3, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri14]], 0, 0, 1, 8, [[LD_i64_]], 7, 2, $noreg :: (store (s8) into %ir.p + 7, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri15:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 48
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri15]], 0, 0, 1, 8, [[LD_i64_]], 6, 0, $noreg :: (store (s8) into %ir.p + 6, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri15]], 0, 0, 1, 8, [[LD_i64_]], 6, 2, $noreg :: (store (s8) into %ir.p + 6, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri16:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 40
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri16]], 0, 0, 1, 8, [[LD_i64_]], 5, 0, $noreg :: (store (s8) into %ir.p + 5, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri16]], 0, 0, 1, 8, [[LD_i64_]], 5, 2, $noreg :: (store (s8) into %ir.p + 5, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri17:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 32
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri17]], 0, 0, 1, 8, [[LD_i64_]], 4, 0, $noreg :: (store (s8) into %ir.p + 4, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri17]], 0, 0, 1, 8, [[LD_i64_]], 4, 2, $noreg :: (store (s8) into %ir.p + 4, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri18:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 24
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri18]], 0, 0, 1, 8, [[LD_i64_]], 3, 0, $noreg :: (store (s8) into %ir.p + 3, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri18]], 0, 0, 1, 8, [[LD_i64_]], 3, 2, $noreg :: (store (s8) into %ir.p + 3, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri19:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 16
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri19]], 0, 0, 1, 8, [[LD_i64_]], 2, 0, $noreg :: (store (s8) into %ir.p + 2, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri19]], 0, 0, 1, 8, [[LD_i64_]], 2, 2, $noreg :: (store (s8) into %ir.p + 2, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri20:%[0-9]+]]:b64 = SRL64_ri [[LDV_i64_v2_]], 8
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri20]], 0, 0, 1, 8, [[LD_i64_]], 1, 0, $noreg :: (store (s8) into %ir.p + 1, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri20]], 0, 0, 1, 8, [[LD_i64_]], 1, 2, $noreg :: (store (s8) into %ir.p + 1, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: Return
store <3 x i64> %v, ptr addrspace(1) %p, align 1, !alias.scope !2, !noalias !4, !mem.cache_hint !5
ret void
@@ -88,21 +88,21 @@ define void @unaligned_i64_store(ptr addrspace(1) %p, i64 %v) {
; CHECK: bb.0 (%ir-block.0):
; CHECK-NEXT: [[LD_i64_:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_i64_store_param_0>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
; CHECK-NEXT: [[LD_i64_1:%[0-9]+]]:b64 = LD_i64 0, 0, 102, 3, 64, -1, <mcsymbol unaligned_i64_store_param_1>, 0, 0, $noreg :: (dereferenceable invariant load (s64), addrspace 102)
- ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 0, 0, $noreg :: (store (s8) into %ir.p, addrspace 1)
+ ; CHECK-NEXT: ST_i64 [[LD_i64_1]], 0, 0, 1, 8, [[LD_i64_]], 0, 2, $noreg :: (store (s8) into %ir.p, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 56
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 7, 0, $noreg :: (store (s8) into %ir.p + 7, !alias.scope !0, !noalias !3, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri]], 0, 0, 1, 8, [[LD_i64_]], 7, 2, $noreg :: (store (s8) into %ir.p + 7, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri1:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 48
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 6, 0, $noreg :: (store (s8) into %ir.p + 6, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri1]], 0, 0, 1, 8, [[LD_i64_]], 6, 2, $noreg :: (store (s8) into %ir.p + 6, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri2:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 40
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 5, 0, $noreg :: (store (s8) into %ir.p + 5, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri2]], 0, 0, 1, 8, [[LD_i64_]], 5, 2, $noreg :: (store (s8) into %ir.p + 5, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri3:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 32
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 4, 0, $noreg :: (store (s8) into %ir.p + 4, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri3]], 0, 0, 1, 8, [[LD_i64_]], 4, 2, $noreg :: (store (s8) into %ir.p + 4, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri4:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 24
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 3, 0, $noreg :: (store (s8) into %ir.p + 3, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri4]], 0, 0, 1, 8, [[LD_i64_]], 3, 2, $noreg :: (store (s8) into %ir.p + 3, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri5:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 16
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 2, 0, $noreg :: (store (s8) into %ir.p + 2, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri5]], 0, 0, 1, 8, [[LD_i64_]], 2, 2, $noreg :: (store (s8) into %ir.p + 2, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: [[SRL64_ri6:%[0-9]+]]:b64 = SRL64_ri [[LD_i64_1]], 8
- ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 1, 0, $noreg :: (store (s8) into %ir.p + 1, addrspace 1)
+ ; CHECK-NEXT: ST_i64 killed [[SRL64_ri6]], 0, 0, 1, 8, [[LD_i64_]], 1, 2, $noreg :: (store (s8) into %ir.p + 1, !alias.scope !0, !noalias !3, !mem.cache_hint !6, addrspace 1)
; CHECK-NEXT: Return
store i64 %v, ptr addrspace(1) %p, align 1, !alias.scope !2, !noalias !4, !mem.cache_hint !5
ret void
diff --git a/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll b/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
index 7a2d74d0c9321..8e288d3371cfe 100644
--- a/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
+++ b/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
@@ -9,10 +9,10 @@ define fp128 @soften_f128_load(ptr %p) {
; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x11
; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
; CHECK-NEXT: [[COPY2:%[0-9]+]]:gpr = COPY [[COPY1]]
- ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 16)
- ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 16)
- ; CHECK-NEXT: [[LW2:%[0-9]+]]:gpr = LW [[COPY]], 8 :: (load (s32) from %ir.p + 8, align 8, basealign 16)
- ; CHECK-NEXT: [[LW3:%[0-9]+]]:gpr = LW [[COPY]], 12 :: (load (s32) from %ir.p + 12, basealign 16)
+ ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: [[LW2:%[0-9]+]]:gpr = LW [[COPY]], 8 :: (load (s32) from %ir.p + 8, align 8, basealign 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: [[LW3:%[0-9]+]]:gpr = LW [[COPY]], 12 :: (load (s32) from %ir.p + 12, basealign 16, !mem.cache_hint !1)
; CHECK-NEXT: SW killed [[LW3]], [[COPY2]], 12 :: (store (s32) into unknown-address + 12, basealign 16)
; CHECK-NEXT: SW killed [[LW2]], [[COPY2]], 8 :: (store (s32) into unknown-address + 8, align 8, basealign 16)
; CHECK-NEXT: SW killed [[LW1]], [[COPY2]], 4 :: (store (s32) into unknown-address + 4, basealign 16)
@@ -29,7 +29,7 @@ define double @soften_f16_extload(ptr %p) {
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x10
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def dead $x2, implicit $x2
- ; CHECK-NEXT: [[LH:%[0-9]+]]:gpr = LH [[COPY]], 0 :: (load (s16) from %ir.p)
+ ; CHECK-NEXT: [[LH:%[0-9]+]]:gpr = LH [[COPY]], 0 :: (load (s16) from %ir.p, !mem.cache_hint !1)
; CHECK-NEXT: $x10 = COPY [[LH]]
; CHECK-NEXT: PseudoCALL target-flags(riscv-call) &__extendhfsf2, csr_ilp32_lp64, implicit-def dead $x1, implicit $x10, implicit-def $x2, implicit-def $x10
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def dead $x2, implicit $x2
@@ -56,10 +56,10 @@ define i128 @expand_i128_load(ptr %p) {
; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x11
; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
; CHECK-NEXT: [[COPY2:%[0-9]+]]:gpr = COPY [[COPY1]]
- ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 16)
- ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 16)
- ; CHECK-NEXT: [[LW2:%[0-9]+]]:gpr = LW [[COPY]], 8 :: (load (s32) from %ir.p + 8, align 8, basealign 16)
- ; CHECK-NEXT: [[LW3:%[0-9]+]]:gpr = LW [[COPY]], 12 :: (load (s32) from %ir.p + 12, basealign 16)
+ ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: [[LW2:%[0-9]+]]:gpr = LW [[COPY]], 8 :: (load (s32) from %ir.p + 8, align 8, basealign 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: [[LW3:%[0-9]+]]:gpr = LW [[COPY]], 12 :: (load (s32) from %ir.p + 12, basealign 16, !mem.cache_hint !1)
; CHECK-NEXT: SW killed [[LW3]], [[COPY2]], 12 :: (store (s32) into unknown-address + 12, basealign 8)
; CHECK-NEXT: SW killed [[LW2]], [[COPY2]], 8 :: (store (s32) into unknown-address + 8, align 8)
; CHECK-NEXT: SW killed [[LW1]], [[COPY2]], 4 :: (store (s32) into unknown-address + 4, basealign 8)
@@ -80,10 +80,10 @@ define void @expand_i128_store(ptr %p, i128 %v) {
; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32))
; CHECK-NEXT: [[LW2:%[0-9]+]]:gpr = LW [[COPY]], 8 :: (load (s32))
; CHECK-NEXT: [[LW3:%[0-9]+]]:gpr = LW [[COPY]], 12 :: (load (s32))
- ; CHECK-NEXT: SW killed [[LW3]], [[COPY1]], 12 :: (store (s32) into %ir.p + 12, basealign 16)
- ; CHECK-NEXT: SW killed [[LW2]], [[COPY1]], 8 :: (store (s32) into %ir.p + 8, align 8, basealign 16)
- ; CHECK-NEXT: SW killed [[LW1]], [[COPY1]], 4 :: (store (s32) into %ir.p + 4, basealign 16)
- ; CHECK-NEXT: SW killed [[LW]], [[COPY1]], 0 :: (store (s32) into %ir.p, align 16)
+ ; CHECK-NEXT: SW killed [[LW3]], [[COPY1]], 12 :: (store (s32) into %ir.p + 12, basealign 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: SW killed [[LW2]], [[COPY1]], 8 :: (store (s32) into %ir.p + 8, align 8, basealign 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: SW killed [[LW1]], [[COPY1]], 4 :: (store (s32) into %ir.p + 4, basealign 16, !mem.cache_hint !1)
+ ; CHECK-NEXT: SW killed [[LW]], [[COPY1]], 0 :: (store (s32) into %ir.p, align 16, !mem.cache_hint !1)
; CHECK-NEXT: PseudoRET
store i128 %v, ptr %p, align 16, !mem.cache_hint !1
ret void
@@ -97,8 +97,8 @@ define i128 @expand_i128_extload(ptr %p) {
; CHECK-NEXT: [[COPY:%[0-9]+]]:gpr = COPY $x11
; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
; CHECK-NEXT: [[COPY2:%[0-9]+]]:gpr = COPY [[COPY1]]
- ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 8)
- ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 8)
+ ; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32) from %ir.p, align 8, !mem.cache_hint !1)
+ ; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32) from %ir.p + 4, basealign 8, !mem.cache_hint !1)
; CHECK-NEXT: [[COPY3:%[0-9]+]]:gpr = COPY $x0
; CHECK-NEXT: SW [[COPY3]], [[COPY2]], 12 :: (store (s32) into unknown-address + 12, basealign 8)
; CHECK-NEXT: SW [[COPY3]], [[COPY2]], 8 :: (store (s32) into unknown-address + 8, align 8)
@@ -119,8 +119,8 @@ define void @expand_i128_truncstore(ptr %p, i128 %v) {
; CHECK-NEXT: [[COPY1:%[0-9]+]]:gpr = COPY $x10
; CHECK-NEXT: [[LW:%[0-9]+]]:gpr = LW [[COPY]], 0 :: (load (s32))
; CHECK-NEXT: [[LW1:%[0-9]+]]:gpr = LW [[COPY]], 4 :: (load (s32))
- ; CHECK-NEXT: SW killed [[LW1]], [[COPY1]], 4 :: (store (s32) into %ir.p + 4, basealign 8)
- ; CHECK-NEXT: SW killed [[LW]], [[COPY1]], 0 :: (store (s32) into %ir.p, align 8)
+ ; CHECK-NEXT: SW killed [[LW1]], [[COPY1]], 4 :: (store (s32) into %ir.p + 4, basealign 8, !mem.cache_hint !1)
+ ; CHECK-NEXT: SW killed [[LW]], [[COPY1]], 0 :: (store (s32) into %ir.p, align 8, !mem.cache_hint !1)
; CHECK-NEXT: PseudoRET
%trunc = trunc i128 %v to i64
store i64 %trunc, ptr %p, align 8, !mem.cache_hint !1
>From 368168a005fb478d173618ab0d0c63dbff81a0d7 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Tue, 22 Sep 2026 19:16:53 +0000
Subject: [PATCH 3/7] format
---
.../SelectionDAG/LegalizeFloatTypes.cpp | 19 +++--
.../SelectionDAG/LegalizeVectorTypes.cpp | 11 ++-
.../CodeGen/SelectionDAG/TargetLowering.cpp | 80 ++++++++-----------
3 files changed, 49 insertions(+), 61 deletions(-)
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
index 1b9b68705195a..c1b0c5771b249 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
@@ -928,11 +928,10 @@ SDValue DAGTypeLegalizer::SoftenFloatRes_LOAD(SDNode *N) {
}
// Do a non-extending load followed by FP_EXTEND.
- NewL = DAG.getLoad(
- L->getAddressingMode(), ISD::NON_EXTLOAD, L->getMemoryVT(), dl,
- L->getChain(), L->getBasePtr(), L->getOffset(), L->getPointerInfo(),
- L->getMemoryVT(), L->getBaseAlign(), MMOFlags,
- L->getNonRangeMMOMetadata());
+ NewL = DAG.getLoad(L->getAddressingMode(), ISD::NON_EXTLOAD, L->getMemoryVT(),
+ dl, L->getChain(), L->getBasePtr(), L->getOffset(),
+ L->getPointerInfo(), L->getMemoryVT(), L->getBaseAlign(),
+ MMOFlags, L->getNonRangeMMOMetadata());
// Legalized the chain result - switch anything that used the old chain to
// use the new one.
ReplaceValueWith(SDValue(N, 1), NewL.getValue(1));
@@ -2819,11 +2818,11 @@ SDValue DAGTypeLegalizer::SoftPromoteHalfRes_LOAD(SDNode *N) {
// Load the value as an integer value with the same number of bits.
assert(L->getExtensionType() == ISD::NON_EXTLOAD && "Unexpected extension!");
- SDValue NewL = DAG.getLoad(
- L->getAddressingMode(), L->getExtensionType(), MVT::i16, SDLoc(N),
- L->getChain(), L->getBasePtr(), L->getOffset(), L->getPointerInfo(),
- MVT::i16, L->getBaseAlign(), L->getMemOperand()->getFlags(),
- L->getNonRangeMMOMetadata());
+ SDValue NewL =
+ DAG.getLoad(L->getAddressingMode(), L->getExtensionType(), MVT::i16,
+ SDLoc(N), L->getChain(), L->getBasePtr(), L->getOffset(),
+ L->getPointerInfo(), MVT::i16, L->getBaseAlign(),
+ L->getMemOperand()->getFlags(), L->getNonRangeMMOMetadata());
// Legalize the chain result by replacing uses of the old value chain with the
// new one
ReplaceValueWith(SDValue(N, 1), NewL.getValue(1));
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
index 8b0dccdac71fc..a6c68e43bb5f8 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
@@ -1205,13 +1205,12 @@ SDValue DAGTypeLegalizer::ScalarizeVecOp_STORE(StoreSDNode *N, unsigned OpNo){
N->getChain(), dl, GetScalarizedVector(N->getOperand(1)),
N->getBasePtr(), N->getPointerInfo(),
N->getMemoryVT().getVectorElementType(), N->getBaseAlign(),
- N->getMemOperand()->getFlags(),
- N->getNonRangeMMOMetadata());
+ N->getMemOperand()->getFlags(), N->getNonRangeMMOMetadata());
- return DAG.getStore(
- N->getChain(), dl, GetScalarizedVector(N->getOperand(1)), N->getBasePtr(),
- N->getPointerInfo(), N->getBaseAlign(), N->getMemOperand()->getFlags(),
- N->getNonRangeMMOMetadata());
+ return DAG.getStore(N->getChain(), dl, GetScalarizedVector(N->getOperand(1)),
+ N->getBasePtr(), N->getPointerInfo(), N->getBaseAlign(),
+ N->getMemOperand()->getFlags(),
+ N->getNonRangeMMOMetadata());
}
/// If the value to store is a vector that needs to be scalarized, it must be
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index 67e14373bfc03..cc3f6f83470ad 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -11779,10 +11779,9 @@ SDValue TargetLowering::scalarizeVectorStore(StoreSDNode *ST,
CurrVal = DAG.getNode(ISD::OR, SL, IntVT, CurrVal, ShiftedElt);
}
- return DAG.getStore(
- Chain, SL, CurrVal, BasePtr, ST->getPointerInfo(), ST->getBaseAlign(),
- ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata());
+ return DAG.getStore(Chain, SL, CurrVal, BasePtr, ST->getPointerInfo(),
+ ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata());
}
// Store Stride in bytes
@@ -11885,8 +11884,7 @@ TargetLowering::expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const {
SDValue Load = DAG.getExtLoad(
ISD::EXTLOAD, dl, RegVT, Chain, Ptr,
LD->getPointerInfo().getWithOffset(Offset), MemVT, LD->getBaseAlign(),
- LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ LD->getMemOperand()->getFlags(), LD->getNonRangeMMOMetadata());
// Follow the load with a store to the stack slot. Remember the store.
// On big-endian machines this requires a truncating store to ensure
// that the bits end up in the right place.
@@ -11927,29 +11925,25 @@ TargetLowering::expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const {
// Load the value in two parts
SDValue Lo, Hi;
if (DAG.getDataLayout().isLittleEndian()) {
- Lo = DAG.getExtLoad(
- ISD::ZEXTLOAD, dl, VT, Chain, Ptr, LD->getPointerInfo(), NewLoadedVT,
- Alignment, LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, VT, Chain, Ptr, LD->getPointerInfo(),
+ NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
- Hi = DAG.getExtLoad(
- HiExtType, dl, VT, Chain, Ptr,
- LD->getPointerInfo().getWithOffset(IncrementSize), NewLoadedVT,
- Alignment, LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ Hi = DAG.getExtLoad(HiExtType, dl, VT, Chain, Ptr,
+ LD->getPointerInfo().getWithOffset(IncrementSize),
+ NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
} else {
- Hi = DAG.getExtLoad(
- HiExtType, dl, VT, Chain, Ptr, LD->getPointerInfo(), NewLoadedVT,
- Alignment, LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ Hi = DAG.getExtLoad(HiExtType, dl, VT, Chain, Ptr, LD->getPointerInfo(),
+ NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
- Lo = DAG.getExtLoad(
- ISD::ZEXTLOAD, dl, VT, Chain, Ptr,
- LD->getPointerInfo().getWithOffset(IncrementSize), NewLoadedVT,
- Alignment, LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, VT, Chain, Ptr,
+ LD->getPointerInfo().getWithOffset(IncrementSize),
+ NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
+ LD->getNonRangeMMOMetadata());
}
// aggregate the two parts
@@ -11989,10 +11983,9 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
// same size, then a (misaligned) int store.
// FIXME: Does not handle truncating floating point stores!
SDValue Result = DAG.getNode(ISD::BITCAST, dl, intVT, Val);
- Result = DAG.getStore(
- Chain, dl, Result, Ptr, ST->getPointerInfo(), Alignment,
- ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata());
+ Result = DAG.getStore(Chain, dl, Result, Ptr, ST->getPointerInfo(),
+ Alignment, ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata());
return Result;
}
// Do a (aligned) store to a stack slot, then copy from the stack slot
@@ -12028,11 +12021,10 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
RegVT, dl, Store, StackPtr,
MachinePointerInfo::getFixedStack(MF, FrameIndex, Offset));
// Store it to the final location. Remember the store.
- Stores.push_back(DAG.getStore(Load.getValue(1), dl, Load, Ptr,
- ST->getPointerInfo().getWithOffset(Offset),
- ST->getBaseAlign(),
- ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata()));
+ Stores.push_back(DAG.getStore(
+ Load.getValue(1), dl, Load, Ptr,
+ ST->getPointerInfo().getWithOffset(Offset), ST->getBaseAlign(),
+ ST->getMemOperand()->getFlags(), ST->getNonRangeMMOMetadata()));
// Increment the pointers.
Offset += RegBytes;
StackPtr = DAG.getObjectPtrOffset(dl, StackPtr, StackPtrIncrement);
@@ -12050,11 +12042,11 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
ISD::EXTLOAD, dl, RegVT, Store, StackPtr,
MachinePointerInfo::getFixedStack(MF, FrameIndex, Offset), LoadMemVT);
- Stores.push_back(DAG.getTruncStore(
- Load.getValue(1), dl, Load, Ptr,
- ST->getPointerInfo().getWithOffset(Offset), LoadMemVT,
- ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata()));
+ Stores.push_back(
+ DAG.getTruncStore(Load.getValue(1), dl, Load, Ptr,
+ ST->getPointerInfo().getWithOffset(Offset), LoadMemVT,
+ ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
+ ST->getNonRangeMMOMetadata()));
// The order of the stores doesn't matter - say it with a TokenFactor.
SDValue Result = DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Stores);
return Result;
@@ -12083,18 +12075,16 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
// Store the two parts
SDValue Store1, Store2;
- Store1 = DAG.getTruncStore(Chain, dl,
- DAG.getDataLayout().isLittleEndian() ? Lo : Hi,
- Ptr, ST->getPointerInfo(), NewStoredVT, Alignment,
- ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata());
+ Store1 = DAG.getTruncStore(
+ Chain, dl, DAG.getDataLayout().isLittleEndian() ? Lo : Hi, Ptr,
+ ST->getPointerInfo(), NewStoredVT, Alignment,
+ ST->getMemOperand()->getFlags(), ST->getNonRangeMMOMetadata());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
Store2 = DAG.getTruncStore(
Chain, dl, DAG.getDataLayout().isLittleEndian() ? Hi : Lo, Ptr,
ST->getPointerInfo().getWithOffset(IncrementSize), NewStoredVT, Alignment,
- ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata());
+ ST->getMemOperand()->getFlags(), ST->getNonRangeMMOMetadata());
SDValue Result =
DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Store1, Store2);
>From ffee8ed2c6e2c9773656a9ddc8c35fe3b8945759 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Wed, 30 Sep 2026 06:36:16 +0000
Subject: [PATCH 4/7] keep todo
---
llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll | 6 ++++++
1 file changed, 6 insertions(+)
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll b/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
index ac86b770dd79f..8a67b3bbfe615 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-transforms.ll
@@ -86,6 +86,12 @@ define i32 @test_load_l2_normal(ptr addrspace(1) %p) {
ret i32 %v
}
+;-----------------------------------------------------------------------------
+; TODO: Preserve cache hints across DAGCombiner-created memory rewrites.
+; This documents the current store-of-concat-trunc behavior: copied MMOs for
+; the split stores do not retain !mem.cache_hint metadata yet.
+;-----------------------------------------------------------------------------
+
define void @test_dagcombine_store_concat_trunc_v8i32(ptr addrspace(1) %p, <4 x i64> %a, <4 x i64> %b) {
; CHECK-LABEL: test_dagcombine_store_concat_trunc_v8i32(
; CHECK: mov.b64 %rd10, 12345;
>From 75e79588fd203999a0e7242e3286521ec612d7c2 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Wed, 30 Sep 2026 18:22:17 +0000
Subject: [PATCH 5/7] change name
---
llvm/include/llvm/CodeGen/SelectionDAGNodes.h | 7 ++--
llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp | 4 +--
.../SelectionDAG/LegalizeFloatTypes.cpp | 14 ++++----
.../SelectionDAG/LegalizeIntegerTypes.cpp | 4 +--
.../SelectionDAG/LegalizeTypesGeneric.cpp | 4 +--
.../SelectionDAG/LegalizeVectorTypes.cpp | 16 ++++-----
.../CodeGen/SelectionDAG/TargetLowering.cpp | 34 +++++++++----------
7 files changed, 41 insertions(+), 42 deletions(-)
diff --git a/llvm/include/llvm/CodeGen/SelectionDAGNodes.h b/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
index a3d2f2c882642..b7bed921fcd4b 100644
--- a/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
+++ b/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
@@ -1493,10 +1493,9 @@ class MemSDNode : public SDNode {
return getMemOperand()->getMemCacheHint();
}
- /// Returns LLVM IR metadata carried by this memory access, except for range
- /// metadata. Range metadata describes a loaded value and cannot be blindly
- /// transferred when an access is split or its type changes.
- MMOMetadata getNonRangeMMOMetadata() const {
+ /// Returns metadata that can be copied unchanged to an access covering
+ /// all or part of this access's bytes, even if the value type changes.
+ MMOMetadata getMMOMetadataForSubAccess() const {
return MMOMetadata(getAAInfo(), /*Ranges=*/nullptr, getMemCacheHint());
}
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
index 09441f1858fe9..090ccbd727ab5 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
@@ -496,7 +496,7 @@ void SelectionDAGLegalize::LegalizeStoreOps(SDNode *Node) {
SDLoc dl(Node);
MachineMemOperand::Flags MMOFlags = ST->getMemOperand()->getFlags();
- MMOMetadata Metadata = ST->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = ST->getMMOMetadataForSubAccess();
if (!ST->isTruncatingStore()) {
LLVM_DEBUG(dbgs() << "Legalizing store operation\n");
@@ -737,7 +737,7 @@ void SelectionDAGLegalize::LegalizeLoadOps(SDNode *Node) {
EVT SrcVT = LD->getMemoryVT();
TypeSize SrcWidth = SrcVT.getSizeInBits();
MachineMemOperand::Flags MMOFlags = LD->getMemOperand()->getFlags();
- MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = LD->getMMOMetadataForSubAccess();
if (SrcWidth != SrcVT.getStoreSizeInBits() &&
// Some targets pretend to have an i1 loading operation, and actually
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
index c1b0c5771b249..5850a836d8e6b 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
@@ -920,7 +920,7 @@ SDValue DAGTypeLegalizer::SoftenFloatRes_LOAD(SDNode *N) {
NewL = DAG.getLoad(L->getAddressingMode(), ISD::EXTLOAD, NVT, dl,
L->getChain(), L->getBasePtr(), L->getOffset(),
L->getPointerInfo(), MemVT, L->getBaseAlign(), MMOFlags,
- L->getNonRangeMMOMetadata());
+ L->getMMOMetadataForSubAccess());
// Legalized the chain result - switch anything that used the old chain to
// use the new one.
ReplaceValueWith(SDValue(N, 1), NewL.getValue(1));
@@ -931,7 +931,7 @@ SDValue DAGTypeLegalizer::SoftenFloatRes_LOAD(SDNode *N) {
NewL = DAG.getLoad(L->getAddressingMode(), ISD::NON_EXTLOAD, L->getMemoryVT(),
dl, L->getChain(), L->getBasePtr(), L->getOffset(),
L->getPointerInfo(), L->getMemoryVT(), L->getBaseAlign(),
- MMOFlags, L->getNonRangeMMOMetadata());
+ MMOFlags, L->getMMOMetadataForSubAccess());
// Legalized the chain result - switch anything that used the old chain to
// use the new one.
ReplaceValueWith(SDValue(N, 1), NewL.getValue(1));
@@ -2818,11 +2818,11 @@ SDValue DAGTypeLegalizer::SoftPromoteHalfRes_LOAD(SDNode *N) {
// Load the value as an integer value with the same number of bits.
assert(L->getExtensionType() == ISD::NON_EXTLOAD && "Unexpected extension!");
- SDValue NewL =
- DAG.getLoad(L->getAddressingMode(), L->getExtensionType(), MVT::i16,
- SDLoc(N), L->getChain(), L->getBasePtr(), L->getOffset(),
- L->getPointerInfo(), MVT::i16, L->getBaseAlign(),
- L->getMemOperand()->getFlags(), L->getNonRangeMMOMetadata());
+ SDValue NewL = DAG.getLoad(L->getAddressingMode(), L->getExtensionType(),
+ MVT::i16, SDLoc(N), L->getChain(), L->getBasePtr(),
+ L->getOffset(), L->getPointerInfo(), MVT::i16,
+ L->getBaseAlign(), L->getMemOperand()->getFlags(),
+ L->getMMOMetadataForSubAccess());
// Legalize the chain result by replacing uses of the old value chain with the
// new one
ReplaceValueWith(SDValue(N, 1), NewL.getValue(1));
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
index 4ffb2a37340f9..6ed97ad1849d4 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
@@ -4445,7 +4445,7 @@ void DAGTypeLegalizer::ExpandIntRes_LOAD(LoadSDNode *N,
SDValue Ptr = N->getBasePtr();
ISD::LoadExtType ExtType = N->getExtensionType();
MachineMemOperand::Flags MMOFlags = N->getMemOperand()->getFlags();
- MMOMetadata Metadata = N->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = N->getMMOMetadataForSubAccess();
SDLoc dl(N);
assert(NVT.isByteSized() && "Expanded type not byte sized!");
@@ -6008,7 +6008,7 @@ SDValue DAGTypeLegalizer::ExpandIntOp_STORE(StoreSDNode *N, unsigned OpNo) {
SDValue Ch = N->getChain();
SDValue Ptr = N->getBasePtr();
MachineMemOperand::Flags MMOFlags = N->getMemOperand()->getFlags();
- MMOMetadata Metadata = N->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = N->getMMOMetadataForSubAccess();
SDLoc dl(N);
SDValue Lo, Hi;
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeTypesGeneric.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeTypesGeneric.cpp
index 2e3a6495cbe05..ee2c3c1d6b707 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeTypesGeneric.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeTypesGeneric.cpp
@@ -254,7 +254,7 @@ void DAGTypeLegalizer::ExpandRes_NormalLoad(SDNode *N, SDValue &Lo,
EVT NVT = TLI.getTypeToTransformTo(*DAG.getContext(), ValueVT);
SDValue Chain = LD->getChain();
SDValue Ptr = LD->getBasePtr();
- MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = LD->getMMOMetadataForSubAccess();
assert(NVT.isByteSized() && "Expanded type not byte sized!");
@@ -483,7 +483,7 @@ SDValue DAGTypeLegalizer::ExpandOp_NormalStore(SDNode *N, unsigned OpNo) {
EVT NVT = TLI.getTypeToTransformTo(*DAG.getContext(), ValueVT);
SDValue Chain = St->getChain();
SDValue Ptr = St->getBasePtr();
- MMOMetadata Metadata = St->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = St->getMMOMetadataForSubAccess();
assert(NVT.isByteSized() && "Expanded type not byte sized!");
unsigned IncrementSize = NVT.getSizeInBits() / 8;
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
index b2f261040ce7f..9be87b90e4534 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
@@ -602,7 +602,7 @@ SDValue DAGTypeLegalizer::ScalarizeVecRes_LOAD(LoadSDNode *N) {
N->getBasePtr(), DAG.getPOISON(N->getBasePtr().getValueType()),
N->getPointerInfo(), N->getMemoryVT().getVectorElementType(),
N->getBaseAlign(), N->getMemOperand()->getFlags(),
- N->getNonRangeMMOMetadata());
+ N->getMMOMetadataForSubAccess());
// Legalize the chain result - switch anything that used the old chain to
// use the new one.
@@ -1228,12 +1228,12 @@ SDValue DAGTypeLegalizer::ScalarizeVecOp_STORE(StoreSDNode *N, unsigned OpNo){
N->getChain(), dl, GetScalarizedVector(N->getOperand(1)),
N->getBasePtr(), N->getPointerInfo(),
N->getMemoryVT().getVectorElementType(), N->getBaseAlign(),
- N->getMemOperand()->getFlags(), N->getNonRangeMMOMetadata());
+ N->getMemOperand()->getFlags(), N->getMMOMetadataForSubAccess());
return DAG.getStore(N->getChain(), dl, GetScalarizedVector(N->getOperand(1)),
N->getBasePtr(), N->getPointerInfo(), N->getBaseAlign(),
N->getMemOperand()->getFlags(),
- N->getNonRangeMMOMetadata());
+ N->getMMOMetadataForSubAccess());
}
/// If the value to store is a vector that needs to be scalarized, it must be
@@ -2458,7 +2458,7 @@ void DAGTypeLegalizer::SplitVecRes_LOAD(LoadSDNode *LD, SDValue &Lo,
SDValue Offset = DAG.getPOISON(Ptr.getValueType());
EVT MemoryVT = LD->getMemoryVT();
MachineMemOperand::Flags MMOFlags = LD->getMemOperand()->getFlags();
- MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = LD->getMMOMetadataForSubAccess();
EVT LoMemVT, HiMemVT;
std::tie(LoMemVT, HiMemVT) = DAG.GetSplitDestVTs(MemoryVT);
@@ -4797,7 +4797,7 @@ SDValue DAGTypeLegalizer::SplitVecOp_STORE(StoreSDNode *N, unsigned OpNo) {
EVT MemoryVT = N->getMemoryVT();
Align Alignment = N->getBaseAlign();
MachineMemOperand::Flags MMOFlags = N->getMemOperand()->getFlags();
- MMOMetadata Metadata = N->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = N->getMMOMetadataForSubAccess();
SDValue Lo, Hi;
GetSplitVector(N->getOperand(1), Lo, Hi);
@@ -9232,7 +9232,7 @@ SDValue DAGTypeLegalizer::GenWidenVectorLoads(SmallVectorImpl<SDValue> &LdChain,
SDValue Chain = LD->getChain();
SDValue BasePtr = LD->getBasePtr();
MachineMemOperand::Flags MMOFlags = LD->getMemOperand()->getFlags();
- MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = LD->getMMOMetadataForSubAccess();
TypeSize LdWidth = LdVT.getSizeInBits();
TypeSize WidenWidth = WidenVT.getSizeInBits();
@@ -9389,7 +9389,7 @@ DAGTypeLegalizer::GenWidenVectorExtLoads(SmallVectorImpl<SDValue> &LdChain,
SDValue Chain = LD->getChain();
SDValue BasePtr = LD->getBasePtr();
MachineMemOperand::Flags MMOFlags = LD->getMemOperand()->getFlags();
- MMOMetadata Metadata = LD->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = LD->getMMOMetadataForSubAccess();
if (LdVT.isScalableVector())
return SDValue();
@@ -9432,7 +9432,7 @@ bool DAGTypeLegalizer::GenWidenVectorStores(SmallVectorImpl<SDValue> &StChain,
SDValue Chain = ST->getChain();
SDValue BasePtr = ST->getBasePtr();
MachineMemOperand::Flags MMOFlags = ST->getMemOperand()->getFlags();
- MMOMetadata Metadata = ST->getNonRangeMMOMetadata();
+ MMOMetadata Metadata = ST->getMMOMetadataForSubAccess();
SDValue ValOp = GetWidenedVector(ST->getValue());
SDLoc dl(ST);
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index fb0fb11694e0f..02f3676752c1d 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -11857,7 +11857,7 @@ TargetLowering::scalarizeVectorLoad(LoadSDNode *LD,
SDValue Load = DAG.getExtLoad(
ISD::EXTLOAD, SL, LoadVT, Chain, BasePTR, LD->getPointerInfo(),
SrcIntVT, LD->getBaseAlign(), LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ LD->getMMOMetadataForSubAccess());
SmallVector<SDValue, 8> Vals;
for (unsigned Idx = 0; Idx < NumElem; ++Idx) {
@@ -11893,7 +11893,7 @@ TargetLowering::scalarizeVectorLoad(LoadSDNode *LD,
ExtType, SL, DstEltVT, Chain, BasePTR,
LD->getPointerInfo().getWithOffset(Idx * Stride), SrcEltVT,
LD->getBaseAlign(), LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ LD->getMMOMetadataForSubAccess());
BasePTR = DAG.getObjectPtrOffset(SL, BasePTR, TypeSize::getFixed(Stride));
@@ -11955,7 +11955,7 @@ SDValue TargetLowering::scalarizeVectorStore(StoreSDNode *ST,
return DAG.getStore(Chain, SL, CurrVal, BasePtr, ST->getPointerInfo(),
ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata());
+ ST->getMMOMetadataForSubAccess());
}
// Store Stride in bytes
@@ -11974,7 +11974,7 @@ SDValue TargetLowering::scalarizeVectorStore(StoreSDNode *ST,
SDValue Store = DAG.getTruncStore(
Chain, SL, Elt, Ptr, ST->getPointerInfo().getWithOffset(Idx * Stride),
MemSclVT, ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata());
+ ST->getMMOMetadataForSubAccess());
Stores.push_back(Store);
}
@@ -12040,7 +12040,7 @@ TargetLowering::expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const {
SDValue Load = DAG.getLoad(
RegVT, dl, Chain, Ptr, LD->getPointerInfo().getWithOffset(Offset),
LD->getBaseAlign(), LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ LD->getMMOMetadataForSubAccess());
// Follow the load with a store to the stack slot. Remember the store.
Stores.push_back(DAG.getStore(
Load.getValue(1), dl, Load, StackPtr,
@@ -12058,7 +12058,7 @@ TargetLowering::expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const {
SDValue Load = DAG.getExtLoad(
ISD::EXTLOAD, dl, RegVT, Chain, Ptr,
LD->getPointerInfo().getWithOffset(Offset), MemVT, LD->getBaseAlign(),
- LD->getMemOperand()->getFlags(), LD->getNonRangeMMOMetadata());
+ LD->getMemOperand()->getFlags(), LD->getMMOMetadataForSubAccess());
// Follow the load with a store to the stack slot. Remember the store.
// On big-endian machines this requires a truncating store to ensure
// that the bits end up in the right place.
@@ -12101,23 +12101,23 @@ TargetLowering::expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const {
if (DAG.getDataLayout().isLittleEndian()) {
Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, VT, Chain, Ptr, LD->getPointerInfo(),
NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ LD->getMMOMetadataForSubAccess());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
Hi = DAG.getExtLoad(HiExtType, dl, VT, Chain, Ptr,
LD->getPointerInfo().getWithOffset(IncrementSize),
NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ LD->getMMOMetadataForSubAccess());
} else {
Hi = DAG.getExtLoad(HiExtType, dl, VT, Chain, Ptr, LD->getPointerInfo(),
NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ LD->getMMOMetadataForSubAccess());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
Lo = DAG.getExtLoad(ISD::ZEXTLOAD, dl, VT, Chain, Ptr,
LD->getPointerInfo().getWithOffset(IncrementSize),
NewLoadedVT, Alignment, LD->getMemOperand()->getFlags(),
- LD->getNonRangeMMOMetadata());
+ LD->getMMOMetadataForSubAccess());
}
// aggregate the two parts
@@ -12159,7 +12159,7 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
SDValue Result = DAG.getNode(ISD::BITCAST, dl, intVT, Val);
Result = DAG.getStore(Chain, dl, Result, Ptr, ST->getPointerInfo(),
Alignment, ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata());
+ ST->getMMOMetadataForSubAccess());
return Result;
}
// Do a (aligned) store to a stack slot, then copy from the stack slot
@@ -12198,7 +12198,7 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
Stores.push_back(DAG.getStore(
Load.getValue(1), dl, Load, Ptr,
ST->getPointerInfo().getWithOffset(Offset), ST->getBaseAlign(),
- ST->getMemOperand()->getFlags(), ST->getNonRangeMMOMetadata()));
+ ST->getMemOperand()->getFlags(), ST->getMMOMetadataForSubAccess()));
// Increment the pointers.
Offset += RegBytes;
StackPtr = DAG.getObjectPtrOffset(dl, StackPtr, StackPtrIncrement);
@@ -12220,7 +12220,7 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
DAG.getTruncStore(Load.getValue(1), dl, Load, Ptr,
ST->getPointerInfo().getWithOffset(Offset), LoadMemVT,
ST->getBaseAlign(), ST->getMemOperand()->getFlags(),
- ST->getNonRangeMMOMetadata()));
+ ST->getMMOMetadataForSubAccess()));
// The order of the stores doesn't matter - say it with a TokenFactor.
SDValue Result = DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Stores);
return Result;
@@ -12252,13 +12252,13 @@ SDValue TargetLowering::expandUnalignedStore(StoreSDNode *ST,
Store1 = DAG.getTruncStore(
Chain, dl, DAG.getDataLayout().isLittleEndian() ? Lo : Hi, Ptr,
ST->getPointerInfo(), NewStoredVT, Alignment,
- ST->getMemOperand()->getFlags(), ST->getNonRangeMMOMetadata());
+ ST->getMemOperand()->getFlags(), ST->getMMOMetadataForSubAccess());
Ptr = DAG.getObjectPtrOffset(dl, Ptr, TypeSize::getFixed(IncrementSize));
Store2 = DAG.getTruncStore(
Chain, dl, DAG.getDataLayout().isLittleEndian() ? Hi : Lo, Ptr,
ST->getPointerInfo().getWithOffset(IncrementSize), NewStoredVT, Alignment,
- ST->getMemOperand()->getFlags(), ST->getNonRangeMMOMetadata());
+ ST->getMemOperand()->getFlags(), ST->getMMOMetadataForSubAccess());
SDValue Result =
DAG.getNode(ISD::TokenFactor, dl, MVT::Other, Store1, Store2);
@@ -14481,13 +14481,13 @@ SDValue TargetLowering::scalarizeExtractedVectorLoad(EVT ResultVT,
Load = DAG.getExtLoad(ExtType, DL, ResultVT, OriginalLoad->getChain(),
NewPtr, MPI, VecEltVT, Alignment,
OriginalLoad->getMemOperand()->getFlags(),
- OriginalLoad->getNonRangeMMOMetadata());
+ OriginalLoad->getMMOMetadataForSubAccess());
DAG.makeEquivalentMemoryOrdering(OriginalLoad, Load);
} else {
// The result type is narrower or the same width as the vector element
Load = DAG.getLoad(VecEltVT, DL, OriginalLoad->getChain(), NewPtr, MPI,
Alignment, OriginalLoad->getMemOperand()->getFlags(),
- OriginalLoad->getNonRangeMMOMetadata());
+ OriginalLoad->getMMOMetadataForSubAccess());
DAG.makeEquivalentMemoryOrdering(OriginalLoad, Load);
if (ResultVT.bitsLT(VecEltVT))
Load = DAG.getNode(ISD::TRUNCATE, DL, ResultVT, Load);
>From 2460373c3e27a462b9e50ffa42f647917e9d63e7 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Wed, 30 Sep 2026 19:29:20 +0000
Subject: [PATCH 6/7] [SelectionDAG] Check cache hint metadata contents in MIR
tests
---
llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll | 2 ++
llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll | 2 ++
2 files changed, 4 insertions(+)
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll b/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
index cc92dbc6b8d58..abf470b853e28 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-unaligned-store-aa.ll
@@ -2,6 +2,8 @@
; RUN: llc -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 \
; RUN: -stop-after=finalize-isel %s -o - | FileCheck %s
+; CHECK-DAG: !6 = !{!"nvvm.l1_eviction", !"first"}
+
define void @unaligned_f64_store(ptr addrspace(1) %p, double %v) {
; CHECK-LABEL: name: unaligned_f64_store
; CHECK: bb.0 (%ir-block.0):
diff --git a/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll b/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
index 8e288d3371cfe..b79ce9a951676 100644
--- a/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
+++ b/llvm/test/CodeGen/RISCV/mem-cache-hint-legalize-types.ll
@@ -1,6 +1,8 @@
; NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
; RUN: llc -mtriple=riscv32 -stop-after=finalize-isel %s -o - | FileCheck %s
+; CHECK-DAG: !1 = !{!"test.cache_hint", !"value"}
+
define fp128 @soften_f128_load(ptr %p) {
; CHECK-LABEL: name: soften_f128_load
; CHECK: bb.0 (%ir-block.0):
>From 9c1d580ad3697cf3344e4866eb61913c4785afbc Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Fri, 2 Oct 2026 18:58:28 +0000
Subject: [PATCH 7/7] adjust comment
---
llvm/include/llvm/CodeGen/SelectionDAGNodes.h | 2 ++
1 file changed, 2 insertions(+)
diff --git a/llvm/include/llvm/CodeGen/SelectionDAGNodes.h b/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
index b7bed921fcd4b..ff6331d6b9b38 100644
--- a/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
+++ b/llvm/include/llvm/CodeGen/SelectionDAGNodes.h
@@ -1495,6 +1495,8 @@ class MemSDNode : public SDNode {
/// Returns metadata that can be copied unchanged to an access covering
/// all or part of this access's bytes, even if the value type changes.
+ /// Metadata describing the loaded value (e.g. !range) is excluded, since
+ /// it may not apply to a value loaded with a different type or byte range.
MMOMetadata getMMOMetadataForSubAccess() const {
return MMOMetadata(getAAInfo(), /*Ranges=*/nullptr, getMemCacheHint());
}
More information about the llvm-commits
mailing list