llvm/test/CodeGen/NVPTX/cache-hint-transforms.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|atom(?:\.[A-Za-z0-9_:]+)*\.global|ld\.param\.b32)" --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 %}

; Test cache hint handling across shared pointers, CSE, forwarding, and legalization.

;-----------------------------------------------------------------------------
; Multiple loads sharing same pointer
;-----------------------------------------------------------------------------

; Non-volatile loads can CSE. Matching cache hints are preserved on the merged
; DAG node, but conflicting hints are dropped.
define i32 @test_cse_loads_same_cache_hint(ptr addrspace(1) %p) {
; CHECK-LABEL: test_cse_loads_same_cache_hint(
; CHECK:    ld.global.L1::evict_first.b32 %r1, [%rd1];
  %v1 = load i32, ptr addrspace(1) %p, !mem.cache_hint !2
  %v2 = load i32, ptr addrspace(1) %p, !mem.cache_hint !2
  %sum = add i32 %v1, %v2
  ret i32 %sum
}

define i32 @test_cse_loads_conflicting_cache_hints(ptr addrspace(1) %p) {
; CHECK-LABEL: test_cse_loads_conflicting_cache_hints(
; CHECK:    ld.global.b32 %r1, [%rd1];
  %v1 = load i32, ptr addrspace(1) %p, !mem.cache_hint !2
  %v2 = load i32, ptr addrspace(1) %p, !mem.cache_hint !3
  %sum = add i32 %v1, %v2
  ret i32 %sum
}

; NVPTXForwardParams rewrites eligible byval loads to ld.param. Since cache
; hints are not allowed on ld.param, we must drop them.
define i32 @test_forward_param_drops_l1_hint(ptr byval(i32) %a) {
; CHECK-LABEL: test_forward_param_drops_l1_hint(
; CHECK:    ld.param.b32 %r1, [test_forward_param_drops_l1_hint_param_0];
  %v = load i32, ptr %a, !mem.cache_hint !4
  ret i32 %v
}

define i32 @test_forward_param_drops_l2_cache_hint(ptr byval(i32) %a) {
; CHECK-LABEL: test_forward_param_drops_l2_cache_hint(
; CHECK:    ld.param.b32 %r1, [test_forward_param_drops_l2_cache_hint_param_0];
  %v = load i32, ptr %a, !mem.cache_hint !5
  ret i32 %v
}

;-----------------------------------------------------------------------------
; Valid edge cases
;-----------------------------------------------------------------------------

; Test with custom hint key order - should still work
define i32 @test_load_reordered_metadata(ptr addrspace(1) %p) {
; CHECK-LABEL: test_load_reordered_metadata(
; CHECK:    ld.global.L1::evict_last.b32 %r1, [%rd1];
  %v = load i32, ptr addrspace(1) %p, !mem.cache_hint !6
  ret i32 %v
}

;-----------------------------------------------------------------------------
; nvvm.l2_cache_hint with alternate integer value width
;-----------------------------------------------------------------------------

; nvvm.l2_cache_hint with i32 instead of i64 - should still work
; as mdconst::dyn_extract<ConstantInt> accepts any integer type
define i32 @test_load_cache_hint_i32_value(ptr addrspace(1) %p) {
; CHECK-LABEL: test_load_cache_hint_i32_value(
; 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 !7
  ret i32 %v
}

; Test "normal" eviction - should not emit any qualifier (default behavior)
define i32 @test_load_l1_normal(ptr addrspace(1) %p) {
; CHECK-LABEL: test_load_l1_normal(
; CHECK:    ld.global.b32 %r1, [%rd1];
  %v = load i32, ptr addrspace(1) %p, !mem.cache_hint !8
  ret i32 %v
}

define i32 @test_load_l2_normal(ptr addrspace(1) %p) {
; CHECK-LABEL: test_load_l2_normal(
; CHECK:    ld.global.b32 %r1, [%rd1];
  %v = load i32, ptr addrspace(1) %p, !mem.cache_hint !9
  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};
  %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>
  store <8 x i32> %c, ptr addrspace(1) %p, align 16, !mem.cache_hint !10
  ret void
}

;-----------------------------------------------------------------------------
; 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.
;-----------------------------------------------------------------------------

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];
;
; 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];
  %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};
;
; 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};
  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];
  %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;
;
; 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;
  store <3 x i64> %v, ptr addrspace(1) %p, align 8, !mem.cache_hint !14
  ret void
}

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

!0 = !{i32 0, !15}
!15 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l1_eviction", !"last"}
!1 = !{i32 0, !16}
!16 = !{!"nvvm.l2_cache_hint", i64 12345, !"nvvm.l1_eviction", !"first"}
!2 = !{i32 0, !17}
!17 = !{!"nvvm.l1_eviction", !"first"}
!3 = !{i32 0, !18}
!18 = !{!"nvvm.l1_eviction", !"last"}
!4 = !{i32 0, !19}
!19 = !{!"nvvm.l1_eviction", !"first"}
!5 = !{i32 0, !20}
!20 = !{!"nvvm.l2_cache_hint", i64 12345}
!6 = !{i32 0, !21}
!21 = !{!"nvvm.l1_eviction", !"last", !"nvvm.l2_eviction", !"first"}
!7 = !{i32 0, !22}
!22 = !{!"nvvm.l2_cache_hint", i32 12345}
!8 = !{i32 0, !23}
!23 = !{!"nvvm.l1_eviction", !"normal"}
!9 = !{i32 0, !24}
!24 = !{!"nvvm.l2_eviction", !"normal"}
!10 = !{i32 1, !25}
!25 = !{!"nvvm.l2_cache_hint", i64 12345}
!11 = !{i32 0, !26}
!26 = !{!"nvvm.l2_cache_hint", i64 12345}
!12 = !{i32 1, !27}
!27 = !{!"nvvm.l2_cache_hint", i64 12345}
!13 = !{i32 0, !28}
!28 = !{!"nvvm.l1_eviction", !"last"}
!14 = !{i32 1, !29}
!29 = !{!"nvvm.l2_eviction", !"first"}