[llvm] [SDAG][NVPTX] Cache control metadata support and lowering (PR #204067)
Kerry McLaughlin via llvm-commits
llvm-commits at lists.llvm.org
Fri Jul 3 06:43:40 PDT 2026
================
@@ -0,0 +1,293 @@
+; 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|ld(?:\.[A-Za-z0-9_:]+)*\.L2::cache_hint|st(?:\.[A-Za-z0-9_:]+)*\.L2::cache_hint|atom(?:\.[A-Za-z0-9_:]+)*\.global)" --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | %ptxas-verify %}
+
+; Test L2::cache_hint metadata lowering with constant cache-policy operands.
+
+;-----------------------------------------------------------------------------
+; L2::cache_hint with constant cache-policy operand (metadata-based)
+;-----------------------------------------------------------------------------
+
+define i32 @test_load_cache_hint_i32(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_i32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+ %v = load i32, ptr addrspace(1) %p, !mem.cache_hint !0
+ ret i32 %v
+}
+
+define i64 @test_load_cache_hint_i64(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_i64(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.b64 %rd3, [%rd1], %rd2;
+ %v = load i64, ptr addrspace(1) %p, !mem.cache_hint !1
+ ret i64 %v
+}
+
+define float @test_load_cache_hint_f32(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_f32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+ %v = load float, ptr addrspace(1) %p, !mem.cache_hint !2
+ ret float %v
+}
+
+define void @test_store_cache_hint_i32(ptr addrspace(1) %p, i32 %v) {
+; CHECK-LABEL: test_store_cache_hint_i32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+ store i32 %v, ptr addrspace(1) %p, !mem.cache_hint !3
+ ret void
+}
+
+define void @test_store_cache_hint_i64(ptr addrspace(1) %p, i64 %v) {
+; CHECK-LABEL: test_store_cache_hint_i64(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L2::cache_hint.b64 [%rd1], %rd3, %rd2;
+ store i64 %v, ptr addrspace(1) %p, !mem.cache_hint !4
+ ret void
+}
+
+define void @test_store_cache_hint_f32(ptr addrspace(1) %p, float %v) {
+; CHECK-LABEL: test_store_cache_hint_f32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+ store float %v, ptr addrspace(1) %p, !mem.cache_hint !5
+ ret void
+}
+
+;-----------------------------------------------------------------------------
+; L2::cache_hint with vector types
+;-----------------------------------------------------------------------------
+
+define <2 x i32> @test_load_cache_hint_v2i32(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_v2i32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.v2.b32 {%r1, %r2}, [%rd1], %rd2;
+ %v = load <2 x i32>, ptr addrspace(1) %p, !mem.cache_hint !6
+ ret <2 x i32> %v
+}
+
+define <4 x i32> @test_load_cache_hint_v4i32(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_v4i32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1], %rd2;
+ %v = load <4 x i32>, ptr addrspace(1) %p, !mem.cache_hint !7
+ ret <4 x i32> %v
+}
+
+define <2 x i64> @test_load_cache_hint_v2i64(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_v2i64(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.v2.b64 {%rd3, %rd4}, [%rd1], %rd2;
+ %v = load <2 x i64>, ptr addrspace(1) %p, !mem.cache_hint !8
+ ret <2 x i64> %v
+}
+
+define <2 x float> @test_load_cache_hint_v2f32(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_v2f32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.v2.b32 {%r1, %r2}, [%rd1], %rd2;
+ %v = load <2 x float>, ptr addrspace(1) %p, !mem.cache_hint !9
+ ret <2 x float> %v
+}
+
+define <2 x double> @test_load_cache_hint_v2f64(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_v2f64(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.v2.b64 {%rd3, %rd4}, [%rd1], %rd2;
+ %v = load <2 x double>, ptr addrspace(1) %p, !mem.cache_hint !10
+ ret <2 x double> %v
+}
+
+define void @test_store_cache_hint_v2i32(ptr addrspace(1) %p, <2 x i32> %v) {
+; CHECK-LABEL: test_store_cache_hint_v2i32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L2::cache_hint.v2.b32 [%rd1], {%r1, %r2}, %rd2;
+ store <2 x i32> %v, ptr addrspace(1) %p, !mem.cache_hint !11
+ ret void
+}
+
+define void @test_store_cache_hint_v4i32(ptr addrspace(1) %p, <4 x i32> %v) {
+; CHECK-LABEL: test_store_cache_hint_v4i32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L2::cache_hint.v4.b32 [%rd1], {%r1, %r2, %r3, %r4}, %rd2;
+ store <4 x i32> %v, ptr addrspace(1) %p, !mem.cache_hint !12
+ ret void
+}
+
+define void @test_store_cache_hint_v2i64(ptr addrspace(1) %p, <2 x i64> %v) {
+; CHECK-LABEL: test_store_cache_hint_v2i64(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L2::cache_hint.v2.b64 [%rd1], {%rd3, %rd4}, %rd2;
+ store <2 x i64> %v, ptr addrspace(1) %p, !mem.cache_hint !13
+ ret void
+}
+
+define void @test_store_cache_hint_v2f32(ptr addrspace(1) %p, <2 x float> %v) {
+; CHECK-LABEL: test_store_cache_hint_v2f32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L2::cache_hint.v2.b32 [%rd1], {%r1, %r2}, %rd2;
+ store <2 x float> %v, ptr addrspace(1) %p, !mem.cache_hint !14
+ ret void
+}
+
+define void @test_store_cache_hint_v2f64(ptr addrspace(1) %p, <2 x double> %v) {
+; CHECK-LABEL: test_store_cache_hint_v2f64(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L2::cache_hint.v2.b64 [%rd1], {%rd3, %rd4}, %rd2;
+ store <2 x double> %v, ptr addrspace(1) %p, !mem.cache_hint !15
+ ret void
+}
+
+;-----------------------------------------------------------------------------
+; L2::cache_hint with generic addressing
+;-----------------------------------------------------------------------------
+
+define i32 @test_generic_load_cache_hint_i32(ptr %p) {
+; CHECK-LABEL: test_generic_load_cache_hint_i32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+ %v = load i32, ptr %p, !mem.cache_hint !48
+ ret i32 %v
+}
+
+define void @test_generic_store_cache_hint_i32(ptr %p, i32 %v) {
+; CHECK-LABEL: test_generic_store_cache_hint_i32(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+ store i32 %v, ptr %p, !mem.cache_hint !49
+ ret void
+}
+
+;-----------------------------------------------------------------------------
+; L2::cache_hint combined with other hints (valid qualifiers are preserved)
+;-----------------------------------------------------------------------------
+
+; L2::cache_hint + L1 eviction: both qualifiers should be emitted
+define i32 @test_load_cache_hint_with_l1(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_with_l1(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L1::evict_first.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+ %v = load i32, ptr addrspace(1) %p, !mem.cache_hint !16
+ ret i32 %v
+}
+
+; L2::cache_hint + L2 eviction: cache policy is emitted, scalar L2 eviction is dropped
+define i32 @test_load_cache_hint_with_l2_eviction(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_with_l2_eviction(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+ %v = load i32, ptr addrspace(1) %p, !mem.cache_hint !17
+ ret i32 %v
+}
+
+; L2::cache_hint + L2 prefetch: both qualifiers should be emitted
+define i32 @test_load_cache_hint_with_prefetch(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_with_prefetch(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L2::cache_hint.L2::128B.b32 %r1, [%rd1], %rd2;
+ %v = load i32, ptr addrspace(1) %p, !mem.cache_hint !18
+ ret i32 %v
+}
+
+; L2::cache_hint + all other hints: valid load qualifiers are emitted
+define i32 @test_load_cache_hint_with_all(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_with_all(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L1::evict_last.L2::cache_hint.L2::256B.b32 %r1, [%rd1], %rd2;
+ %v = load i32, ptr addrspace(1) %p, !mem.cache_hint !19
+ ret i32 %v
+}
+
+; Store: L2::cache_hint + L1 eviction
+define void @test_store_cache_hint_with_l1(ptr addrspace(1) %p, i32 %v) {
+; CHECK-LABEL: test_store_cache_hint_with_l1(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L1::evict_unchanged.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+ store i32 %v, ptr addrspace(1) %p, !mem.cache_hint !20
+ ret void
+}
+
+; Store: L2::cache_hint + L1 + L2 eviction (scalar L2 eviction is dropped)
+define void @test_store_cache_hint_with_all(ptr addrspace(1) %p, i32 %v) {
+; CHECK-LABEL: test_store_cache_hint_with_all(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L1::no_allocate.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+ store i32 %v, ptr addrspace(1) %p, !mem.cache_hint !21
+ ret void
+}
+
+; Vector load: L2::cache_hint + L1 eviction
+define <2 x i32> @test_load_cache_hint_v2i32_with_l1(ptr addrspace(1) %p) {
+; CHECK-LABEL: test_load_cache_hint_v2i32_with_l1(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: ld.global.L1::evict_first.L2::cache_hint.v2.b32 {%r1, %r2}, [%rd1], %rd2;
+ %v = load <2 x i32>, ptr addrspace(1) %p, !mem.cache_hint !22
+ ret <2 x i32> %v
+}
+
+; Vector store: cache policy and L1 are emitted; L2 eviction/prefetch are dropped
+define void @test_store_cache_hint_v2i32_with_all(ptr addrspace(1) %p, <2 x i32> %v) {
+; CHECK-LABEL: test_store_cache_hint_v2i32_with_all(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.global.L1::evict_last.L2::cache_hint.v2.b32 [%rd1], {%r1, %r2}, %rd2;
+ store <2 x i32> %v, ptr addrspace(1) %p, !mem.cache_hint !23
+ ret void
+}
+
+;-----------------------------------------------------------------------------
+; Metadata definitions
+;-----------------------------------------------------------------------------
+
+!0 = !{i32 0, !24}
+!24 = !{!"nvvm.l2_cache_hint", i64 12345}
+!1 = !{i32 0, !25}
+!25 = !{!"nvvm.l2_cache_hint", i64 12345}
+!2 = !{i32 0, !26}
+!26 = !{!"nvvm.l2_cache_hint", i64 12345}
+!3 = !{i32 1, !27}
+!27 = !{!"nvvm.l2_cache_hint", i64 12345}
+!4 = !{i32 1, !28}
+!28 = !{!"nvvm.l2_cache_hint", i64 12345}
+!5 = !{i32 1, !29}
+!29 = !{!"nvvm.l2_cache_hint", i64 12345}
+!6 = !{i32 0, !30}
+!30 = !{!"nvvm.l2_cache_hint", i64 12345}
+!7 = !{i32 0, !31}
+!31 = !{!"nvvm.l2_cache_hint", i64 12345}
+!8 = !{i32 0, !32}
+!32 = !{!"nvvm.l2_cache_hint", i64 12345}
+!9 = !{i32 0, !33}
+!33 = !{!"nvvm.l2_cache_hint", i64 12345}
+!10 = !{i32 0, !34}
+!34 = !{!"nvvm.l2_cache_hint", i64 12345}
+!11 = !{i32 1, !35}
+!35 = !{!"nvvm.l2_cache_hint", i64 12345}
+!12 = !{i32 1, !36}
+!36 = !{!"nvvm.l2_cache_hint", i64 12345}
+!13 = !{i32 1, !37}
+!37 = !{!"nvvm.l2_cache_hint", i64 12345}
+!14 = !{i32 1, !38}
+!38 = !{!"nvvm.l2_cache_hint", i64 12345}
+!15 = !{i32 1, !39}
+!39 = !{!"nvvm.l2_cache_hint", i64 12345}
+!16 = !{i32 0, !40}
----------------
kmclaughlin-arm wrote:
Many of these metadata definitions are identical, can they be reduced to one of each type?
https://github.com/llvm/llvm-project/pull/204067
More information about the llvm-commits
mailing list