82 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_50 | FileCheck %s3; RUN: %if ptxas-sm_50 %{ llc < %s -mcpu=sm_50 | %ptxas-verify -arch=sm_50 %}4 5target triple = "nvptx64-nvidia-cuda"6 7define float @uitofp_trunc_nuw(i32 %x, i32 %y) {8; CHECK-LABEL: uitofp_trunc_nuw(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<5>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b32 %r1, [uitofp_trunc_nuw_param_0];14; CHECK-NEXT: ld.param.b32 %r2, [uitofp_trunc_nuw_param_1];15; CHECK-NEXT: add.s32 %r3, %r1, %r2;16; CHECK-NEXT: cvt.rn.f32.u32 %r4, %r3;17; CHECK-NEXT: st.param.b32 [func_retval0], %r4;18; CHECK-NEXT: ret;19 %v = add i32 %x, %y20 %t = trunc nuw i32 %v to i1621 %f = uitofp i16 %t to float22 ret float %f23}24 25define float @sitofp_trunc_nsw(i32 %x, i32 %y) {26; CHECK-LABEL: sitofp_trunc_nsw(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<5>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.param.b32 %r1, [sitofp_trunc_nsw_param_0];32; CHECK-NEXT: ld.param.b32 %r2, [sitofp_trunc_nsw_param_1];33; CHECK-NEXT: add.s32 %r3, %r1, %r2;34; CHECK-NEXT: cvt.rn.f32.s32 %r4, %r3;35; CHECK-NEXT: st.param.b32 [func_retval0], %r4;36; CHECK-NEXT: ret;37 %v = add i32 %x, %y38 %t = trunc nsw i32 %v to i1639 %f = sitofp i16 %t to float40 ret float %f41}42 43define float @uitofp_trunc_nsw(i32 %x, i32 %y) {44; CHECK-LABEL: uitofp_trunc_nsw(45; CHECK: {46; CHECK-NEXT: .reg .b16 %rs<2>;47; CHECK-NEXT: .reg .b32 %r<5>;48; CHECK-EMPTY:49; CHECK-NEXT: // %bb.0:50; CHECK-NEXT: ld.param.b32 %r1, [uitofp_trunc_nsw_param_0];51; CHECK-NEXT: ld.param.b32 %r2, [uitofp_trunc_nsw_param_1];52; CHECK-NEXT: add.s32 %r3, %r1, %r2;53; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;54; CHECK-NEXT: cvt.rn.f32.u16 %r4, %rs1;55; CHECK-NEXT: st.param.b32 [func_retval0], %r4;56; CHECK-NEXT: ret;57 %v = add i32 %x, %y58 %t = trunc nsw i32 %v to i1659 %f = uitofp i16 %t to float60 ret float %f61}62 63define float @sitofp_trunc_nuw(i32 %x, i32 %y) {64; CHECK-LABEL: sitofp_trunc_nuw(65; CHECK: {66; CHECK-NEXT: .reg .b16 %rs<2>;67; CHECK-NEXT: .reg .b32 %r<5>;68; CHECK-EMPTY:69; CHECK-NEXT: // %bb.0:70; CHECK-NEXT: ld.param.b32 %r1, [sitofp_trunc_nuw_param_0];71; CHECK-NEXT: ld.param.b32 %r2, [sitofp_trunc_nuw_param_1];72; CHECK-NEXT: add.s32 %r3, %r1, %r2;73; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;74; CHECK-NEXT: cvt.rn.f32.s16 %r4, %rs1;75; CHECK-NEXT: st.param.b32 [func_retval0], %r4;76; CHECK-NEXT: ret;77 %v = add i32 %x, %y78 %t = trunc nuw i32 %v to i1679 %f = sitofp i16 %t to float80 ret float %f81}82