[llvm] [NVPTX] Expand fp/int conversions involving integers wider than 64 bits (PR #201679)
Lukas Stephan via llvm-commits
llvm-commits at lists.llvm.org
Fri Jun 12 03:41:30 PDT 2026
https://github.com/kulst updated https://github.com/llvm/llvm-project/pull/201679
>From 6525040632f7a8b5eef6ce0dee29f53edebfd1f4 Mon Sep 17 00:00:00 2001
From: kulst <kulst at mailbox.org>
Date: Thu, 4 Jun 2026 19:05:27 +0200
Subject: [PATCH] [NVPTX] Expand fp/int conversions involving integers wider
than 64 bits
NVPTX does not support direct PTX conversion instructions between floating
point types and integer types wider than 64 bits.
Previously, such conversions could reach instruction selection and fail with
an "unsupported library call operation" error.
Set the maximum supported fp/int conversion width to 64 bits so larger
conversions are expanded before instruction selection.
Add regression coverage for 128-bit integer/floating-point conversions.
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 1 +
llvm/test/CodeGen/NVPTX/convert-fp-i128.ll | 849 ++++++++++++++++++++
2 files changed, 850 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/convert-fp-i128.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 17d9f857312d6..e9beb22808ca8 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1120,6 +1120,7 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setMinCmpXchgSizeInBits(STI.getMinCmpXchgSizeInBits());
setMaxAtomicSizeInBitsSupported(STI.hasAtomSwap128() ? 128 : 64);
setMaxDivRemBitWidthSupported(64);
+ setMaxLargeFPConvertBitWidthSupported(64);
// Custom lowering for tcgen05.ld vector operands
setOperationAction(ISD::INTRINSIC_W_CHAIN,
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp-i128.ll b/llvm/test/CodeGen/NVPTX/convert-fp-i128.ll
new file mode 100644
index 0000000000000..f9eff6a8d6cd1
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp-i128.ll
@@ -0,0 +1,849 @@
+; 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
+}
More information about the llvm-commits
mailing list