1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s | FileCheck %s 3; RUN: %if ptxas %{ llc < %s | %ptxas-verify %} 4 5target triple = "nvptx64-nvidia-cuda" 6 7define i32 @flo_1(i32 %a) { 8; CHECK-LABEL: flo_1( 9; CHECK: { 10; CHECK-NEXT: .reg .b32 %r<3>; 11; CHECK-EMPTY: 12; CHECK-NEXT: // %bb.0: 13; CHECK-NEXT: ld.param.u32 %r1, [flo_1_param_0]; 14; CHECK-NEXT: bfind.s32 %r2, %r1; 15; CHECK-NEXT: st.param.b32 [func_retval0], %r2; 16; CHECK-NEXT: ret; 17 %r = call i32 @llvm.nvvm.flo.s.i32(i32 %a, i1 false) 18 ret i32 %r 19} 20 21 22define i32 @flo_2(i32 %a) { 23; CHECK-LABEL: flo_2( 24; CHECK: { 25; CHECK-NEXT: .reg .b32 %r<3>; 26; CHECK-EMPTY: 27; CHECK-NEXT: // %bb.0: 28; CHECK-NEXT: ld.param.u32 %r1, [flo_2_param_0]; 29; CHECK-NEXT: bfind.shiftamt.s32 %r2, %r1; 30; CHECK-NEXT: st.param.b32 [func_retval0], %r2; 31; CHECK-NEXT: ret; 32 %r = call i32 @llvm.nvvm.flo.s.i32(i32 %a, i1 true) 33 ret i32 %r 34} 35 36define i32 @flo_3(i32 %a) { 37; CHECK-LABEL: flo_3( 38; CHECK: { 39; CHECK-NEXT: .reg .b32 %r<3>; 40; CHECK-EMPTY: 41; CHECK-NEXT: // %bb.0: 42; CHECK-NEXT: ld.param.u32 %r1, [flo_3_param_0]; 43; CHECK-NEXT: bfind.u32 %r2, %r1; 44; CHECK-NEXT: st.param.b32 [func_retval0], %r2; 45; CHECK-NEXT: ret; 46 %r = call i32 @llvm.nvvm.flo.u.i32(i32 %a, i1 false) 47 ret i32 %r 48} 49 50 51define i32 @flo_4(i32 %a) { 52; CHECK-LABEL: flo_4( 53; CHECK: { 54; CHECK-NEXT: .reg .b32 %r<3>; 55; CHECK-EMPTY: 56; CHECK-NEXT: // %bb.0: 57; CHECK-NEXT: ld.param.u32 %r1, [flo_4_param_0]; 58; CHECK-NEXT: bfind.shiftamt.u32 %r2, %r1; 59; CHECK-NEXT: st.param.b32 [func_retval0], %r2; 60; CHECK-NEXT: ret; 61 %r = call i32 @llvm.nvvm.flo.u.i32(i32 %a, i1 true) 62 ret i32 %r 63} 64 65 66 67define i32 @flo_5(i64 %a) { 68; CHECK-LABEL: flo_5( 69; CHECK: { 70; CHECK-NEXT: .reg .b32 %r<2>; 71; CHECK-NEXT: .reg .b64 %rd<2>; 72; CHECK-EMPTY: 73; CHECK-NEXT: // %bb.0: 74; CHECK-NEXT: ld.param.u64 %rd1, [flo_5_param_0]; 75; CHECK-NEXT: bfind.s64 %r1, %rd1; 76; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 77; CHECK-NEXT: ret; 78 %r = call i32 @llvm.nvvm.flo.s.i64(i64 %a, i1 false) 79 ret i32 %r 80} 81 82 83define i32 @flo_6(i64 %a) { 84; CHECK-LABEL: flo_6( 85; CHECK: { 86; CHECK-NEXT: .reg .b32 %r<2>; 87; CHECK-NEXT: .reg .b64 %rd<2>; 88; CHECK-EMPTY: 89; CHECK-NEXT: // %bb.0: 90; CHECK-NEXT: ld.param.u64 %rd1, [flo_6_param_0]; 91; CHECK-NEXT: bfind.shiftamt.s64 %r1, %rd1; 92; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 93; CHECK-NEXT: ret; 94 %r = call i32 @llvm.nvvm.flo.s.i64(i64 %a, i1 true) 95 ret i32 %r 96} 97 98define i32 @flo_7(i64 %a) { 99; CHECK-LABEL: flo_7( 100; CHECK: { 101; CHECK-NEXT: .reg .b32 %r<2>; 102; CHECK-NEXT: .reg .b64 %rd<2>; 103; CHECK-EMPTY: 104; CHECK-NEXT: // %bb.0: 105; CHECK-NEXT: ld.param.u64 %rd1, [flo_7_param_0]; 106; CHECK-NEXT: bfind.u64 %r1, %rd1; 107; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 108; CHECK-NEXT: ret; 109 %r = call i32 @llvm.nvvm.flo.u.i64(i64 %a, i1 false) 110 ret i32 %r 111} 112 113 114define i32 @flo_8(i64 %a) { 115; CHECK-LABEL: flo_8( 116; CHECK: { 117; CHECK-NEXT: .reg .b32 %r<2>; 118; CHECK-NEXT: .reg .b64 %rd<2>; 119; CHECK-EMPTY: 120; CHECK-NEXT: // %bb.0: 121; CHECK-NEXT: ld.param.u64 %rd1, [flo_8_param_0]; 122; CHECK-NEXT: bfind.shiftamt.u64 %r1, %rd1; 123; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 124; CHECK-NEXT: ret; 125 %r = call i32 @llvm.nvvm.flo.u.i64(i64 %a, i1 true) 126 ret i32 %r 127} 128 129declare i32 @llvm.nvvm.flo.s.i32(i32, i1) 130declare i32 @llvm.nvvm.flo.u.i32(i32, i1) 131declare i32 @llvm.nvvm.flo.s.i64(i64, i1) 132declare i32 @llvm.nvvm.flo.u.i64(i64, i1) 133