llvm/test/CodeGen/NVPTX/cache-hint-cache-policy.ll

; 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-sm_80 && ptxas-isa-7.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | %ptxas-verify -arch=sm_80 %}

; 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 !0
  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 !0
  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 !1
  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 !1
  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 !1
  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 !0
  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 !0
  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 !0
  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 !0
  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 !0
  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 !1
  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 !1
  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 !1
  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 !1
  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 !1
  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 !0
  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 !1
  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 !3
  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 !5
  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 !7
  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 !9
  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 !11
  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 !13
  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 !15
  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 !17
  ret void
}

;-----------------------------------------------------------------------------
; Metadata definitions
;-----------------------------------------------------------------------------

!0 = !{i32 0, !2}
!1 = !{i32 1, !2}
!2 = !{!"nvvm.l2_cache_hint", i64 12345}
!3 = !{i32 0, !4}
!4 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l1_eviction", !"first"}
!5 = !{i32 0, !6}
!6 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l2_eviction", !"last"}
!7 = !{i32 0, !8}
!8 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l2_prefetch_size", !"128B"}
!9 = !{i32 0, !10}
!10 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l1_eviction", !"last", !"nvvm.l2_eviction", !"first", !"nvvm.l2_prefetch_size", !"256B"}
!11 = !{i32 1, !12}
!12 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l1_eviction", !"unchanged"}
!13 = !{i32 1, !14}
!14 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l1_eviction", !"no_allocate", !"nvvm.l2_eviction", !"last"}
!15 = !{i32 0, !16}
!16 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l1_eviction", !"first"}
!17 = !{i32 1, !18}
!18 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l1_eviction", !"last", !"nvvm.l2_eviction", !"first", !"nvvm.l2_prefetch_size", !"64B"}