44 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s | FileCheck %s3; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}4 5target triple = "nvptx64-nvidia-cuda"6 7define i1 @and_ord(float %a, float %b) {8; CHECK-LABEL: and_ord(9; CHECK: {10; CHECK-NEXT: .reg .pred %p<2>;11; CHECK-NEXT: .reg .b32 %r<4>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b32 %r1, [and_ord_param_0];15; CHECK-NEXT: ld.param.b32 %r2, [and_ord_param_1];16; CHECK-NEXT: setp.num.f32 %p1, %r1, %r2;17; CHECK-NEXT: selp.b32 %r3, -1, 0, %p1;18; CHECK-NEXT: st.param.b32 [func_retval0], %r3;19; CHECK-NEXT: ret;20 %c = fcmp ord float %a, 0.021 %d = fcmp ord float %b, 0.022 %e = and i1 %c, %d23 ret i1 %e24}25 26define i1 @or_uno(float %a, float %b) {27; CHECK-LABEL: or_uno(28; CHECK: {29; CHECK-NEXT: .reg .pred %p<2>;30; CHECK-NEXT: .reg .b32 %r<4>;31; CHECK-EMPTY:32; CHECK-NEXT: // %bb.0:33; CHECK-NEXT: ld.param.b32 %r1, [or_uno_param_0];34; CHECK-NEXT: ld.param.b32 %r2, [or_uno_param_1];35; CHECK-NEXT: setp.nan.f32 %p1, %r1, %r2;36; CHECK-NEXT: selp.b32 %r3, -1, 0, %p1;37; CHECK-NEXT: st.param.b32 [func_retval0], %r3;38; CHECK-NEXT: ret;39 %c = fcmp uno float %a, 0.040 %d = fcmp uno float %b, 0.041 %e = or i1 %c, %d42 ret i1 %e43}44