133 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 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.b32 %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 %r19}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.b32 %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 %r34}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.b32 %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 %r48}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.b32 %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 %r63}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.b64 %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 %r80}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.b64 %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 %r96}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.b64 %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 %r111}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.b64 %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 %r127}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