llvm/test/CodeGen/NVPTX/convert-fp-i128.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx | FileCheck %s
; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s
; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 | %ptxas-verify %}

; Regression test for 128-bit integer/floating-point conversions.
;
; PTX does not have direct conversion instructions between fp values and
; s128/u128 values. These conversions must be expanded before instruction
; selection. In particular, they should be expanded inline rather than lowered
; to compiler-rt helper calls.

define i128 @cvt_u128_f16(half %x) {
; CHECK-LABEL: cvt_u128_f16(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b16 %rs1, [cvt_u128_f16_param_0];
; CHECK-NEXT:    cvt.rzi.u32.f16 %r1, %rs1;
; CHECK-NEXT:    cvt.u64.u32 %rd1, %r1;
; CHECK-NEXT:    st.param.v2.b64 [func_retval0], {%rd1, 0};
; CHECK-NEXT:    ret;
  %a = fptoui half %x to i128
  ret i128 %a
}

define i128 @cvt_u128_f32(float %x) {
; CHECK-LABEL: cvt_u128_f32(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<4>;
; CHECK-NEXT:    .reg .b32 %r<10>;
; CHECK-NEXT:    .reg .b64 %rd<6>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %fp-to-i-entry
; CHECK-NEXT:    ld.param.b32 %r3, [cvt_u128_f32_param_0];
; CHECK-NEXT:    bfe.u32 %r1, %r3, 23, 8;
; CHECK-NEXT:    setp.lt.u32 %p1, %r1, 127;
; CHECK-NEXT:    mov.b64 %rd5, 0;
; CHECK-NEXT:    mov.b64 %rd4, %rd5;
; CHECK-NEXT:    @%p1 bra $L__BB1_4;
; CHECK-NEXT:  // %bb.1: // %fp-to-i-if-check.exp.size
; CHECK-NEXT:    and.b32 %r4, %r3, 8388607;
; CHECK-NEXT:    or.b32 %r2, %r4, 8388608;
; CHECK-NEXT:    setp.gt.u32 %p2, %r1, 149;
; CHECK-NEXT:    @%p2 bra $L__BB1_3;
; CHECK-NEXT:  // %bb.2: // %fp-to-i-if-exp.small
; CHECK-NEXT:    sub.s32 %r8, 150, %r1;
; CHECK-NEXT:    shr.u32 %r9, %r2, %r8;
; CHECK-NEXT:    cvt.u64.u32 %rd4, %r9;
; CHECK-NEXT:    bra.uni $L__BB1_4;
; CHECK-NEXT:  $L__BB1_3: // %fp-to-i-if-exp.large
; CHECK-NEXT:    add.s32 %r5, %r1, -150;
; CHECK-NEXT:    cvt.u64.u32 %rd1, %r2;
; CHECK-NEXT:    sub.s32 %r6, 214, %r1;
; CHECK-NEXT:    shr.u64 %rd2, %rd1, %r6;
; CHECK-NEXT:    add.s32 %r7, %r1, -214;
; CHECK-NEXT:    shl.b64 %rd3, %rd1, %r7;
; CHECK-NEXT:    setp.gt.s32 %p3, %r5, 63;
; CHECK-NEXT:    selp.b64 %rd5, %rd3, %rd2, %p3;
; CHECK-NEXT:    shl.b64 %rd4, %rd1, %r5;
; CHECK-NEXT:  $L__BB1_4: // %fp-to-i-cleanup
; CHECK-NEXT:    st.param.v2.b64 [func_retval0], {%rd4, %rd5};
; CHECK-NEXT:    ret;
  %a = fptoui float %x to i128
  ret i128 %a
}

define i128 @cvt_u128_f64(double %x) {
; CHECK-LABEL: cvt_u128_f64(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<4>;
; CHECK-NEXT:    .reg .b32 %r<7>;
; CHECK-NEXT:    .reg .b64 %rd<10>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %fp-to-i-entry
; CHECK-NEXT:    ld.param.b64 %rd3, [cvt_u128_f64_param_0];
; CHECK-NEXT:    shr.u64 %rd4, %rd3, 52;
; CHECK-NEXT:    and.b64 %rd1, %rd4, 2047;
; CHECK-NEXT:    setp.lt.u64 %p1, %rd1, 1023;
; CHECK-NEXT:    mov.b64 %rd9, 0;
; CHECK-NEXT:    mov.b64 %rd8, %rd9;
; CHECK-NEXT:    @%p1 bra $L__BB2_4;
; CHECK-NEXT:  // %bb.1: // %fp-to-i-if-check.exp.size
; CHECK-NEXT:    and.b64 %rd5, %rd3, 4503599627370495;
; CHECK-NEXT:    or.b64 %rd2, %rd5, 4503599627370496;
; CHECK-NEXT:    setp.gt.u64 %p2, %rd1, 1074;
; CHECK-NEXT:    @%p2 bra $L__BB2_3;
; CHECK-NEXT:  // %bb.2: // %fp-to-i-if-exp.small
; CHECK-NEXT:    cvt.u32.u64 %r5, %rd1;
; CHECK-NEXT:    sub.s32 %r6, 1075, %r5;
; CHECK-NEXT:    shr.u64 %rd8, %rd2, %r6;
; CHECK-NEXT:    bra.uni $L__BB2_4;
; CHECK-NEXT:  $L__BB2_3: // %fp-to-i-if-exp.large
; CHECK-NEXT:    cvt.u32.u64 %r1, %rd1;
; CHECK-NEXT:    sub.s32 %r2, 1139, %r1;
; CHECK-NEXT:    shr.u64 %rd6, %rd2, %r2;
; CHECK-NEXT:    add.s32 %r3, %r1, -1139;
; CHECK-NEXT:    shl.b64 %rd7, %rd2, %r3;
; CHECK-NEXT:    add.s32 %r4, %r1, -1075;
; CHECK-NEXT:    setp.gt.s32 %p3, %r4, 63;
; CHECK-NEXT:    selp.b64 %rd9, %rd7, %rd6, %p3;
; CHECK-NEXT:    shl.b64 %rd8, %rd2, %r4;
; CHECK-NEXT:  $L__BB2_4: // %fp-to-i-cleanup
; CHECK-NEXT:    st.param.v2.b64 [func_retval0], {%rd8, %rd9};
; CHECK-NEXT:    ret;
  %a = fptoui double %x to i128
  ret i128 %a
}

define half @cvt_f16_u128(i128 %x) {
; CHECK-LABEL: cvt_f16_u128(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<10>;
; CHECK-NEXT:    .reg .b16 %rs<2>;
; CHECK-NEXT:    .reg .b32 %r<19>;
; CHECK-NEXT:    .reg .b64 %rd<28>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %itofp-entry
; CHECK-NEXT:    ld.param.v2.b64 {%rd27, %rd1}, [cvt_f16_u128_param_0];
; CHECK-NEXT:    or.b64 %rd2, %rd27, %rd1;
; CHECK-NEXT:    setp.eq.b64 %p1, %rd2, 0;
; CHECK-NEXT:    mov.b16 %rs1, 0x0000;
; CHECK-NEXT:    @%p1 bra $L__BB3_10;
; CHECK-NEXT:  // %bb.1: // %itofp-if-end
; CHECK-NEXT:    setp.ne.b64 %p2, %rd1, 0;
; CHECK-NEXT:    clz.b64 %r3, %rd27;
; CHECK-NEXT:    cvt.u64.u32 %rd3, %r3;
; CHECK-NEXT:    add.s64 %rd4, %rd3, 64;
; CHECK-NEXT:    clz.b64 %r4, %rd1;
; CHECK-NEXT:    cvt.u64.u32 %rd5, %r4;
; CHECK-NEXT:    selp.b64 %rd6, %rd5, %rd4, %p2;
; CHECK-NEXT:    cvt.u32.u64 %r1, %rd6;
; CHECK-NEXT:    sub.s32 %r2, 128, %r1;
; CHECK-NEXT:    sub.s32 %r18, 127, %r1;
; CHECK-NEXT:    setp.lt.s32 %p3, %r2, 25;
; CHECK-NEXT:    @%p3 bra $L__BB3_8;
; CHECK-NEXT:  // %bb.2: // %itofp-if-then4
; CHECK-NEXT:    setp.eq.b32 %p4, %r2, 26;
; CHECK-NEXT:    @%p4 bra $L__BB3_6;
; CHECK-NEXT:  // %bb.3: // %itofp-if-then4
; CHECK-NEXT:    setp.ne.b32 %p5, %r2, 25;
; CHECK-NEXT:    @!%p5 bra $L__BB3_4;
; CHECK-NEXT:  // %bb.5: // %itofp-sw-default
; CHECK-NEXT:    sub.s32 %r6, 102, %r1;
; CHECK-NEXT:    shr.u64 %rd8, %rd27, %r6;
; CHECK-NEXT:    sub.s32 %r7, 64, %r6;
; CHECK-NEXT:    shl.b64 %rd9, %rd1, %r7;
; CHECK-NEXT:    or.b64 %rd10, %rd8, %rd9;
; CHECK-NEXT:    sub.s32 %r8, 38, %r1;
; CHECK-NEXT:    shr.u64 %rd11, %rd1, %r8;
; CHECK-NEXT:    setp.gt.s32 %p6, %r6, 63;
; CHECK-NEXT:    selp.b64 %rd12, %rd11, %rd10, %p6;
; CHECK-NEXT:    add.s32 %r9, %r1, 26;
; CHECK-NEXT:    shr.u64 %rd13, %rd27, %r8;
; CHECK-NEXT:    shl.b64 %rd14, %rd1, %r9;
; CHECK-NEXT:    or.b64 %rd15, %rd14, %rd13;
; CHECK-NEXT:    add.s32 %r10, %r1, -38;
; CHECK-NEXT:    shl.b64 %rd16, %rd27, %r10;
; CHECK-NEXT:    setp.gt.s32 %p7, %r9, 63;
; CHECK-NEXT:    selp.b64 %rd17, %rd16, %rd15, %p7;
; CHECK-NEXT:    shl.b64 %rd18, %rd27, %r9;
; CHECK-NEXT:    or.b64 %rd19, %rd18, %rd17;
; CHECK-NEXT:    setp.ne.b64 %p8, %rd19, 0;
; CHECK-NEXT:    selp.b64 %rd20, 1, 0, %p8;
; CHECK-NEXT:    or.b64 %rd27, %rd12, %rd20;
; CHECK-NEXT:  $L__BB3_6: // %itofp-sw-epilog
; CHECK-NEXT:    cvt.u32.u64 %r11, %rd27;
; CHECK-NEXT:    bfe.u32 %r12, %r11, 2, 1;
; CHECK-NEXT:    cvt.u64.u32 %rd21, %r12;
; CHECK-NEXT:    or.b64 %rd22, %rd27, %rd21;
; CHECK-NEXT:    add.s64 %rd23, %rd22, 1;
; CHECK-NEXT:    shr.u64 %rd24, %rd23, 2;
; CHECK-NEXT:    and.b64 %rd25, %rd23, 67108864;
; CHECK-NEXT:    setp.eq.b64 %p9, %rd25, 0;
; CHECK-NEXT:    cvt.u32.u64 %r17, %rd24;
; CHECK-NEXT:    @!%p9 bra $L__BB3_7;
; CHECK-NEXT:  $L__BB3_9: // %itofp-if-end26
; CHECK-NEXT:    shl.b32 %r13, %r18, 23;
; CHECK-NEXT:    and.b32 %r14, %r17, 8388607;
; CHECK-NEXT:    or.b32 %r15, %r13, %r14;
; CHECK-NEXT:    add.s32 %r16, %r15, 1065353216;
; CHECK-NEXT:    cvt.rn.f16.f32 %rs1, %r16;
; CHECK-NEXT:  $L__BB3_10: // %itofp-return
; CHECK-NEXT:    st.param.b16 [func_retval0], %rs1;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB3_8: // %itofp-if-else
; CHECK-NEXT:    add.s32 %r5, %r1, -104;
; CHECK-NEXT:    shl.b64 %rd7, %rd27, %r5;
; CHECK-NEXT:    cvt.u32.u64 %r17, %rd7;
; CHECK-NEXT:    bra.uni $L__BB3_9;
; CHECK-NEXT:  $L__BB3_7: // %itofp-if-then20
; CHECK-NEXT:    shr.u64 %rd26, %rd23, 3;
; CHECK-NEXT:    cvt.u32.u64 %r17, %rd26;
; CHECK-NEXT:    mov.b32 %r18, %r2;
; CHECK-NEXT:    bra.uni $L__BB3_9;
; CHECK-NEXT:  $L__BB3_4: // %itofp-sw-bb
; CHECK-NEXT:    shl.b64 %rd27, %rd27, 1;
; CHECK-NEXT:    bra.uni $L__BB3_6;
  %a = uitofp i128 %x to half
  ret half %a
}

define float @cvt_f32_u128(i128 %x) {
; CHECK-LABEL: cvt_f32_u128(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<10>;
; CHECK-NEXT:    .reg .b32 %r<19>;
; CHECK-NEXT:    .reg .b64 %rd<28>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %itofp-entry
; CHECK-NEXT:    ld.param.v2.b64 {%rd27, %rd1}, [cvt_f32_u128_param_0];
; CHECK-NEXT:    or.b64 %rd2, %rd27, %rd1;
; CHECK-NEXT:    setp.eq.b64 %p1, %rd2, 0;
; CHECK-NEXT:    mov.b32 %r18, 0f00000000;
; CHECK-NEXT:    @%p1 bra $L__BB4_10;
; CHECK-NEXT:  // %bb.1: // %itofp-if-end
; CHECK-NEXT:    setp.ne.b64 %p2, %rd1, 0;
; CHECK-NEXT:    clz.b64 %r3, %rd27;
; CHECK-NEXT:    cvt.u64.u32 %rd3, %r3;
; CHECK-NEXT:    add.s64 %rd4, %rd3, 64;
; CHECK-NEXT:    clz.b64 %r4, %rd1;
; CHECK-NEXT:    cvt.u64.u32 %rd5, %r4;
; CHECK-NEXT:    selp.b64 %rd6, %rd5, %rd4, %p2;
; CHECK-NEXT:    cvt.u32.u64 %r1, %rd6;
; CHECK-NEXT:    sub.s32 %r2, 128, %r1;
; CHECK-NEXT:    sub.s32 %r17, 127, %r1;
; CHECK-NEXT:    setp.lt.s32 %p3, %r2, 25;
; CHECK-NEXT:    @%p3 bra $L__BB4_8;
; CHECK-NEXT:  // %bb.2: // %itofp-if-then4
; CHECK-NEXT:    setp.eq.b32 %p4, %r2, 26;
; CHECK-NEXT:    @%p4 bra $L__BB4_6;
; CHECK-NEXT:  // %bb.3: // %itofp-if-then4
; CHECK-NEXT:    setp.ne.b32 %p5, %r2, 25;
; CHECK-NEXT:    @!%p5 bra $L__BB4_4;
; CHECK-NEXT:  // %bb.5: // %itofp-sw-default
; CHECK-NEXT:    sub.s32 %r6, 102, %r1;
; CHECK-NEXT:    shr.u64 %rd8, %rd27, %r6;
; CHECK-NEXT:    sub.s32 %r7, 64, %r6;
; CHECK-NEXT:    shl.b64 %rd9, %rd1, %r7;
; CHECK-NEXT:    or.b64 %rd10, %rd8, %rd9;
; CHECK-NEXT:    sub.s32 %r8, 38, %r1;
; CHECK-NEXT:    shr.u64 %rd11, %rd1, %r8;
; CHECK-NEXT:    setp.gt.s32 %p6, %r6, 63;
; CHECK-NEXT:    selp.b64 %rd12, %rd11, %rd10, %p6;
; CHECK-NEXT:    add.s32 %r9, %r1, 26;
; CHECK-NEXT:    shr.u64 %rd13, %rd27, %r8;
; CHECK-NEXT:    shl.b64 %rd14, %rd1, %r9;
; CHECK-NEXT:    or.b64 %rd15, %rd14, %rd13;
; CHECK-NEXT:    add.s32 %r10, %r1, -38;
; CHECK-NEXT:    shl.b64 %rd16, %rd27, %r10;
; CHECK-NEXT:    setp.gt.s32 %p7, %r9, 63;
; CHECK-NEXT:    selp.b64 %rd17, %rd16, %rd15, %p7;
; CHECK-NEXT:    shl.b64 %rd18, %rd27, %r9;
; CHECK-NEXT:    or.b64 %rd19, %rd18, %rd17;
; CHECK-NEXT:    setp.ne.b64 %p8, %rd19, 0;
; CHECK-NEXT:    selp.b64 %rd20, 1, 0, %p8;
; CHECK-NEXT:    or.b64 %rd27, %rd12, %rd20;
; CHECK-NEXT:  $L__BB4_6: // %itofp-sw-epilog
; CHECK-NEXT:    cvt.u32.u64 %r11, %rd27;
; CHECK-NEXT:    bfe.u32 %r12, %r11, 2, 1;
; CHECK-NEXT:    cvt.u64.u32 %rd21, %r12;
; CHECK-NEXT:    or.b64 %rd22, %rd27, %rd21;
; CHECK-NEXT:    add.s64 %rd23, %rd22, 1;
; CHECK-NEXT:    shr.u64 %rd24, %rd23, 2;
; CHECK-NEXT:    and.b64 %rd25, %rd23, 67108864;
; CHECK-NEXT:    setp.eq.b64 %p9, %rd25, 0;
; CHECK-NEXT:    cvt.u32.u64 %r16, %rd24;
; CHECK-NEXT:    @!%p9 bra $L__BB4_7;
; CHECK-NEXT:  $L__BB4_9: // %itofp-if-end26
; CHECK-NEXT:    shl.b32 %r13, %r17, 23;
; CHECK-NEXT:    and.b32 %r14, %r16, 8388607;
; CHECK-NEXT:    or.b32 %r15, %r13, %r14;
; CHECK-NEXT:    add.s32 %r18, %r15, 1065353216;
; CHECK-NEXT:  $L__BB4_10: // %itofp-return
; CHECK-NEXT:    st.param.b32 [func_retval0], %r18;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB4_8: // %itofp-if-else
; CHECK-NEXT:    add.s32 %r5, %r1, -104;
; CHECK-NEXT:    shl.b64 %rd7, %rd27, %r5;
; CHECK-NEXT:    cvt.u32.u64 %r16, %rd7;
; CHECK-NEXT:    bra.uni $L__BB4_9;
; CHECK-NEXT:  $L__BB4_7: // %itofp-if-then20
; CHECK-NEXT:    shr.u64 %rd26, %rd23, 3;
; CHECK-NEXT:    cvt.u32.u64 %r16, %rd26;
; CHECK-NEXT:    mov.b32 %r17, %r2;
; CHECK-NEXT:    bra.uni $L__BB4_9;
; CHECK-NEXT:  $L__BB4_4: // %itofp-sw-bb
; CHECK-NEXT:    shl.b64 %rd27, %rd27, 1;
; CHECK-NEXT:    bra.uni $L__BB4_6;
  %a = uitofp i128 %x to float
  ret float %a
}

define double @cvt_f64_u128(i128 %x) {
; CHECK-LABEL: cvt_f64_u128(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<10>;
; CHECK-NEXT:    .reg .b32 %r<35>;
; CHECK-NEXT:    .reg .b64 %rd<34>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %itofp-entry
; CHECK-NEXT:    ld.param.v2.b64 {%rd30, %rd31}, [cvt_f64_u128_param_0];
; CHECK-NEXT:    or.b64 %rd2, %rd30, %rd31;
; CHECK-NEXT:    setp.eq.b64 %p1, %rd2, 0;
; CHECK-NEXT:    mov.b64 %rd33, 0d0000000000000000;
; CHECK-NEXT:    @%p1 bra $L__BB5_10;
; CHECK-NEXT:  // %bb.1: // %itofp-if-end
; CHECK-NEXT:    setp.ne.b64 %p2, %rd31, 0;
; CHECK-NEXT:    clz.b64 %r3, %rd30;
; CHECK-NEXT:    cvt.u64.u32 %rd3, %r3;
; CHECK-NEXT:    add.s64 %rd4, %rd3, 64;
; CHECK-NEXT:    clz.b64 %r4, %rd31;
; CHECK-NEXT:    cvt.u64.u32 %rd5, %r4;
; CHECK-NEXT:    selp.b64 %rd6, %rd5, %rd4, %p2;
; CHECK-NEXT:    cvt.u32.u64 %r1, %rd6;
; CHECK-NEXT:    sub.s32 %r2, 128, %r1;
; CHECK-NEXT:    sub.s32 %r34, 127, %r1;
; CHECK-NEXT:    setp.lt.s32 %p3, %r2, 54;
; CHECK-NEXT:    @%p3 bra $L__BB5_8;
; CHECK-NEXT:  // %bb.2: // %itofp-if-then4
; CHECK-NEXT:    setp.eq.b32 %p4, %r2, 55;
; CHECK-NEXT:    @%p4 bra $L__BB5_6;
; CHECK-NEXT:  // %bb.3: // %itofp-if-then4
; CHECK-NEXT:    setp.ne.b32 %p5, %r2, 54;
; CHECK-NEXT:    @!%p5 bra $L__BB5_4;
; CHECK-NEXT:  // %bb.5: // %itofp-sw-default
; CHECK-NEXT:    sub.s32 %r12, 73, %r1;
; CHECK-NEXT:    shr.u64 %rd7, %rd30, %r12;
; CHECK-NEXT:    sub.s32 %r13, 64, %r12;
; CHECK-NEXT:    shl.b64 %rd8, %rd31, %r13;
; CHECK-NEXT:    or.b64 %rd9, %rd7, %rd8;
; CHECK-NEXT:    sub.s32 %r14, 9, %r1;
; CHECK-NEXT:    shr.u64 %rd10, %rd31, %r14;
; CHECK-NEXT:    setp.gt.s32 %p6, %r12, 63;
; CHECK-NEXT:    selp.b64 %rd11, %rd10, %rd9, %p6;
; CHECK-NEXT:    shr.u64 %rd1, %rd31, %r12;
; CHECK-NEXT:    add.s32 %r15, %r1, 55;
; CHECK-NEXT:    shr.u64 %rd12, %rd30, %r14;
; CHECK-NEXT:    shl.b64 %rd13, %rd31, %r15;
; CHECK-NEXT:    or.b64 %rd14, %rd13, %rd12;
; CHECK-NEXT:    add.s32 %r16, %r1, -9;
; CHECK-NEXT:    shl.b64 %rd15, %rd30, %r16;
; CHECK-NEXT:    setp.gt.s32 %p7, %r15, 63;
; CHECK-NEXT:    selp.b64 %rd16, %rd15, %rd14, %p7;
; CHECK-NEXT:    shl.b64 %rd17, %rd30, %r15;
; CHECK-NEXT:    or.b64 %rd18, %rd17, %rd16;
; CHECK-NEXT:    setp.ne.b64 %p8, %rd18, 0;
; CHECK-NEXT:    selp.b64 %rd19, 1, 0, %p8;
; CHECK-NEXT:    or.b64 %rd30, %rd11, %rd19;
; CHECK-NEXT:    mov.b64 %rd31, %rd1;
; CHECK-NEXT:  $L__BB5_6: // %itofp-sw-epilog
; CHECK-NEXT:    cvt.u32.u64 %r17, %rd30;
; CHECK-NEXT:    bfe.u32 %r18, %r17, 2, 1;
; CHECK-NEXT:    cvt.u64.u32 %rd20, %r18;
; CHECK-NEXT:    or.b64 %rd21, %rd30, %rd20;
; CHECK-NEXT:    add.cc.s64 %rd22, %rd21, 1;
; CHECK-NEXT:    addc.cc.s64 %rd23, %rd31, 0;
; CHECK-NEXT:    mov.b64 {%r19, %r20}, %rd22;
; CHECK-NEXT:    shf.l.wrap.b32 %r21, %r19, %r20, 30;
; CHECK-NEXT:    mov.b64 {%r22, %r23}, %rd23;
; CHECK-NEXT:    shf.l.wrap.b32 %r24, %r20, %r22, 30;
; CHECK-NEXT:    mov.b64 %rd32, {%r21, %r24};
; CHECK-NEXT:    and.b64 %rd24, %rd22, 36028797018963968;
; CHECK-NEXT:    setp.eq.b64 %p9, %rd24, 0;
; CHECK-NEXT:    shf.l.wrap.b32 %r25, %r22, %r23, 30;
; CHECK-NEXT:    mov.b64 %rd25, {%r24, %r25};
; CHECK-NEXT:    cvt.u32.u64 %r33, %rd25;
; CHECK-NEXT:    @!%p9 bra $L__BB5_7;
; CHECK-NEXT:  $L__BB5_9: // %itofp-if-end26
; CHECK-NEXT:    shl.b32 %r29, %r34, 20;
; CHECK-NEXT:    and.b32 %r30, %r33, 1048575;
; CHECK-NEXT:    or.b32 %r31, %r29, %r30;
; CHECK-NEXT:    add.s32 %r32, %r31, 1072693248;
; CHECK-NEXT:    cvt.u64.u32 %rd27, %r32;
; CHECK-NEXT:    shl.b64 %rd28, %rd27, 32;
; CHECK-NEXT:    and.b64 %rd29, %rd32, 4294967295;
; CHECK-NEXT:    or.b64 %rd33, %rd28, %rd29;
; CHECK-NEXT:  $L__BB5_10: // %itofp-return
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd33;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB5_8: // %itofp-if-else
; CHECK-NEXT:    add.s32 %r5, %r1, -75;
; CHECK-NEXT:    shl.b64 %rd32, %rd30, %r5;
; CHECK-NEXT:    { .reg .b32 tmp; mov.b64 {tmp, %r33}, %rd32; }
; CHECK-NEXT:    bra.uni $L__BB5_9;
; CHECK-NEXT:  $L__BB5_7: // %itofp-if-then20
; CHECK-NEXT:    shf.l.wrap.b32 %r26, %r19, %r20, 29;
; CHECK-NEXT:    shf.l.wrap.b32 %r27, %r20, %r22, 29;
; CHECK-NEXT:    mov.b64 %rd32, {%r26, %r27};
; CHECK-NEXT:    shf.l.wrap.b32 %r28, %r22, %r23, 29;
; CHECK-NEXT:    mov.b64 %rd26, {%r27, %r28};
; CHECK-NEXT:    cvt.u32.u64 %r33, %rd26;
; CHECK-NEXT:    mov.b32 %r34, %r2;
; CHECK-NEXT:    bra.uni $L__BB5_9;
; CHECK-NEXT:  $L__BB5_4: // %itofp-sw-bb
; CHECK-NEXT:    mov.b64 {%r6, %r7}, %rd30;
; CHECK-NEXT:    mov.b64 {%r8, %r9}, %rd31;
; CHECK-NEXT:    shf.l.wrap.b32 %r10, %r7, %r8, 1;
; CHECK-NEXT:    shf.l.wrap.b32 %r11, %r8, %r9, 1;
; CHECK-NEXT:    mov.b64 %rd31, {%r10, %r11};
; CHECK-NEXT:    shl.b64 %rd30, %rd30, 1;
; CHECK-NEXT:    bra.uni $L__BB5_6;
  %a = uitofp i128 %x to double
  ret double %a
}

define i128 @cvt_s128_f16(half %x) {
; CHECK-LABEL: cvt_s128_f16(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b16 %rs1, [cvt_s128_f16_param_0];
; CHECK-NEXT:    cvt.rzi.s32.f16 %r1, %rs1;
; CHECK-NEXT:    cvt.s64.s32 %rd1, %r1;
; CHECK-NEXT:    shr.s64 %rd2, %rd1, 63;
; CHECK-NEXT:    st.param.v2.b64 [func_retval0], {%rd1, %rd2};
; CHECK-NEXT:    ret;
  %a = fptosi half %x to i128
  ret i128 %a
}

define i128 @cvt_s128_f32(float %x) {
; CHECK-LABEL: cvt_s128_f32(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<4>;
; CHECK-NEXT:    .reg .b32 %r<11>;
; CHECK-NEXT:    .reg .b64 %rd<14>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %fp-to-i-entry
; CHECK-NEXT:    ld.param.b32 %r3, [cvt_s128_f32_param_0];
; CHECK-NEXT:    bfe.u32 %r1, %r3, 23, 8;
; CHECK-NEXT:    setp.lt.u32 %p1, %r1, 127;
; CHECK-NEXT:    mov.b64 %rd12, 0;
; CHECK-NEXT:    mov.b64 %rd13, %rd12;
; CHECK-NEXT:    @%p1 bra $L__BB7_4;
; CHECK-NEXT:  // %bb.1: // %fp-to-i-if-check.exp.size
; CHECK-NEXT:    shr.s32 %r4, %r3, 31;
; CHECK-NEXT:    cvt.s64.s32 %rd2, %r4;
; CHECK-NEXT:    or.b64 %rd1, %rd2, 1;
; CHECK-NEXT:    and.b32 %r5, %r3, 8388607;
; CHECK-NEXT:    or.b32 %r2, %r5, 8388608;
; CHECK-NEXT:    setp.gt.u32 %p2, %r1, 149;
; CHECK-NEXT:    @%p2 bra $L__BB7_3;
; CHECK-NEXT:  // %bb.2: // %fp-to-i-if-exp.small
; CHECK-NEXT:    sub.s32 %r9, 150, %r1;
; CHECK-NEXT:    shr.u32 %r10, %r2, %r9;
; CHECK-NEXT:    cvt.u64.u32 %rd10, %r10;
; CHECK-NEXT:    mul.hi.u64 %rd11, %rd10, %rd1;
; CHECK-NEXT:    mad.lo.s64 %rd13, %rd10, %rd2, %rd11;
; CHECK-NEXT:    mul.lo.s64 %rd12, %rd10, %rd1;
; CHECK-NEXT:    bra.uni $L__BB7_4;
; CHECK-NEXT:  $L__BB7_3: // %fp-to-i-if-exp.large
; CHECK-NEXT:    add.s32 %r6, %r1, -150;
; CHECK-NEXT:    cvt.u64.u32 %rd3, %r2;
; CHECK-NEXT:    sub.s32 %r7, 214, %r1;
; CHECK-NEXT:    shr.u64 %rd4, %rd3, %r7;
; CHECK-NEXT:    add.s32 %r8, %r1, -214;
; CHECK-NEXT:    shl.b64 %rd5, %rd3, %r8;
; CHECK-NEXT:    setp.gt.s32 %p3, %r6, 63;
; CHECK-NEXT:    selp.b64 %rd6, %rd5, %rd4, %p3;
; CHECK-NEXT:    shl.b64 %rd7, %rd3, %r6;
; CHECK-NEXT:    mul.hi.u64 %rd8, %rd7, %rd1;
; CHECK-NEXT:    mad.lo.s64 %rd9, %rd7, %rd2, %rd8;
; CHECK-NEXT:    mad.lo.s64 %rd13, %rd6, %rd1, %rd9;
; CHECK-NEXT:    mul.lo.s64 %rd12, %rd7, %rd1;
; CHECK-NEXT:  $L__BB7_4: // %fp-to-i-cleanup
; CHECK-NEXT:    st.param.v2.b64 [func_retval0], {%rd12, %rd13};
; CHECK-NEXT:    ret;
  %a = fptosi float %x to i128
  ret i128 %a
}

define i128 @cvt_s128_f64(double %x) {
; CHECK-LABEL: cvt_s128_f64(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<4>;
; CHECK-NEXT:    .reg .b32 %r<7>;
; CHECK-NEXT:    .reg .b64 %rd<18>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %fp-to-i-entry
; CHECK-NEXT:    ld.param.b64 %rd5, [cvt_s128_f64_param_0];
; CHECK-NEXT:    shr.u64 %rd6, %rd5, 52;
; CHECK-NEXT:    and.b64 %rd3, %rd6, 2047;
; CHECK-NEXT:    setp.lt.u64 %p1, %rd3, 1023;
; CHECK-NEXT:    mov.b64 %rd16, 0;
; CHECK-NEXT:    mov.b64 %rd17, %rd16;
; CHECK-NEXT:    @%p1 bra $L__BB8_4;
; CHECK-NEXT:  // %bb.1: // %fp-to-i-if-check.exp.size
; CHECK-NEXT:    shr.s64 %rd2, %rd5, 63;
; CHECK-NEXT:    or.b64 %rd1, %rd2, 1;
; CHECK-NEXT:    and.b64 %rd7, %rd5, 4503599627370495;
; CHECK-NEXT:    or.b64 %rd4, %rd7, 4503599627370496;
; CHECK-NEXT:    setp.gt.u64 %p2, %rd3, 1074;
; CHECK-NEXT:    @%p2 bra $L__BB8_3;
; CHECK-NEXT:  // %bb.2: // %fp-to-i-if-exp.small
; CHECK-NEXT:    cvt.u32.u64 %r5, %rd3;
; CHECK-NEXT:    sub.s32 %r6, 1075, %r5;
; CHECK-NEXT:    shr.u64 %rd14, %rd4, %r6;
; CHECK-NEXT:    mul.hi.u64 %rd15, %rd14, %rd1;
; CHECK-NEXT:    mad.lo.s64 %rd17, %rd14, %rd2, %rd15;
; CHECK-NEXT:    mul.lo.s64 %rd16, %rd14, %rd1;
; CHECK-NEXT:    bra.uni $L__BB8_4;
; CHECK-NEXT:  $L__BB8_3: // %fp-to-i-if-exp.large
; CHECK-NEXT:    cvt.u32.u64 %r1, %rd3;
; CHECK-NEXT:    sub.s32 %r2, 1139, %r1;
; CHECK-NEXT:    shr.u64 %rd8, %rd4, %r2;
; CHECK-NEXT:    add.s32 %r3, %r1, -1139;
; CHECK-NEXT:    shl.b64 %rd9, %rd4, %r3;
; CHECK-NEXT:    add.s32 %r4, %r1, -1075;
; CHECK-NEXT:    setp.gt.s32 %p3, %r4, 63;
; CHECK-NEXT:    selp.b64 %rd10, %rd9, %rd8, %p3;
; CHECK-NEXT:    shl.b64 %rd11, %rd4, %r4;
; CHECK-NEXT:    mul.hi.u64 %rd12, %rd11, %rd1;
; CHECK-NEXT:    mad.lo.s64 %rd13, %rd11, %rd2, %rd12;
; CHECK-NEXT:    mad.lo.s64 %rd17, %rd10, %rd1, %rd13;
; CHECK-NEXT:    mul.lo.s64 %rd16, %rd11, %rd1;
; CHECK-NEXT:  $L__BB8_4: // %fp-to-i-cleanup
; CHECK-NEXT:    st.param.v2.b64 [func_retval0], {%rd16, %rd17};
; CHECK-NEXT:    ret;
  %a = fptosi double %x to i128
  ret i128 %a
}

define half @cvt_f16_s128(i128 %x) {
; CHECK-LABEL: cvt_f16_s128(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<11>;
; CHECK-NEXT:    .reg .b16 %rs<2>;
; CHECK-NEXT:    .reg .b32 %r<22>;
; CHECK-NEXT:    .reg .b64 %rd<33>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %itofp-entry
; CHECK-NEXT:    ld.param.v2.b64 {%rd1, %rd2}, [cvt_f16_s128_param_0];
; CHECK-NEXT:    or.b64 %rd5, %rd1, %rd2;
; CHECK-NEXT:    setp.eq.b64 %p1, %rd5, 0;
; CHECK-NEXT:    mov.b16 %rs1, 0x0000;
; CHECK-NEXT:    @%p1 bra $L__BB9_10;
; CHECK-NEXT:  // %bb.1: // %itofp-if-end
; CHECK-NEXT:    shr.s64 %rd3, %rd2, 63;
; CHECK-NEXT:    sub.cc.s64 %rd6, 0, %rd1;
; CHECK-NEXT:    subc.cc.s64 %rd7, 0, %rd2;
; CHECK-NEXT:    setp.lt.s64 %p2, %rd2, 0;
; CHECK-NEXT:    selp.b64 %rd32, %rd6, %rd1, %p2;
; CHECK-NEXT:    selp.b64 %rd4, %rd7, %rd2, %p2;
; CHECK-NEXT:    setp.ne.b64 %p3, %rd4, 0;
; CHECK-NEXT:    clz.b64 %r3, %rd4;
; CHECK-NEXT:    cvt.u64.u32 %rd8, %r3;
; CHECK-NEXT:    clz.b64 %r4, %rd32;
; CHECK-NEXT:    cvt.u64.u32 %rd9, %r4;
; CHECK-NEXT:    add.s64 %rd10, %rd9, 64;
; CHECK-NEXT:    selp.b64 %rd11, %rd8, %rd10, %p3;
; CHECK-NEXT:    cvt.u32.u64 %r1, %rd11;
; CHECK-NEXT:    sub.s32 %r2, 128, %r1;
; CHECK-NEXT:    sub.s32 %r21, 127, %r1;
; CHECK-NEXT:    setp.lt.s32 %p4, %r2, 25;
; CHECK-NEXT:    @%p4 bra $L__BB9_8;
; CHECK-NEXT:  // %bb.2: // %itofp-if-then4
; CHECK-NEXT:    setp.eq.b32 %p5, %r2, 26;
; CHECK-NEXT:    @%p5 bra $L__BB9_6;
; CHECK-NEXT:  // %bb.3: // %itofp-if-then4
; CHECK-NEXT:    setp.ne.b32 %p6, %r2, 25;
; CHECK-NEXT:    @!%p6 bra $L__BB9_4;
; CHECK-NEXT:  // %bb.5: // %itofp-sw-default
; CHECK-NEXT:    sub.s32 %r6, 102, %r1;
; CHECK-NEXT:    shr.u64 %rd13, %rd32, %r6;
; CHECK-NEXT:    sub.s32 %r7, 64, %r6;
; CHECK-NEXT:    shl.b64 %rd14, %rd4, %r7;
; CHECK-NEXT:    or.b64 %rd15, %rd13, %rd14;
; CHECK-NEXT:    sub.s32 %r8, 38, %r1;
; CHECK-NEXT:    shr.u64 %rd16, %rd4, %r8;
; CHECK-NEXT:    setp.gt.s32 %p7, %r6, 63;
; CHECK-NEXT:    selp.b64 %rd17, %rd16, %rd15, %p7;
; CHECK-NEXT:    add.s32 %r9, %r1, 26;
; CHECK-NEXT:    shr.u64 %rd18, %rd32, %r8;
; CHECK-NEXT:    shl.b64 %rd19, %rd4, %r9;
; CHECK-NEXT:    or.b64 %rd20, %rd19, %rd18;
; CHECK-NEXT:    add.s32 %r10, %r1, -38;
; CHECK-NEXT:    shl.b64 %rd21, %rd32, %r10;
; CHECK-NEXT:    setp.gt.s32 %p8, %r9, 63;
; CHECK-NEXT:    selp.b64 %rd22, %rd21, %rd20, %p8;
; CHECK-NEXT:    shl.b64 %rd23, %rd32, %r9;
; CHECK-NEXT:    or.b64 %rd24, %rd23, %rd22;
; CHECK-NEXT:    setp.ne.b64 %p9, %rd24, 0;
; CHECK-NEXT:    selp.b64 %rd25, 1, 0, %p9;
; CHECK-NEXT:    or.b64 %rd32, %rd17, %rd25;
; CHECK-NEXT:  $L__BB9_6: // %itofp-sw-epilog
; CHECK-NEXT:    cvt.u32.u64 %r11, %rd32;
; CHECK-NEXT:    bfe.u32 %r12, %r11, 2, 1;
; CHECK-NEXT:    cvt.u64.u32 %rd26, %r12;
; CHECK-NEXT:    or.b64 %rd27, %rd32, %rd26;
; CHECK-NEXT:    add.s64 %rd28, %rd27, 1;
; CHECK-NEXT:    shr.u64 %rd29, %rd28, 2;
; CHECK-NEXT:    and.b64 %rd30, %rd28, 67108864;
; CHECK-NEXT:    setp.eq.b64 %p10, %rd30, 0;
; CHECK-NEXT:    cvt.u32.u64 %r20, %rd29;
; CHECK-NEXT:    @!%p10 bra $L__BB9_7;
; CHECK-NEXT:  $L__BB9_9: // %itofp-if-end26
; CHECK-NEXT:    cvt.u32.u64 %r13, %rd3;
; CHECK-NEXT:    and.b32 %r14, %r13, -2147483648;
; CHECK-NEXT:    shl.b32 %r15, %r21, 23;
; CHECK-NEXT:    add.s32 %r16, %r15, 1065353216;
; CHECK-NEXT:    and.b32 %r17, %r20, 8388607;
; CHECK-NEXT:    or.b32 %r18, %r17, %r14;
; CHECK-NEXT:    or.b32 %r19, %r18, %r16;
; CHECK-NEXT:    cvt.rn.f16.f32 %rs1, %r19;
; CHECK-NEXT:  $L__BB9_10: // %itofp-return
; CHECK-NEXT:    st.param.b16 [func_retval0], %rs1;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB9_8: // %itofp-if-else
; CHECK-NEXT:    add.s32 %r5, %r1, -104;
; CHECK-NEXT:    shl.b64 %rd12, %rd32, %r5;
; CHECK-NEXT:    cvt.u32.u64 %r20, %rd12;
; CHECK-NEXT:    bra.uni $L__BB9_9;
; CHECK-NEXT:  $L__BB9_7: // %itofp-if-then20
; CHECK-NEXT:    shr.u64 %rd31, %rd28, 3;
; CHECK-NEXT:    cvt.u32.u64 %r20, %rd31;
; CHECK-NEXT:    mov.b32 %r21, %r2;
; CHECK-NEXT:    bra.uni $L__BB9_9;
; CHECK-NEXT:  $L__BB9_4: // %itofp-sw-bb
; CHECK-NEXT:    shl.b64 %rd32, %rd32, 1;
; CHECK-NEXT:    bra.uni $L__BB9_6;
  %a = sitofp i128 %x to half
  ret half %a
}

define float @cvt_f32_s128(i128 %x) {
; CHECK-LABEL: cvt_f32_s128(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<11>;
; CHECK-NEXT:    .reg .b32 %r<22>;
; CHECK-NEXT:    .reg .b64 %rd<33>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %itofp-entry
; CHECK-NEXT:    ld.param.v2.b64 {%rd1, %rd2}, [cvt_f32_s128_param_0];
; CHECK-NEXT:    or.b64 %rd5, %rd1, %rd2;
; CHECK-NEXT:    setp.eq.b64 %p1, %rd5, 0;
; CHECK-NEXT:    mov.b32 %r21, 0f00000000;
; CHECK-NEXT:    @%p1 bra $L__BB10_10;
; CHECK-NEXT:  // %bb.1: // %itofp-if-end
; CHECK-NEXT:    shr.s64 %rd3, %rd2, 63;
; CHECK-NEXT:    sub.cc.s64 %rd6, 0, %rd1;
; CHECK-NEXT:    subc.cc.s64 %rd7, 0, %rd2;
; CHECK-NEXT:    setp.lt.s64 %p2, %rd2, 0;
; CHECK-NEXT:    selp.b64 %rd32, %rd6, %rd1, %p2;
; CHECK-NEXT:    selp.b64 %rd4, %rd7, %rd2, %p2;
; CHECK-NEXT:    setp.ne.b64 %p3, %rd4, 0;
; CHECK-NEXT:    clz.b64 %r3, %rd4;
; CHECK-NEXT:    cvt.u64.u32 %rd8, %r3;
; CHECK-NEXT:    clz.b64 %r4, %rd32;
; CHECK-NEXT:    cvt.u64.u32 %rd9, %r4;
; CHECK-NEXT:    add.s64 %rd10, %rd9, 64;
; CHECK-NEXT:    selp.b64 %rd11, %rd8, %rd10, %p3;
; CHECK-NEXT:    cvt.u32.u64 %r1, %rd11;
; CHECK-NEXT:    sub.s32 %r2, 128, %r1;
; CHECK-NEXT:    sub.s32 %r20, 127, %r1;
; CHECK-NEXT:    setp.lt.s32 %p4, %r2, 25;
; CHECK-NEXT:    @%p4 bra $L__BB10_8;
; CHECK-NEXT:  // %bb.2: // %itofp-if-then4
; CHECK-NEXT:    setp.eq.b32 %p5, %r2, 26;
; CHECK-NEXT:    @%p5 bra $L__BB10_6;
; CHECK-NEXT:  // %bb.3: // %itofp-if-then4
; CHECK-NEXT:    setp.ne.b32 %p6, %r2, 25;
; CHECK-NEXT:    @!%p6 bra $L__BB10_4;
; CHECK-NEXT:  // %bb.5: // %itofp-sw-default
; CHECK-NEXT:    sub.s32 %r6, 102, %r1;
; CHECK-NEXT:    shr.u64 %rd13, %rd32, %r6;
; CHECK-NEXT:    sub.s32 %r7, 64, %r6;
; CHECK-NEXT:    shl.b64 %rd14, %rd4, %r7;
; CHECK-NEXT:    or.b64 %rd15, %rd13, %rd14;
; CHECK-NEXT:    sub.s32 %r8, 38, %r1;
; CHECK-NEXT:    shr.u64 %rd16, %rd4, %r8;
; CHECK-NEXT:    setp.gt.s32 %p7, %r6, 63;
; CHECK-NEXT:    selp.b64 %rd17, %rd16, %rd15, %p7;
; CHECK-NEXT:    add.s32 %r9, %r1, 26;
; CHECK-NEXT:    shr.u64 %rd18, %rd32, %r8;
; CHECK-NEXT:    shl.b64 %rd19, %rd4, %r9;
; CHECK-NEXT:    or.b64 %rd20, %rd19, %rd18;
; CHECK-NEXT:    add.s32 %r10, %r1, -38;
; CHECK-NEXT:    shl.b64 %rd21, %rd32, %r10;
; CHECK-NEXT:    setp.gt.s32 %p8, %r9, 63;
; CHECK-NEXT:    selp.b64 %rd22, %rd21, %rd20, %p8;
; CHECK-NEXT:    shl.b64 %rd23, %rd32, %r9;
; CHECK-NEXT:    or.b64 %rd24, %rd23, %rd22;
; CHECK-NEXT:    setp.ne.b64 %p9, %rd24, 0;
; CHECK-NEXT:    selp.b64 %rd25, 1, 0, %p9;
; CHECK-NEXT:    or.b64 %rd32, %rd17, %rd25;
; CHECK-NEXT:  $L__BB10_6: // %itofp-sw-epilog
; CHECK-NEXT:    cvt.u32.u64 %r11, %rd32;
; CHECK-NEXT:    bfe.u32 %r12, %r11, 2, 1;
; CHECK-NEXT:    cvt.u64.u32 %rd26, %r12;
; CHECK-NEXT:    or.b64 %rd27, %rd32, %rd26;
; CHECK-NEXT:    add.s64 %rd28, %rd27, 1;
; CHECK-NEXT:    shr.u64 %rd29, %rd28, 2;
; CHECK-NEXT:    and.b64 %rd30, %rd28, 67108864;
; CHECK-NEXT:    setp.eq.b64 %p10, %rd30, 0;
; CHECK-NEXT:    cvt.u32.u64 %r19, %rd29;
; CHECK-NEXT:    @!%p10 bra $L__BB10_7;
; CHECK-NEXT:  $L__BB10_9: // %itofp-if-end26
; CHECK-NEXT:    cvt.u32.u64 %r13, %rd3;
; CHECK-NEXT:    and.b32 %r14, %r13, -2147483648;
; CHECK-NEXT:    shl.b32 %r15, %r20, 23;
; CHECK-NEXT:    add.s32 %r16, %r15, 1065353216;
; CHECK-NEXT:    and.b32 %r17, %r19, 8388607;
; CHECK-NEXT:    or.b32 %r18, %r17, %r14;
; CHECK-NEXT:    or.b32 %r21, %r18, %r16;
; CHECK-NEXT:  $L__BB10_10: // %itofp-return
; CHECK-NEXT:    st.param.b32 [func_retval0], %r21;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB10_8: // %itofp-if-else
; CHECK-NEXT:    add.s32 %r5, %r1, -104;
; CHECK-NEXT:    shl.b64 %rd12, %rd32, %r5;
; CHECK-NEXT:    cvt.u32.u64 %r19, %rd12;
; CHECK-NEXT:    bra.uni $L__BB10_9;
; CHECK-NEXT:  $L__BB10_7: // %itofp-if-then20
; CHECK-NEXT:    shr.u64 %rd31, %rd28, 3;
; CHECK-NEXT:    cvt.u32.u64 %r19, %rd31;
; CHECK-NEXT:    mov.b32 %r20, %r2;
; CHECK-NEXT:    bra.uni $L__BB10_9;
; CHECK-NEXT:  $L__BB10_4: // %itofp-sw-bb
; CHECK-NEXT:    shl.b64 %rd32, %rd32, 1;
; CHECK-NEXT:    bra.uni $L__BB10_6;
  %a = sitofp i128 %x to float
  ret float %a
}

define double @cvt_f64_s128(i128 %x) {
; CHECK-LABEL: cvt_f64_s128(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<11>;
; CHECK-NEXT:    .reg .b32 %r<36>;
; CHECK-NEXT:    .reg .b64 %rd<37>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %itofp-entry
; CHECK-NEXT:    ld.param.v2.b64 {%rd1, %rd2}, [cvt_f64_s128_param_0];
; CHECK-NEXT:    or.b64 %rd5, %rd1, %rd2;
; CHECK-NEXT:    setp.eq.b64 %p1, %rd5, 0;
; CHECK-NEXT:    mov.b64 %rd36, 0d0000000000000000;
; CHECK-NEXT:    @%p1 bra $L__BB11_10;
; CHECK-NEXT:  // %bb.1: // %itofp-if-end
; CHECK-NEXT:    shr.s64 %rd3, %rd2, 63;
; CHECK-NEXT:    sub.cc.s64 %rd6, 0, %rd1;
; CHECK-NEXT:    subc.cc.s64 %rd7, 0, %rd2;
; CHECK-NEXT:    setp.lt.s64 %p2, %rd2, 0;
; CHECK-NEXT:    selp.b64 %rd33, %rd6, %rd1, %p2;
; CHECK-NEXT:    selp.b64 %rd34, %rd7, %rd2, %p2;
; CHECK-NEXT:    setp.ne.b64 %p3, %rd34, 0;
; CHECK-NEXT:    clz.b64 %r3, %rd34;
; CHECK-NEXT:    cvt.u64.u32 %rd8, %r3;
; CHECK-NEXT:    clz.b64 %r4, %rd33;
; CHECK-NEXT:    cvt.u64.u32 %rd9, %r4;
; CHECK-NEXT:    add.s64 %rd10, %rd9, 64;
; CHECK-NEXT:    selp.b64 %rd11, %rd8, %rd10, %p3;
; CHECK-NEXT:    cvt.u32.u64 %r1, %rd11;
; CHECK-NEXT:    sub.s32 %r2, 128, %r1;
; CHECK-NEXT:    sub.s32 %r35, 127, %r1;
; CHECK-NEXT:    setp.lt.s32 %p4, %r2, 54;
; CHECK-NEXT:    @%p4 bra $L__BB11_8;
; CHECK-NEXT:  // %bb.2: // %itofp-if-then4
; CHECK-NEXT:    setp.eq.b32 %p5, %r2, 55;
; CHECK-NEXT:    @%p5 bra $L__BB11_6;
; CHECK-NEXT:  // %bb.3: // %itofp-if-then4
; CHECK-NEXT:    setp.ne.b32 %p6, %r2, 54;
; CHECK-NEXT:    @!%p6 bra $L__BB11_4;
; CHECK-NEXT:  // %bb.5: // %itofp-sw-default
; CHECK-NEXT:    sub.s32 %r12, 73, %r1;
; CHECK-NEXT:    shr.u64 %rd12, %rd33, %r12;
; CHECK-NEXT:    sub.s32 %r13, 64, %r12;
; CHECK-NEXT:    shl.b64 %rd13, %rd34, %r13;
; CHECK-NEXT:    or.b64 %rd14, %rd12, %rd13;
; CHECK-NEXT:    sub.s32 %r14, 9, %r1;
; CHECK-NEXT:    shr.u64 %rd15, %rd34, %r14;
; CHECK-NEXT:    setp.gt.s32 %p7, %r12, 63;
; CHECK-NEXT:    selp.b64 %rd16, %rd15, %rd14, %p7;
; CHECK-NEXT:    shr.u64 %rd4, %rd34, %r12;
; CHECK-NEXT:    add.s32 %r15, %r1, 55;
; CHECK-NEXT:    shr.u64 %rd17, %rd33, %r14;
; CHECK-NEXT:    shl.b64 %rd18, %rd34, %r15;
; CHECK-NEXT:    or.b64 %rd19, %rd18, %rd17;
; CHECK-NEXT:    add.s32 %r16, %r1, -9;
; CHECK-NEXT:    shl.b64 %rd20, %rd33, %r16;
; CHECK-NEXT:    setp.gt.s32 %p8, %r15, 63;
; CHECK-NEXT:    selp.b64 %rd21, %rd20, %rd19, %p8;
; CHECK-NEXT:    shl.b64 %rd22, %rd33, %r15;
; CHECK-NEXT:    or.b64 %rd23, %rd22, %rd21;
; CHECK-NEXT:    setp.ne.b64 %p9, %rd23, 0;
; CHECK-NEXT:    selp.b64 %rd24, 1, 0, %p9;
; CHECK-NEXT:    or.b64 %rd33, %rd16, %rd24;
; CHECK-NEXT:    mov.b64 %rd34, %rd4;
; CHECK-NEXT:  $L__BB11_6: // %itofp-sw-epilog
; CHECK-NEXT:    cvt.u32.u64 %r17, %rd33;
; CHECK-NEXT:    bfe.u32 %r18, %r17, 2, 1;
; CHECK-NEXT:    cvt.u64.u32 %rd25, %r18;
; CHECK-NEXT:    or.b64 %rd26, %rd33, %rd25;
; CHECK-NEXT:    add.cc.s64 %rd27, %rd26, 1;
; CHECK-NEXT:    addc.cc.s64 %rd28, %rd34, 0;
; CHECK-NEXT:    mov.b64 {%r19, %r20}, %rd28;
; CHECK-NEXT:    mov.b64 {%r21, %r22}, %rd27;
; CHECK-NEXT:    shf.l.wrap.b32 %r23, %r22, %r19, 30;
; CHECK-NEXT:    shf.l.wrap.b32 %r24, %r21, %r22, 30;
; CHECK-NEXT:    mov.b64 %rd35, {%r24, %r23};
; CHECK-NEXT:    and.b64 %rd29, %rd27, 36028797018963968;
; CHECK-NEXT:    setp.eq.b64 %p10, %rd29, 0;
; CHECK-NEXT:    { .reg .b32 tmp; mov.b64 {tmp, %r34}, %rd35; }
; CHECK-NEXT:    @!%p10 bra $L__BB11_7;
; CHECK-NEXT:  $L__BB11_9: // %itofp-if-end26
; CHECK-NEXT:    cvt.u32.u64 %r27, %rd3;
; CHECK-NEXT:    and.b32 %r28, %r27, -2147483648;
; CHECK-NEXT:    shl.b32 %r29, %r35, 20;
; CHECK-NEXT:    add.s32 %r30, %r29, 1072693248;
; CHECK-NEXT:    and.b32 %r31, %r34, 1048575;
; CHECK-NEXT:    or.b32 %r32, %r31, %r28;
; CHECK-NEXT:    or.b32 %r33, %r32, %r30;
; CHECK-NEXT:    cvt.u64.u32 %rd30, %r33;
; CHECK-NEXT:    shl.b64 %rd31, %rd30, 32;
; CHECK-NEXT:    and.b64 %rd32, %rd35, 4294967295;
; CHECK-NEXT:    or.b64 %rd36, %rd31, %rd32;
; CHECK-NEXT:  $L__BB11_10: // %itofp-return
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd36;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB11_8: // %itofp-if-else
; CHECK-NEXT:    add.s32 %r5, %r1, -75;
; CHECK-NEXT:    shl.b64 %rd35, %rd33, %r5;
; CHECK-NEXT:    { .reg .b32 tmp; mov.b64 {tmp, %r34}, %rd35; }
; CHECK-NEXT:    bra.uni $L__BB11_9;
; CHECK-NEXT:  $L__BB11_7: // %itofp-if-then20
; CHECK-NEXT:    shf.l.wrap.b32 %r25, %r22, %r19, 29;
; CHECK-NEXT:    shf.l.wrap.b32 %r26, %r21, %r22, 29;
; CHECK-NEXT:    mov.b64 %rd35, {%r26, %r25};
; CHECK-NEXT:    { .reg .b32 tmp; mov.b64 {tmp, %r34}, %rd35; }
; CHECK-NEXT:    mov.b32 %r35, %r2;
; CHECK-NEXT:    bra.uni $L__BB11_9;
; CHECK-NEXT:  $L__BB11_4: // %itofp-sw-bb
; CHECK-NEXT:    mov.b64 {%r6, %r7}, %rd33;
; CHECK-NEXT:    mov.b64 {%r8, %r9}, %rd34;
; CHECK-NEXT:    shf.l.wrap.b32 %r10, %r7, %r8, 1;
; CHECK-NEXT:    shf.l.wrap.b32 %r11, %r8, %r9, 1;
; CHECK-NEXT:    mov.b64 %rd34, {%r10, %r11};
; CHECK-NEXT:    shl.b64 %rd33, %rd33, 1;
; CHECK-NEXT:    bra.uni $L__BB11_6;
  %a = sitofp i128 %x to double
  ret double %a
}