131 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 | FileCheck %s3; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 | FileCheck %s4; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 | FileCheck %s5 6define i8 @cvt_u8_f32(float %x) {7; CHECK-LABEL: cvt_u8_f32(8; CHECK: {9; CHECK-NEXT: .reg .b16 %rs<2>;10; CHECK-NEXT: .reg .b32 %r<3>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b32 %r1, [cvt_u8_f32_param_0];14; CHECK-NEXT: cvt.rzi.u16.f32 %rs1, %r1;15; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;16; CHECK-NEXT: st.param.b32 [func_retval0], %r2;17; CHECK-NEXT: ret;18 %a = fptoui float %x to i819 ret i8 %a20}21 22define i8 @cvt_u8_f64(double %x) {23; CHECK-LABEL: cvt_u8_f64(24; CHECK: {25; CHECK-NEXT: .reg .b16 %rs<2>;26; CHECK-NEXT: .reg .b32 %r<2>;27; CHECK-NEXT: .reg .b64 %rd<2>;28; CHECK-EMPTY:29; CHECK-NEXT: // %bb.0:30; CHECK-NEXT: ld.param.b64 %rd1, [cvt_u8_f64_param_0];31; CHECK-NEXT: cvt.rzi.u16.f64 %rs1, %rd1;32; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;33; CHECK-NEXT: st.param.b32 [func_retval0], %r1;34; CHECK-NEXT: ret;35 %a = fptoui double %x to i836 ret i8 %a37}38 39define float @cvt_f32_i8(i8 %x) {40; CHECK-LABEL: cvt_f32_i8(41; CHECK: {42; CHECK-NEXT: .reg .b16 %rs<2>;43; CHECK-NEXT: .reg .b32 %r<2>;44; CHECK-EMPTY:45; CHECK-NEXT: // %bb.0:46; CHECK-NEXT: ld.param.b8 %rs1, [cvt_f32_i8_param_0];47; CHECK-NEXT: cvt.rn.f32.u16 %r1, %rs1;48; CHECK-NEXT: st.param.b32 [func_retval0], %r1;49; CHECK-NEXT: ret;50 %a = uitofp i8 %x to float51 ret float %a52}53 54define double @cvt_f64_i8(i8 %x) {55; CHECK-LABEL: cvt_f64_i8(56; CHECK: {57; CHECK-NEXT: .reg .b16 %rs<2>;58; CHECK-NEXT: .reg .b64 %rd<2>;59; CHECK-EMPTY:60; CHECK-NEXT: // %bb.0:61; CHECK-NEXT: ld.param.b8 %rs1, [cvt_f64_i8_param_0];62; CHECK-NEXT: cvt.rn.f64.u16 %rd1, %rs1;63; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;64; CHECK-NEXT: ret;65 %a = uitofp i8 %x to double66 ret double %a67}68 69define float @cvt_f32_s8(i8 %x) {70; CHECK-LABEL: cvt_f32_s8(71; CHECK: {72; CHECK-NEXT: .reg .b16 %rs<2>;73; CHECK-NEXT: .reg .b32 %r<2>;74; CHECK-EMPTY:75; CHECK-NEXT: // %bb.0:76; CHECK-NEXT: ld.param.s8 %rs1, [cvt_f32_s8_param_0];77; CHECK-NEXT: cvt.rn.f32.s16 %r1, %rs1;78; CHECK-NEXT: st.param.b32 [func_retval0], %r1;79; CHECK-NEXT: ret;80 %a = sitofp i8 %x to float81 ret float %a82}83 84define double @cvt_f64_s8(i8 %x) {85; CHECK-LABEL: cvt_f64_s8(86; CHECK: {87; CHECK-NEXT: .reg .b16 %rs<2>;88; CHECK-NEXT: .reg .b64 %rd<2>;89; CHECK-EMPTY:90; CHECK-NEXT: // %bb.0:91; CHECK-NEXT: ld.param.s8 %rs1, [cvt_f64_s8_param_0];92; CHECK-NEXT: cvt.rn.f64.s16 %rd1, %rs1;93; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;94; CHECK-NEXT: ret;95 %a = sitofp i8 %x to double96 ret double %a97}98 99define i8 @cvt_s8_f32(float %x) {100; CHECK-LABEL: cvt_s8_f32(101; CHECK: {102; CHECK-NEXT: .reg .b16 %rs<2>;103; CHECK-NEXT: .reg .b32 %r<3>;104; CHECK-EMPTY:105; CHECK-NEXT: // %bb.0:106; CHECK-NEXT: ld.param.b32 %r1, [cvt_s8_f32_param_0];107; CHECK-NEXT: cvt.rzi.s16.f32 %rs1, %r1;108; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;109; CHECK-NEXT: st.param.b32 [func_retval0], %r2;110; CHECK-NEXT: ret;111 %a = fptosi float %x to i8112 ret i8 %a113}114 115define i8 @cvt_s8_f64(double %x) {116; CHECK-LABEL: cvt_s8_f64(117; CHECK: {118; CHECK-NEXT: .reg .b16 %rs<2>;119; CHECK-NEXT: .reg .b32 %r<2>;120; CHECK-NEXT: .reg .b64 %rd<2>;121; CHECK-EMPTY:122; CHECK-NEXT: // %bb.0:123; CHECK-NEXT: ld.param.b64 %rd1, [cvt_s8_f64_param_0];124; CHECK-NEXT: cvt.rzi.s16.f64 %rs1, %rd1;125; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;126; CHECK-NEXT: st.param.b32 [func_retval0], %r1;127; CHECK-NEXT: ret;128 %a = fptosi double %x to i8129 ret i8 %a130}131