llvm/test/CodeGen/NVPTX/setp-bool.ll

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

target triple = "nvptx64-nvidia-cuda"

; Integer, register operands

define i1 @and_i32(i32 %a, i32 %b, i1 %c) {
; CHECK-LABEL: and_i32(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<3>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b8 %rs1, [and_i32_param_2];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    ld.param.b32 %r1, [and_i32_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [and_i32_param_1];
; CHECK-NEXT:    setp.lt.and.s32 %p2, %r1, %r2, %p1;
; CHECK-NEXT:    selp.b32 %r3, -1, 0, %p2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %cmp = icmp slt i32 %a, %b
  %r = and i1 %cmp, %c
  ret i1 %r
}

define i1 @or_i32(i32 %a, i32 %b, i1 %c) {
; CHECK-LABEL: or_i32(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<3>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b8 %rs1, [or_i32_param_2];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    ld.param.b32 %r1, [or_i32_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [or_i32_param_1];
; CHECK-NEXT:    setp.lt.or.s32 %p2, %r1, %r2, %p1;
; CHECK-NEXT:    selp.b32 %r3, -1, 0, %p2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %cmp = icmp slt i32 %a, %b
  %r = or i1 %cmp, %c
  ret i1 %r
}

define i1 @xor_i32(i32 %a, i32 %b, i1 %c) {
; CHECK-LABEL: xor_i32(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<3>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b8 %rs1, [xor_i32_param_2];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    ld.param.b32 %r1, [xor_i32_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [xor_i32_param_1];
; CHECK-NEXT:    setp.lt.xor.s32 %p2, %r1, %r2, %p1;
; CHECK-NEXT:    selp.b32 %r3, -1, 0, %p2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %cmp = icmp slt i32 %a, %b
  %r = xor i1 %cmp, %c
  ret i1 %r
}

define i1 @and_i16(i16 %a, i16 %b, i1 %c) {
; CHECK-LABEL: and_i16(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<3>;
; CHECK-NEXT:    .reg .b16 %rs<5>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b8 %rs1, [and_i16_param_2];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    ld.param.b16 %rs3, [and_i16_param_0];
; CHECK-NEXT:    ld.param.b16 %rs4, [and_i16_param_1];
; CHECK-NEXT:    setp.lt.and.u16 %p2, %rs3, %rs4, %p1;
; CHECK-NEXT:    selp.b32 %r1, -1, 0, %p2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %cmp = icmp ult i16 %a, %b
  %r = and i1 %cmp, %c
  ret i1 %r
}

; Float

define i1 @and_f32(float %a, float %b, i1 %c) {
; CHECK-LABEL: and_f32(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<3>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b8 %rs1, [and_f32_param_2];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    ld.param.b32 %r1, [and_f32_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [and_f32_param_1];
; CHECK-NEXT:    setp.lt.and.f32 %p2, %r1, %r2, %p1;
; CHECK-NEXT:    selp.b32 %r3, -1, 0, %p2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %cmp = fcmp olt float %a, %b
  %r = and i1 %cmp, %c
  ret i1 %r
}

define i1 @or_f64(double %a, double %b, i1 %c) {
; CHECK-LABEL: or_f64(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<3>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b8 %rs1, [or_f64_param_2];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    ld.param.b64 %rd1, [or_f64_param_0];
; CHECK-NEXT:    ld.param.b64 %rd2, [or_f64_param_1];
; CHECK-NEXT:    setp.gt.or.f64 %p2, %rd1, %rd2, %p1;
; CHECK-NEXT:    selp.b32 %r1, -1, 0, %p2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %cmp = fcmp ogt double %a, %b
  %r = or i1 %cmp, %c
  ret i1 %r
}

define i1 @xor_f16(half %a, half %b, i1 %c) {
; CHECK-LABEL: xor_f16(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<3>;
; CHECK-NEXT:    .reg .b16 %rs<5>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b8 %rs1, [xor_f16_param_2];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    ld.param.b16 %rs3, [xor_f16_param_0];
; CHECK-NEXT:    ld.param.b16 %rs4, [xor_f16_param_1];
; CHECK-NEXT:    setp.lt.xor.f16 %p2, %rs3, %rs4, %p1;
; CHECK-NEXT:    selp.b32 %r1, -1, 0, %p2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %cmp = fcmp olt half %a, %b
  %r = xor i1 %cmp, %c
  ret i1 %r
}

define i1 @and_bf16(bfloat %a, bfloat %b, i1 %c) {
; CHECK-LABEL: and_bf16(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<3>;
; CHECK-NEXT:    .reg .b16 %rs<5>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b8 %rs1, [and_bf16_param_2];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    ld.param.b16 %rs3, [and_bf16_param_0];
; CHECK-NEXT:    ld.param.b16 %rs4, [and_bf16_param_1];
; CHECK-NEXT:    setp.lt.and.bf16 %p2, %rs3, %rs4, %p1;
; CHECK-NEXT:    selp.b32 %r1, -1, 0, %p2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %cmp = fcmp olt bfloat %a, %b
  %r = and i1 %cmp, %c
  ret i1 %r
}