|
| 1 | +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 |
| 2 | +; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s |
| 3 | +; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 | %ptxas-verify %} |
| 4 | + |
| 5 | +target triple = "nvptx64-nvidia-cuda" |
| 6 | + |
| 7 | +define i1 @and_ord(float %a, float %b) { |
| 8 | +; CHECK-LABEL: and_ord( |
| 9 | +; CHECK: { |
| 10 | +; CHECK-NEXT: .reg .pred %p<4>; |
| 11 | +; CHECK-NEXT: .reg .b32 %r<2>; |
| 12 | +; CHECK-NEXT: .reg .f32 %f<3>; |
| 13 | +; CHECK-EMPTY: |
| 14 | +; CHECK-NEXT: // %bb.0: |
| 15 | +; CHECK-NEXT: ld.param.f32 %f1, [and_ord_param_0]; |
| 16 | +; CHECK-NEXT: setp.num.f32 %p1, %f1, %f1; |
| 17 | +; CHECK-NEXT: ld.param.f32 %f2, [and_ord_param_1]; |
| 18 | +; CHECK-NEXT: setp.num.f32 %p2, %f2, %f2; |
| 19 | +; CHECK-NEXT: and.pred %p3, %p1, %p2; |
| 20 | +; CHECK-NEXT: selp.b32 %r1, 1, 0, %p3; |
| 21 | +; CHECK-NEXT: st.param.b32 [func_retval0], %r1; |
| 22 | +; CHECK-NEXT: ret; |
| 23 | + %c = fcmp ord float %a, 0.0 |
| 24 | + %d = fcmp ord float %b, 0.0 |
| 25 | + %e = and i1 %c, %d |
| 26 | + ret i1 %e |
| 27 | +} |
| 28 | + |
| 29 | +define i1 @or_uno(float %a, float %b) { |
| 30 | +; CHECK-LABEL: or_uno( |
| 31 | +; CHECK: { |
| 32 | +; CHECK-NEXT: .reg .pred %p<4>; |
| 33 | +; CHECK-NEXT: .reg .b32 %r<2>; |
| 34 | +; CHECK-NEXT: .reg .f32 %f<3>; |
| 35 | +; CHECK-EMPTY: |
| 36 | +; CHECK-NEXT: // %bb.0: |
| 37 | +; CHECK-NEXT: ld.param.f32 %f1, [or_uno_param_0]; |
| 38 | +; CHECK-NEXT: setp.nan.f32 %p1, %f1, %f1; |
| 39 | +; CHECK-NEXT: ld.param.f32 %f2, [or_uno_param_1]; |
| 40 | +; CHECK-NEXT: setp.nan.f32 %p2, %f2, %f2; |
| 41 | +; CHECK-NEXT: or.pred %p3, %p1, %p2; |
| 42 | +; CHECK-NEXT: selp.b32 %r1, 1, 0, %p3; |
| 43 | +; CHECK-NEXT: st.param.b32 [func_retval0], %r1; |
| 44 | +; CHECK-NEXT: ret; |
| 45 | + %c = fcmp uno float %a, 0.0 |
| 46 | + %d = fcmp uno float %b, 0.0 |
| 47 | + %e = or i1 %c, %d |
| 48 | + ret i1 %e |
| 49 | +} |
0 commit comments