llvm/test/CodeGen/NVPTX/masked-load-3xhalf.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | FileCheck %s
; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | %ptxas-verify -arch=sm_100 %}

; This is testing the lowering behavior of this case from LoadStoreVectorizer/NVPTX/4x2xhalf.ll
; where two 3xhalfs are chained together and extended to 8xhalf.
define void @halfx3_extend_chain(ptr align 16 captures(none) %rd0) {
; CHECK-LABEL: halfx3_extend_chain(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<7>;
; CHECK-NEXT:    .reg .b32 %r<12>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [halfx3_extend_chain_param_0];
; CHECK-NEXT:    .pragma "used_bytes_mask 0xfff";
; CHECK-NEXT:    ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
; CHECK-NEXT:    mov.b32 {%rs1, %rs2}, %r3;
; CHECK-NEXT:    mov.b32 {_, %rs3}, %r2;
; CHECK-NEXT:    mov.b32 %r5, {%rs3, %rs1};
; CHECK-NEXT:    mov.b32 %r6, {%rs2, %rs4};
; CHECK-NEXT:    mov.b32 %r7, 0;
; CHECK-NEXT:    max.f16x2 %r8, %r2, %r7;
; CHECK-NEXT:    max.f16x2 %r9, %r1, %r7;
; CHECK-NEXT:    st.b32 [%rd1], %r9;
; CHECK-NEXT:    mov.b32 {%rs5, _}, %r8;
; CHECK-NEXT:    st.b16 [%rd1+4], %rs5;
; CHECK-NEXT:    max.f16x2 %r10, %r6, %r7;
; CHECK-NEXT:    max.f16x2 %r11, %r5, %r7;
; CHECK-NEXT:    st.b32 [%rd1+6], %r11;
; CHECK-NEXT:    mov.b32 {%rs6, _}, %r10;
; CHECK-NEXT:    st.b16 [%rd1+10], %rs6;
; CHECK-NEXT:    ret;
  %load1 = load <3 x half>, ptr %rd0, align 16
  %p1 = fcmp ogt <3 x half> %load1, zeroinitializer
  %s1 = select <3 x i1> %p1, <3 x half> %load1, <3 x half> zeroinitializer
  store <3 x half> %s1, ptr %rd0, align 16
  %in2 = getelementptr half, ptr %rd0, i64 3
  %load2 = load <3 x half>, ptr %in2, align 4
  %p2 = fcmp ogt <3 x half> %load2, zeroinitializer
  %s2 = select <3 x i1> %p2, <3 x half> %load2, <3 x half> zeroinitializer
  store <3 x half> %s2, ptr %in2, align 4
  ret void
}

; This disables the vectorization by reducing the alignment.
define void @halfx3_no_align(ptr align 4 captures(none) %rd0) {
; CHECK-LABEL: halfx3_no_align(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<7>;
; CHECK-NEXT:    .reg .b32 %r<10>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [halfx3_no_align_param_0];
; CHECK-NEXT:    ld.b16 %rs1, [%rd1+4];
; CHECK-NEXT:    mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT:    ld.b32 %r2, [%rd1];
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    max.f16x2 %r4, %r1, %r3;
; CHECK-NEXT:    max.f16x2 %r5, %r2, %r3;
; CHECK-NEXT:    st.b32 [%rd1], %r5;
; CHECK-NEXT:    mov.b32 {%rs3, _}, %r4;
; CHECK-NEXT:    st.b16 [%rd1+4], %rs3;
; CHECK-NEXT:    ld.b16 %rs4, [%rd1+10];
; CHECK-NEXT:    mov.b32 %r6, {%rs4, %rs5};
; CHECK-NEXT:    ld.b32 %r7, [%rd1+6];
; CHECK-NEXT:    max.f16x2 %r8, %r6, %r3;
; CHECK-NEXT:    max.f16x2 %r9, %r7, %r3;
; CHECK-NEXT:    st.b32 [%rd1+6], %r9;
; CHECK-NEXT:    mov.b32 {%rs6, _}, %r8;
; CHECK-NEXT:    st.b16 [%rd1+10], %rs6;
; CHECK-NEXT:    ret;
  %load1 = load <3 x half>, ptr %rd0, align 4
  %p1 = fcmp ogt <3 x half> %load1, zeroinitializer
  %s1 = select <3 x i1> %p1, <3 x half> %load1, <3 x half> zeroinitializer
  store <3 x half> %s1, ptr %rd0, align 4
  %in2 = getelementptr half, ptr %rd0, i64 3
  %load2 = load <3 x half>, ptr %in2, align 4
  %p2 = fcmp ogt <3 x half> %load2, zeroinitializer
  %s2 = select <3 x i1> %p2, <3 x half> %load2, <3 x half> zeroinitializer
  store <3 x half> %s2, ptr %in2, align 4
  ret void
}