[llvm] [SelectionDAG][NVPTX] Support cache hints on masked loads and stores (PR #215307)
Yonah Goldberg via llvm-commits
llvm-commits at lists.llvm.org
Thu Aug 27 13:56:51 PDT 2026
https://github.com/YonahGoldberg updated https://github.com/llvm/llvm-project/pull/215307
>From 76589e08498eafc30000816a762023261b3b0d3a Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Mon, 10 Aug 2026 14:58:03 +0000
Subject: [PATCH 1/4] masked load/store
---
.../SelectionDAG/LegalizeVectorTypes.cpp | 12 ++-
.../SelectionDAG/SelectionDAGBuilder.cpp | 7 +-
.../Scalar/ScalarizeMaskedMemIntrin.cpp | 14 ++-
.../CodeGen/NVPTX/cache-hint-intrinsics.ll | 100 +++++++++++++++---
4 files changed, 112 insertions(+), 21 deletions(-)
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
index e93e77bd1defd..61593caae1d08 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
@@ -2720,7 +2720,8 @@ void DAGTypeLegalizer::SplitVecRes_MLOAD(MaskedLoadSDNode *MLD,
MachineMemOperand *MMO = DAG.getMachineFunction().getMachineMemOperand(
MLD->getPointerInfo(), MMOFlags, LocationSize::beforeOrAfterPointer(),
- Alignment, MMOMetadata(MLD->getAAInfo(), MLD->getRanges()));
+ Alignment, MMOMetadata(MLD->getAAInfo(), MLD->getRanges(),
+ MLD->getMemCacheHint()));
Lo = DAG.getMaskedLoad(LoVT, dl, Ch, Ptr, Offset, MaskLo, PassThruLo, LoMemVT,
MMO, MLD->getAddressingMode(), ExtType,
@@ -2744,7 +2745,8 @@ void DAGTypeLegalizer::SplitVecRes_MLOAD(MaskedLoadSDNode *MLD,
MMO = DAG.getMachineFunction().getMachineMemOperand(
MPI, MMOFlags, LocationSize::beforeOrAfterPointer(), Alignment,
- MMOMetadata(MLD->getAAInfo(), MLD->getRanges()));
+ MMOMetadata(MLD->getAAInfo(), MLD->getRanges(),
+ MLD->getMemCacheHint()));
Hi = DAG.getMaskedLoad(HiVT, dl, Ch, Ptr, Offset, MaskHi, PassThruHi,
HiMemVT, MMO, MLD->getAddressingMode(), ExtType,
@@ -4648,7 +4650,8 @@ SDValue DAGTypeLegalizer::SplitVecOp_MSTORE(MaskedStoreSDNode *N,
MachineMemOperand *MMO = DAG.getMachineFunction().getMachineMemOperand(
N->getPointerInfo(), MachineMemOperand::MOStore,
LocationSize::beforeOrAfterPointer(), Alignment,
- MMOMetadata(N->getAAInfo(), N->getRanges()));
+ MMOMetadata(N->getAAInfo(), N->getRanges(),
+ N->getMemCacheHint()));
Lo = DAG.getMaskedStore(Ch, DL, DataLo, Ptr, Offset, MaskLo, LoMemVT, MMO,
N->getAddressingMode(), N->isTruncatingStore(),
@@ -4674,7 +4677,8 @@ SDValue DAGTypeLegalizer::SplitVecOp_MSTORE(MaskedStoreSDNode *N,
MMO = DAG.getMachineFunction().getMachineMemOperand(
MPI, MachineMemOperand::MOStore, LocationSize::beforeOrAfterPointer(),
- Alignment, MMOMetadata(N->getAAInfo(), N->getRanges()));
+ Alignment, MMOMetadata(N->getAAInfo(), N->getRanges(),
+ N->getMemCacheHint()));
Hi = DAG.getMaskedStore(Ch, DL, DataHi, Ptr, Offset, MaskHi, HiMemVT, MMO,
N->getAddressingMode(), N->isTruncatingStore(),
diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
index dea008fd252a9..265018a6ff5b5 100644
--- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
@@ -5006,10 +5006,12 @@ void SelectionDAGBuilder::visitMaskedStore(const CallInst &I,
if (I.hasMetadata(LLVMContext::MD_nontemporal))
MMOFlags |= MachineMemOperand::MONonTemporal;
+ const MDNode *MemCacheHint = getMemCacheHintMetadata(I, /*OperandNo=*/1);
+
MachineMemOperand *MMO = DAG.getMachineFunction().getMachineMemOperand(
MachinePointerInfo(PtrOperand), MMOFlags,
LocationSize::upperBound(VT.getStoreSize()), Alignment,
- I.getAAMetadata());
+ MMOMetadata(I.getAAMetadata(), /*Ranges=*/nullptr, MemCacheHint));
const auto &TLI = DAG.getTargetLoweringInfo();
@@ -5153,6 +5155,7 @@ void SelectionDAGBuilder::visitMaskedLoad(const CallInst &I, bool IsExpanding) {
EVT VT = Src0.getValueType();
AAMDNodes AAInfo = I.getAAMetadata();
const MDNode *Ranges = getRangeMetadata(I);
+ const MDNode *MemCacheHint = getMemCacheHintMetadata(I, /*OperandNo=*/0);
// Do not serialize masked loads of constant memory with anything.
MemoryLocation ML = MemoryLocation::getAfter(PtrOperand, AAInfo);
@@ -5169,7 +5172,7 @@ void SelectionDAGBuilder::visitMaskedLoad(const CallInst &I, bool IsExpanding) {
MachineMemOperand *MMO = DAG.getMachineFunction().getMachineMemOperand(
MachinePointerInfo(PtrOperand), MMOFlags,
LocationSize::upperBound(VT.getStoreSize()), Alignment,
- MMOMetadata(AAInfo, Ranges));
+ MMOMetadata(AAInfo, Ranges, MemCacheHint));
const auto &TLI = DAG.getTargetLoweringInfo();
diff --git a/llvm/lib/Transforms/Scalar/ScalarizeMaskedMemIntrin.cpp b/llvm/lib/Transforms/Scalar/ScalarizeMaskedMemIntrin.cpp
index da9ceb4f440e5..87343462778bc 100644
--- a/llvm/lib/Transforms/Scalar/ScalarizeMaskedMemIntrin.cpp
+++ b/llvm/lib/Transforms/Scalar/ScalarizeMaskedMemIntrin.cpp
@@ -186,6 +186,8 @@ static void scalarizeMaskedLoad(const DataLayout &DL, bool HasBranchDivergence,
continue;
Value *Gep = Builder.CreateConstInBoundsGEP1_32(EltTy, Ptr, Idx);
LoadInst *Load = Builder.CreateAlignedLoad(EltTy, Gep, AdjustedAlignVal);
+ Load->setMetadata(LLVMContext::MD_mem_cache_hint,
+ CI->getMetadata(LLVMContext::MD_mem_cache_hint));
VResult = Builder.CreateInsertElement(VResult, Load, Idx);
}
CI->replaceAllUsesWith(VResult);
@@ -267,6 +269,8 @@ static void scalarizeMaskedLoad(const DataLayout &DL, bool HasBranchDivergence,
Builder.SetInsertPoint(CondBlock->getTerminator());
Value *Gep = Builder.CreateConstInBoundsGEP1_32(EltTy, Ptr, Idx);
LoadInst *Load = Builder.CreateAlignedLoad(EltTy, Gep, AdjustedAlignVal);
+ Load->setMetadata(LLVMContext::MD_mem_cache_hint,
+ CI->getMetadata(LLVMContext::MD_mem_cache_hint));
Value *NewVResult = Builder.CreateInsertElement(VResult, Load, Idx);
// Create "else" block, fill it in the next iteration
@@ -352,7 +356,10 @@ static void scalarizeMaskedStore(const DataLayout &DL, bool HasBranchDivergence,
continue;
Value *OneElt = Builder.CreateExtractElement(Src, Idx);
Value *Gep = Builder.CreateConstInBoundsGEP1_32(EltTy, Ptr, Idx);
- Builder.CreateAlignedStore(OneElt, Gep, AdjustedAlignVal);
+ StoreInst *Store =
+ Builder.CreateAlignedStore(OneElt, Gep, AdjustedAlignVal);
+ Store->setMetadata(LLVMContext::MD_mem_cache_hint,
+ CI->getMetadata(LLVMContext::MD_mem_cache_hint));
}
CI->eraseFromParent();
return;
@@ -425,7 +432,10 @@ static void scalarizeMaskedStore(const DataLayout &DL, bool HasBranchDivergence,
Builder.SetInsertPoint(CondBlock->getTerminator());
Value *OneElt = Builder.CreateExtractElement(Src, Idx);
Value *Gep = Builder.CreateConstInBoundsGEP1_32(EltTy, Ptr, Idx);
- Builder.CreateAlignedStore(OneElt, Gep, AdjustedAlignVal);
+ StoreInst *Store =
+ Builder.CreateAlignedStore(OneElt, Gep, AdjustedAlignVal);
+ Store->setMetadata(LLVMContext::MD_mem_cache_hint,
+ CI->getMetadata(LLVMContext::MD_mem_cache_hint));
// Create "else" block, fill it in the next iteration
BasicBlock *NewIfBlock = ThenTerm->getSuccessor(0);
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll b/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
index 19ac0f863a7ac..275d7b1a35f2c 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
@@ -1,6 +1,8 @@
-; 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)" --version 6
-; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | FileCheck %s
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --filter "^\s*(?:\.pragma .*used_bytes_mask|mov\.b64|ld(?:\.[A-Za-z0-9_:]+)*\.global|st(?:\.[A-Za-z0-9_:]+)*\.global|atom(?:\.[A-Za-z0-9_:]+)*\.global)" --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | FileCheck %s --check-prefixes=CHECK,SM80
; RUN: %if ptxas-sm_80 && ptxas-isa-7.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | %ptxas-verify -arch=sm_80 %}
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | FileCheck %s --check-prefixes=CHECK,SM100
+; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | %ptxas-verify -arch=sm_100 %}
; Test !mem.cache_hint metadata on LLVM memory intrinsics.
;
@@ -16,6 +18,8 @@
declare void @llvm.memcpy.p1.p1.i64(ptr addrspace(1), ptr addrspace(1), i64, i1)
declare <2 x i16> @llvm.masked.load.v2i16.p1(ptr addrspace(1), <2 x i1>, <2 x i16>)
declare <4 x i32> @llvm.masked.load.v4i32.p1(ptr addrspace(1), <4 x i1>, <4 x i32>)
+declare <8 x i32> @llvm.masked.load.v8i32.p1(ptr addrspace(1), <8 x i1>, <8 x i32>)
+declare void @llvm.masked.store.v8i32.p1(<8 x i32>, ptr addrspace(1), <8 x i1>)
;-----------------------------------------------------------------------------
; Test memcpy with cache hints on both source and destination
@@ -358,28 +362,85 @@ define void @test_memcpy_all_hints_both(ptr addrspace(1) %dest, ptr addrspace(1)
}
;-----------------------------------------------------------------------------
-; TODO: Preserve cache hints on llvm.masked.load.
-; Masked load pointer operand is operand 0. SelectionDAGBuilder::visitMaskedLoad
-; does not currently thread !mem.cache_hint into the MachineMemOperand, so the
-; generated loads are plain today.
+; Cache hints on masked load/store
;-----------------------------------------------------------------------------
-; TODO: This should eventually lower to ld.global.L1::evict_first.b32.
define <2 x i16> @test_masked_load_l1_hint_v2i16(ptr addrspace(1) %p) {
-; CHECK-LABEL: test_masked_load_l1_hint_v2i16(
-; CHECK: ld.global.b32 %r1, [%rd1];
+; SM80-LABEL: test_masked_load_l1_hint_v2i16(
+; SM80: ld.global.L1::evict_first.b32 %r1, [%rd1];
+;
+; SM100-LABEL: test_masked_load_l1_hint_v2i16(
+; SM100: .pragma "used_bytes_mask 0x3";
+; SM100: ld.global.L1::evict_first.b32 %r1, [%rd1];
%v = call <2 x i16> @llvm.masked.load.v2i16.p1(ptr addrspace(1) align 4 %p, <2 x i1> <i1 true, i1 false>, <2 x i16> poison), !mem.cache_hint !101
ret <2 x i16> %v
}
-; TODO: This should eventually lower to ld.global.L2::cache_hint.v4.b32.
define <4 x i32> @test_masked_load_l2_cache_policy_v4i32(ptr addrspace(1) %p) {
-; CHECK-LABEL: test_masked_load_l2_cache_policy_v4i32(
-; CHECK: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM80-LABEL: test_masked_load_l2_cache_policy_v4i32(
+; SM80: mov.b64 %rd2, 12345;
+; SM80: ld.global.L2::cache_hint.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1], %rd2;
+;
+; SM100-LABEL: test_masked_load_l2_cache_policy_v4i32(
+; SM100: mov.b64 %rd2, 12345;
+; SM100: .pragma "used_bytes_mask 0xff0f";
+; SM100: ld.global.L2::cache_hint.v2.b64 {%rd3, %rd4}, [%rd1], %rd2;
%v = call <4 x i32> @llvm.masked.load.v4i32.p1(ptr addrspace(1) align 16 %p, <4 x i1> <i1 true, i1 false, i1 true, i1 true>, <4 x i32> poison), !mem.cache_hint !102
ret <4 x i32> %v
}
+define <8 x i32> @test_masked_load_split_v8i32(ptr addrspace(1) %p) {
+; SM80-LABEL: test_masked_load_split_v8i32(
+; SM80: ld.global.L1::evict_last.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM80: ld.global.L1::evict_last.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];
+;
+; SM100-LABEL: test_masked_load_split_v8i32(
+; SM100: .pragma "used_bytes_mask 0xf0f0f0f";
+; SM100: ld.global.L1::evict_last.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];
+ %v = call <8 x i32> @llvm.masked.load.v8i32.p1(ptr addrspace(1) align 32 %p, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 true, i1 false, i1 true, i1 false>, <8 x i32> poison), !mem.cache_hint !103
+ ret <8 x i32> %v
+}
+
+define <8 x i32> @test_masked_load_l2_prefetch_v8i32(ptr addrspace(1) %p) {
+; SM80-LABEL: test_masked_load_l2_prefetch_v8i32(
+; SM80: ld.global.L2::256B.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM80: ld.global.L2::256B.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];
+;
+; SM100-LABEL: test_masked_load_l2_prefetch_v8i32(
+; SM100: .pragma "used_bytes_mask 0xf0f0f0f";
+; SM100: ld.global.L2::256B.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];
+ %v = call <8 x i32> @llvm.masked.load.v8i32.p1(ptr addrspace(1) align 32 %p, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 true, i1 false, i1 true, i1 false>, <8 x i32> poison), !mem.cache_hint !105
+ ret <8 x i32> %v
+}
+
+define void @test_masked_store_l2_cache_hint_v8i32(<8 x i32> %v, ptr addrspace(1) %p) {
+; SM80-LABEL: test_masked_store_l2_cache_hint_v8i32(
+; SM80: mov.b64 %rd2, 12345;
+; SM80: st.global.L2::cache_hint.b32 [%rd1], %r5, %rd2;
+; SM80: st.global.L2::cache_hint.b32 [%rd1+8], %r7, %rd2;
+; SM80: st.global.L2::cache_hint.b32 [%rd1+28], %r4, %rd2;
+;
+; SM100-LABEL: test_masked_store_l2_cache_hint_v8i32(
+; SM100: mov.b64 %rd1, 12345;
+; SM100: st.global.L2::cache_hint.v8.b32 [%rd2], {%r5, _, %r7, _, _, _, _, %r4}, %rd1;
+ call void @llvm.masked.store.v8i32.p1(<8 x i32> %v, ptr addrspace(1) align 32 %p, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>), !mem.cache_hint !106
+ ret void
+}
+
+; SM100 supports the 256-bit masked store with l2 eviction. SM80 legalizes to
+; scalar stores, which don't support l2 eviction.
+define void @test_masked_store_l2_eviction_v8i32(<8 x i32> %v, ptr addrspace(1) %p) {
+; SM80-LABEL: test_masked_store_l2_eviction_v8i32(
+; SM80: st.global.b32 [%rd1], %r5;
+; SM80: st.global.b32 [%rd1+8], %r7;
+; SM80: st.global.b32 [%rd1+28], %r4;
+;
+; SM100-LABEL: test_masked_store_l2_eviction_v8i32(
+; SM100: st.global.L2::evict_last.v8.b32 [%rd1], {%r5, _, %r7, _, _, _, _, %r4};
+ call void @llvm.masked.store.v8i32.p1(<8 x i32> %v, ptr addrspace(1) align 32 %p, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>), !mem.cache_hint !104
+ ret void
+}
+
;-----------------------------------------------------------------------------
; Large memcpy tests - verify hints propagate to all expanded load/stores
; LLVM expands memcpy to multiple load/store pairs. Each pair should
@@ -614,9 +675,22 @@ define void @test_memcpy_128bytes_combined(ptr addrspace(1) %dest, ptr addrspace
!214 = !{!"nvvm.l1_eviction", !"first", !"nvvm.l2_eviction", !"last", !"nvvm.l2_prefetch_size", !"256B"}
!215 = !{!"nvvm.l1_eviction", !"last", !"nvvm.l2_eviction", !"first"}
-; Masked load metadata
+; Masked load metadata (pointer operand 0)
!101 = !{i32 0, !216}
!216 = !{!"nvvm.l1_eviction", !"first"}
!102 = !{i32 0, !217}
!217 = !{!"nvvm.l2_cache_hint", i64 12345}
+
+!103 = !{i32 0, !218}
+!218 = !{!"nvvm.l1_eviction", !"last"}
+
+; Masked store metadata (pointer operand 1)
+!104 = !{i32 1, !219}
+!219 = !{!"nvvm.l2_eviction", !"last"}
+
+!105 = !{i32 0, !220}
+!220 = !{!"nvvm.l2_prefetch_size", !"256B"}
+
+!106 = !{i32 1, !221}
+!221 = !{!"nvvm.l2_cache_hint", i64 12345}
>From c3395b95a39d25d73fd429c366ffbc96c6a7f124 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Mon, 10 Aug 2026 14:58:13 +0000
Subject: [PATCH 2/4] format
---
llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp | 11 +++++------
1 file changed, 5 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
index 61593caae1d08..110b5c9c02256 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
@@ -2720,8 +2720,8 @@ void DAGTypeLegalizer::SplitVecRes_MLOAD(MaskedLoadSDNode *MLD,
MachineMemOperand *MMO = DAG.getMachineFunction().getMachineMemOperand(
MLD->getPointerInfo(), MMOFlags, LocationSize::beforeOrAfterPointer(),
- Alignment, MMOMetadata(MLD->getAAInfo(), MLD->getRanges(),
- MLD->getMemCacheHint()));
+ Alignment,
+ MMOMetadata(MLD->getAAInfo(), MLD->getRanges(), MLD->getMemCacheHint()));
Lo = DAG.getMaskedLoad(LoVT, dl, Ch, Ptr, Offset, MaskLo, PassThruLo, LoMemVT,
MMO, MLD->getAddressingMode(), ExtType,
@@ -4650,8 +4650,7 @@ SDValue DAGTypeLegalizer::SplitVecOp_MSTORE(MaskedStoreSDNode *N,
MachineMemOperand *MMO = DAG.getMachineFunction().getMachineMemOperand(
N->getPointerInfo(), MachineMemOperand::MOStore,
LocationSize::beforeOrAfterPointer(), Alignment,
- MMOMetadata(N->getAAInfo(), N->getRanges(),
- N->getMemCacheHint()));
+ MMOMetadata(N->getAAInfo(), N->getRanges(), N->getMemCacheHint()));
Lo = DAG.getMaskedStore(Ch, DL, DataLo, Ptr, Offset, MaskLo, LoMemVT, MMO,
N->getAddressingMode(), N->isTruncatingStore(),
@@ -4677,8 +4676,8 @@ SDValue DAGTypeLegalizer::SplitVecOp_MSTORE(MaskedStoreSDNode *N,
MMO = DAG.getMachineFunction().getMachineMemOperand(
MPI, MachineMemOperand::MOStore, LocationSize::beforeOrAfterPointer(),
- Alignment, MMOMetadata(N->getAAInfo(), N->getRanges(),
- N->getMemCacheHint()));
+ Alignment,
+ MMOMetadata(N->getAAInfo(), N->getRanges(), N->getMemCacheHint()));
Hi = DAG.getMaskedStore(Ch, DL, DataHi, Ptr, Offset, MaskHi, HiMemVT, MMO,
N->getAddressingMode(), N->isTruncatingStore(),
>From caa55cb9bcdc283c8e9822eef0861f0e3569366b Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Mon, 10 Aug 2026 15:08:06 +0000
Subject: [PATCH 3/4] dedup
---
.../CodeGen/NVPTX/cache-hint-intrinsics.ll | 28 ++++++-------------
1 file changed, 9 insertions(+), 19 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll b/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
index 275d7b1a35f2c..5ff100127d505 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
@@ -675,22 +675,12 @@ define void @test_memcpy_128bytes_combined(ptr addrspace(1) %dest, ptr addrspace
!214 = !{!"nvvm.l1_eviction", !"first", !"nvvm.l2_eviction", !"last", !"nvvm.l2_prefetch_size", !"256B"}
!215 = !{!"nvvm.l1_eviction", !"last", !"nvvm.l2_eviction", !"first"}
-; Masked load metadata (pointer operand 0)
-!101 = !{i32 0, !216}
-!216 = !{!"nvvm.l1_eviction", !"first"}
-
-!102 = !{i32 0, !217}
-!217 = !{!"nvvm.l2_cache_hint", i64 12345}
-
-!103 = !{i32 0, !218}
-!218 = !{!"nvvm.l1_eviction", !"last"}
-
-; Masked store metadata (pointer operand 1)
-!104 = !{i32 1, !219}
-!219 = !{!"nvvm.l2_eviction", !"last"}
-
-!105 = !{i32 0, !220}
-!220 = !{!"nvvm.l2_prefetch_size", !"256B"}
-
-!106 = !{i32 1, !221}
-!221 = !{!"nvvm.l2_cache_hint", i64 12345}
+; Masked load metadata
+!101 = !{i32 0, !182}
+!102 = !{i32 0, !184}
+!103 = !{i32 0, !199}
+!105 = !{i32 0, !200}
+
+; Masked store metadata
+!104 = !{i32 1, !191}
+!106 = !{i32 1, !184}
>From 468bbf8dfa4966b833b03610a59a98ff4f8c676b Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Mon, 10 Aug 2026 15:22:18 +0000
Subject: [PATCH 4/4] leave note about prefetch duplication
---
llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll | 4 ++++
1 file changed, 4 insertions(+)
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll b/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
index 5ff100127d505..45b1cfe183aa9 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-intrinsics.ll
@@ -401,6 +401,10 @@ define <8 x i32> @test_masked_load_split_v8i32(ptr addrspace(1) %p) {
ret <8 x i32> %v
}
+; TODO: It makes sense to propagate eviction hints when splitting, but
+; duplicating the prefetch hint might cause a regression. We might need a
+; target hook to refine metadata after splitting. In this case, I think it
+; would be ideal to only preserve the prefetch on the first load.
define <8 x i32> @test_masked_load_l2_prefetch_v8i32(ptr addrspace(1) %p) {
; SM80-LABEL: test_masked_load_l2_prefetch_v8i32(
; SM80: ld.global.L2::256B.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
More information about the llvm-commits
mailing list