From 64f76dea9c95fd59e383e7550de35786781d9662 Mon Sep 17 00:00:00 2001 From: Orlando Cazalet-Hyams Date: Fri, 15 Mar 2024 13:40:44 +0000 Subject: [PATCH 001/782] [NFC] Fix comment in test from #83251 --- llvm/test/Bitcode/dbg-record-roundtrip.ll | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/llvm/test/Bitcode/dbg-record-roundtrip.ll b/llvm/test/Bitcode/dbg-record-roundtrip.ll index e7774a9d80d9..242c3819c860 100644 --- a/llvm/test/Bitcode/dbg-record-roundtrip.ll +++ b/llvm/test/Bitcode/dbg-record-roundtrip.ll @@ -1,7 +1,6 @@ ;; Roundtrip tests. ; RUN: llvm-as --write-experimental-debuginfo-iterators-to-bitcode=true %s -o - | llvm-dis | FileCheck %s ;; Check that verify-uselistorder passes regardless of input format. -;; NOTE: This test fails intermittently ; RUN: llvm-as %s --write-experimental-debuginfo-iterators-to-bitcode=true -o - | verify-uselistorder %s ; RUN: verify-uselistorder %s @@ -15,6 +14,9 @@ ;; Check that llvm-link doesn't explode if we give it different formats to ;; link. +;; NOTE: This test fails intermittently on linux if the llvm-as output is piped +;; into llvm-link in the RUN lines below, unless the verify-uselistorder RUN +;; lines above are removed. Write to a temporary file to avoid that weirdness. ; RUN: llvm-as %s --experimental-debuginfo-iterators=true --write-experimental-debuginfo-iterators-to-bitcode=true -o %t ; RUN: llvm-link %t %s --experimental-debuginfo-iterators=false -o /dev/null ; RUN: llvm-as %s --experimental-debuginfo-iterators=false -o %t -- GitLab From 0ae76a74985b3639ffb99d1dbb857068939547aa Mon Sep 17 00:00:00 2001 From: Orlando Cazalet-Hyams Date: Fri, 15 Mar 2024 14:01:14 +0000 Subject: [PATCH 002/782] [NFC] Fix incorrect RUN line in test from #83251 Note: This wasn't the cause of the strange behaviour mentioned in the NOTE comment in the test. --- llvm/test/Bitcode/dbg-record-roundtrip.ll | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/test/Bitcode/dbg-record-roundtrip.ll b/llvm/test/Bitcode/dbg-record-roundtrip.ll index 242c3819c860..d092e3822d26 100644 --- a/llvm/test/Bitcode/dbg-record-roundtrip.ll +++ b/llvm/test/Bitcode/dbg-record-roundtrip.ll @@ -1,7 +1,7 @@ ;; Roundtrip tests. ; RUN: llvm-as --write-experimental-debuginfo-iterators-to-bitcode=true %s -o - | llvm-dis | FileCheck %s ;; Check that verify-uselistorder passes regardless of input format. -; RUN: llvm-as %s --write-experimental-debuginfo-iterators-to-bitcode=true -o - | verify-uselistorder %s +; RUN: llvm-as %s --write-experimental-debuginfo-iterators-to-bitcode=true -o - | verify-uselistorder ; RUN: verify-uselistorder %s ;; Confirm we're producing RemoveDI records from various tools. -- GitLab From 0b9f19a9880eb786871194af116f223d2ad30c52 Mon Sep 17 00:00:00 2001 From: antoine moynault Date: Fri, 15 Mar 2024 15:07:40 +0100 Subject: [PATCH 003/782] Revert "[compiler-rt] Avoid generating coredumps when piped to a tool" (#85390) This reverts commit 27e5312a8bc8935f9c5620ff061c647d9fbcec85. This commit broke some bots: - clang-aarch64-sve-vla https://lab.llvm.org/buildbot/#/builders/197/builds/13609 - clang-aarch64-sve-vls https://lab.llvm.org/buildbot/#/builders/184/builds/10988 - clang-aarch64-lld-2stage https://lab.llvm.org/buildbot/#/builders/185/builds/6312 https://github.com/llvm/llvm-project/pull/83701 --- .../sanitizer_posix_libcdep.cpp | 19 +------------------ .../sanitizer_common/TestCases/corelimit.cpp | 7 +------ 2 files changed, 2 insertions(+), 24 deletions(-) diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_posix_libcdep.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_posix_libcdep.cpp index 3605d0d666e3..ef1fc3549743 100644 --- a/compiler-rt/lib/sanitizer_common/sanitizer_posix_libcdep.cpp +++ b/compiler-rt/lib/sanitizer_common/sanitizer_posix_libcdep.cpp @@ -104,24 +104,7 @@ static void setlim(int res, rlim_t lim) { void DisableCoreDumperIfNecessary() { if (common_flags()->disable_coredump) { - rlimit rlim; - CHECK_EQ(0, getrlimit(RLIMIT_CORE, &rlim)); - // On Linux, if the kernel.core_pattern sysctl starts with a '|' (i.e. it - // is being piped to a coredump handler such as systemd-coredumpd), the - // kernel ignores RLIMIT_CORE (since we aren't creating a file in the file - // system) except for the magic value of 1, which disables coredumps when - // piping. 1 byte is too small for any kind of valid core dump, so it - // also disables coredumps if kernel.core_pattern creates files directly. - // While most piped coredump handlers do respect the crashing processes' - // RLIMIT_CORE, this is notable not the case for Debian's systemd-coredump - // due to a local patch that changes sysctl.d/50-coredump.conf to ignore - // the specified limit and instead use RLIM_INFINITY. - // - // The alternative to using RLIMIT_CORE=1 would be to use prctl() with the - // PR_SET_DUMPABLE flag, however that also prevents ptrace(), so makes it - // impossible to attach a debugger. - rlim.rlim_cur = Min(SANITIZER_LINUX ? 1 : 0, rlim.rlim_max); - CHECK_EQ(0, setrlimit(RLIMIT_CORE, &rlim)); + setlim(RLIMIT_CORE, 0); } } diff --git a/compiler-rt/test/sanitizer_common/TestCases/corelimit.cpp b/compiler-rt/test/sanitizer_common/TestCases/corelimit.cpp index fed2e1d89cbf..2378a4cfdced 100644 --- a/compiler-rt/test/sanitizer_common/TestCases/corelimit.cpp +++ b/compiler-rt/test/sanitizer_common/TestCases/corelimit.cpp @@ -10,12 +10,7 @@ int main() { getrlimit(RLIMIT_CORE, &lim_core); void *p; if (sizeof(p) == 8) { -#ifdef __linux__ - // See comments in DisableCoreDumperIfNecessary(). - assert(lim_core.rlim_cur == 1); -#else - assert(lim_core.rlim_cur == 0); -#endif + assert(0 == lim_core.rlim_cur); } return 0; } -- GitLab From f4676b6be6ee0d908c92d64936d17bd6fa3fbda8 Mon Sep 17 00:00:00 2001 From: Phoebe Wang Date: Fri, 15 Mar 2024 22:09:56 +0800 Subject: [PATCH 004/782] [X86] Add Support for X86 TLSDESC Relocations (#83136) --- clang/lib/Driver/ToolChains/CommonArgs.cpp | 3 +- clang/test/Driver/tls-dialect.c | 2 +- llvm/lib/Target/X86/X86ISelDAGToDAG.cpp | 16 +- llvm/lib/Target/X86/X86ISelLowering.cpp | 38 +++- llvm/lib/Target/X86/X86ISelLowering.h | 4 + llvm/lib/Target/X86/X86InstrCompiler.td | 10 ++ llvm/lib/Target/X86/X86InstrFragments.td | 3 + llvm/lib/Target/X86/X86MCInstLower.cpp | 33 +++- llvm/test/CodeGen/X86/tls-desc.ll | 199 +++++++++++++++++++++ 9 files changed, 289 insertions(+), 19 deletions(-) create mode 100644 llvm/test/CodeGen/X86/tls-desc.ll diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp index 100e71245394..83015b0cb81a 100644 --- a/clang/lib/Driver/ToolChains/CommonArgs.cpp +++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp @@ -740,7 +740,8 @@ bool tools::isTLSDESCEnabled(const ToolChain &TC, SupportedArgument = V == "desc" || V == "trad"; EnableTLSDESC = V == "desc"; } else if (Triple.isX86()) { - SupportedArgument = V == "gnu"; + SupportedArgument = V == "gnu" || V == "gnu2"; + EnableTLSDESC = V == "gnu2"; } else { Unsupported = true; } diff --git a/clang/test/Driver/tls-dialect.c b/clang/test/Driver/tls-dialect.c index f73915b28ec2..a808dd81531c 100644 --- a/clang/test/Driver/tls-dialect.c +++ b/clang/test/Driver/tls-dialect.c @@ -2,6 +2,7 @@ // RUN: %clang -### --target=riscv64-linux -mtls-dialect=trad %s 2>&1 | FileCheck --check-prefix=NODESC %s // RUN: %clang -### --target=riscv64-linux %s 2>&1 | FileCheck --check-prefix=NODESC %s // RUN: %clang -### --target=x86_64-linux -mtls-dialect=gnu %s 2>&1 | FileCheck --check-prefix=NODESC %s +// RUN: %clang -### --target=x86_64-linux -mtls-dialect=gnu2 %s 2>&1 | FileCheck --check-prefix=DESC %s /// Android supports TLSDESC by default on RISC-V /// TLSDESC is not on by default in Linux, even on RISC-V, and is covered above @@ -18,7 +19,6 @@ /// Unsupported argument // RUN: not %clang -### --target=riscv64-linux -mtls-dialect=gnu2 %s 2>&1 | FileCheck --check-prefix=UNSUPPORTED-ARG %s -// RUN: not %clang -### --target=x86_64-linux -mtls-dialect=gnu2 %s 2>&1 | FileCheck --check-prefix=UNSUPPORTED-ARG %s // DESC: "-cc1" {{.*}}"-enable-tlsdesc" // NODESC-NOT: "-enable-tlsdesc" diff --git a/llvm/lib/Target/X86/X86ISelDAGToDAG.cpp b/llvm/lib/Target/X86/X86ISelDAGToDAG.cpp index 76c6c1645239..4e4241efd63d 100644 --- a/llvm/lib/Target/X86/X86ISelDAGToDAG.cpp +++ b/llvm/lib/Target/X86/X86ISelDAGToDAG.cpp @@ -3090,13 +3090,19 @@ bool X86DAGToDAGISel::selectLEAAddr(SDValue N, bool X86DAGToDAGISel::selectTLSADDRAddr(SDValue N, SDValue &Base, SDValue &Scale, SDValue &Index, SDValue &Disp, SDValue &Segment) { - assert(N.getOpcode() == ISD::TargetGlobalTLSAddress); - auto *GA = cast(N); + assert(N.getOpcode() == ISD::TargetGlobalTLSAddress || + N.getOpcode() == ISD::TargetExternalSymbol); X86ISelAddressMode AM; - AM.GV = GA->getGlobal(); - AM.Disp += GA->getOffset(); - AM.SymbolFlags = GA->getTargetFlags(); + if (auto *GA = dyn_cast(N)) { + AM.GV = GA->getGlobal(); + AM.Disp += GA->getOffset(); + AM.SymbolFlags = GA->getTargetFlags(); + } else { + auto *SA = cast(N); + AM.ES = SA->getSymbol(); + AM.SymbolFlags = SA->getTargetFlags(); + } if (Subtarget->is32Bit()) { AM.Scale = 1; diff --git a/llvm/lib/Target/X86/X86ISelLowering.cpp b/llvm/lib/Target/X86/X86ISelLowering.cpp index 2b5e3c0379a1..dbfcb3752ae8 100644 --- a/llvm/lib/Target/X86/X86ISelLowering.cpp +++ b/llvm/lib/Target/X86/X86ISelLowering.cpp @@ -18592,13 +18592,22 @@ GetTLSADDR(SelectionDAG &DAG, SDValue Chain, GlobalAddressSDNode *GA, MachineFrameInfo &MFI = DAG.getMachineFunction().getFrameInfo(); SDVTList NodeTys = DAG.getVTList(MVT::Other, MVT::Glue); SDLoc dl(GA); - SDValue TGA = DAG.getTargetGlobalAddress(GA->getGlobal(), dl, - GA->getValueType(0), - GA->getOffset(), - OperandFlags); + SDValue TGA; + bool UseTLSDESC = DAG.getTarget().useTLSDESC(); + if (LocalDynamic && UseTLSDESC) { + TGA = DAG.getTargetExternalSymbol("_TLS_MODULE_BASE_", PtrVT, OperandFlags); + auto UI = TGA->use_begin(); + // Reuse existing GetTLSADDR node if we can find it. + if (UI != TGA->use_end()) + return SDValue(*UI->use_begin()->use_begin(), 0); + } else { + TGA = DAG.getTargetGlobalAddress(GA->getGlobal(), dl, GA->getValueType(0), + GA->getOffset(), OperandFlags); + } - X86ISD::NodeType CallType = LocalDynamic ? X86ISD::TLSBASEADDR - : X86ISD::TLSADDR; + X86ISD::NodeType CallType = UseTLSDESC ? X86ISD::TLSDESC + : LocalDynamic ? X86ISD::TLSBASEADDR + : X86ISD::TLSADDR; if (InGlue) { SDValue Ops[] = { Chain, TGA, *InGlue }; @@ -18613,7 +18622,19 @@ GetTLSADDR(SelectionDAG &DAG, SDValue Chain, GlobalAddressSDNode *GA, MFI.setHasCalls(true); SDValue Glue = Chain.getValue(1); - return DAG.getCopyFromReg(Chain, dl, ReturnReg, PtrVT, Glue); + SDValue Ret = DAG.getCopyFromReg(Chain, dl, ReturnReg, PtrVT, Glue); + + if (!UseTLSDESC) + return Ret; + + const X86Subtarget &Subtarget = DAG.getSubtarget(); + unsigned Seg = Subtarget.is64Bit() ? X86AS::FS : X86AS::GS; + + Value *Ptr = Constant::getNullValue(PointerType::get(*DAG.getContext(), Seg)); + SDValue Offset = + DAG.getLoad(PtrVT, dl, DAG.getEntryNode(), DAG.getIntPtrConstant(0, dl), + MachinePointerInfo(Ptr)); + return DAG.getNode(ISD::ADD, dl, PtrVT, Ret, Offset); } // Lower ISD::GlobalTLSAddress using the "general dynamic" model, 32 bit @@ -33426,6 +33447,7 @@ const char *X86TargetLowering::getTargetNodeName(unsigned Opcode) const { NODE_NAME_CASE(TLSADDR) NODE_NAME_CASE(TLSBASEADDR) NODE_NAME_CASE(TLSCALL) + NODE_NAME_CASE(TLSDESC) NODE_NAME_CASE(EH_SJLJ_SETJMP) NODE_NAME_CASE(EH_SJLJ_LONGJMP) NODE_NAME_CASE(EH_SJLJ_SETUP_DISPATCH) @@ -36206,6 +36228,8 @@ X86TargetLowering::EmitInstrWithCustomInserter(MachineInstr &MI, case X86::TLS_base_addr32: case X86::TLS_base_addr64: case X86::TLS_base_addrX32: + case X86::TLS_desc32: + case X86::TLS_desc64: return EmitLoweredTLSAddr(MI, BB); case X86::INDIRECT_THUNK_CALL32: case X86::INDIRECT_THUNK_CALL64: diff --git a/llvm/lib/Target/X86/X86ISelLowering.h b/llvm/lib/Target/X86/X86ISelLowering.h index fe1943b57608..0a1e8ca44273 100644 --- a/llvm/lib/Target/X86/X86ISelLowering.h +++ b/llvm/lib/Target/X86/X86ISelLowering.h @@ -295,6 +295,10 @@ namespace llvm { // thunk at the address from an earlier relocation. TLSCALL, + // Thread Local Storage. A descriptor containing pointer to + // code and to argument to get the TLS offset for the symbol. + TLSDESC, + // Exception Handling helpers. EH_RETURN, diff --git a/llvm/lib/Target/X86/X86InstrCompiler.td b/llvm/lib/Target/X86/X86InstrCompiler.td index f393f86e64aa..ce3b6af4cab4 100644 --- a/llvm/lib/Target/X86/X86InstrCompiler.td +++ b/llvm/lib/Target/X86/X86InstrCompiler.td @@ -507,6 +507,16 @@ def TLS_base_addrX32 : I<0, Pseudo, (outs), (ins i32mem:$sym), Requires<[In64BitMode, NotLP64]>; } +// TLSDESC only clobbers EAX and EFLAGS. ESP is marked as a use to prevent +// stack-pointer assignments that appear immediately before calls from +// potentially appearing dead. +let Defs = [EAX, EFLAGS], usesCustomInserter = 1, Uses = [RSP, SSP] in { + def TLS_desc32 : I<0, Pseudo, (outs), (ins i32mem:$sym), + "# TLS_desc32", [(X86tlsdesc tls32addr:$sym)]>; + def TLS_desc64 : I<0, Pseudo, (outs), (ins i64mem:$sym), + "# TLS_desc64", [(X86tlsdesc tls64addr:$sym)]>; +} + // Darwin TLS Support // For i386, the address of the thunk is passed on the stack, on return the // address of the variable is in %eax. %ecx is trashed during the function diff --git a/llvm/lib/Target/X86/X86InstrFragments.td b/llvm/lib/Target/X86/X86InstrFragments.td index adf527d72f5b..f14c7200af96 100644 --- a/llvm/lib/Target/X86/X86InstrFragments.td +++ b/llvm/lib/Target/X86/X86InstrFragments.td @@ -223,6 +223,9 @@ def X86tlsaddr : SDNode<"X86ISD::TLSADDR", SDT_X86TLSADDR, def X86tlsbaseaddr : SDNode<"X86ISD::TLSBASEADDR", SDT_X86TLSBASEADDR, [SDNPHasChain, SDNPOptInGlue, SDNPOutGlue]>; +def X86tlsdesc : SDNode<"X86ISD::TLSDESC", SDT_X86TLSADDR, + [SDNPHasChain, SDNPOptInGlue, SDNPOutGlue]>; + def X86ehret : SDNode<"X86ISD::EH_RETURN", SDT_X86EHRET, [SDNPHasChain]>; diff --git a/llvm/lib/Target/X86/X86MCInstLower.cpp b/llvm/lib/Target/X86/X86MCInstLower.cpp index 64d4d411e7b4..e2330ff34c17 100644 --- a/llvm/lib/Target/X86/X86MCInstLower.cpp +++ b/llvm/lib/Target/X86/X86MCInstLower.cpp @@ -519,10 +519,8 @@ void X86MCInstLower::Lower(const MachineInstr *MI, MCInst &OutMI) const { void X86AsmPrinter::LowerTlsAddr(X86MCInstLower &MCInstLowering, const MachineInstr &MI) { NoAutoPaddingScope NoPadScope(*OutStreamer); - bool Is64Bits = MI.getOpcode() != X86::TLS_addr32 && - MI.getOpcode() != X86::TLS_base_addr32; - bool Is64BitsLP64 = MI.getOpcode() == X86::TLS_addr64 || - MI.getOpcode() == X86::TLS_base_addr64; + bool Is64Bits = getSubtarget().is64Bit(); + bool Is64BitsLP64 = getSubtarget().isTarget64BitLP64(); MCContext &Ctx = OutStreamer->getContext(); MCSymbolRefExpr::VariantKind SRVK; @@ -539,6 +537,10 @@ void X86AsmPrinter::LowerTlsAddr(X86MCInstLower &MCInstLowering, case X86::TLS_base_addrX32: SRVK = MCSymbolRefExpr::VK_TLSLD; break; + case X86::TLS_desc32: + case X86::TLS_desc64: + SRVK = MCSymbolRefExpr::VK_TLSDESC; + break; default: llvm_unreachable("unexpected opcode"); } @@ -554,7 +556,26 @@ void X86AsmPrinter::LowerTlsAddr(X86MCInstLower &MCInstLowering, bool UseGot = MMI->getModule()->getRtLibUseGOT() && Ctx.getTargetOptions()->X86RelaxRelocations; - if (Is64Bits) { + if (SRVK == MCSymbolRefExpr::VK_TLSDESC) { + const MCSymbolRefExpr *Expr = MCSymbolRefExpr::create( + MCInstLowering.GetSymbolFromOperand(MI.getOperand(3)), + MCSymbolRefExpr::VK_TLSCALL, Ctx); + EmitAndCountInstruction( + MCInstBuilder(Is64BitsLP64 ? X86::LEA64r : X86::LEA32r) + .addReg(Is64BitsLP64 ? X86::RAX : X86::EAX) + .addReg(Is64Bits ? X86::RIP : X86::EBX) + .addImm(1) + .addReg(0) + .addExpr(Sym) + .addReg(0)); + EmitAndCountInstruction( + MCInstBuilder(Is64Bits ? X86::CALL64m : X86::CALL32m) + .addReg(Is64BitsLP64 ? X86::RAX : X86::EAX) + .addImm(1) + .addReg(0) + .addExpr(Expr) + .addReg(0)); + } else if (Is64Bits) { bool NeedsPadding = SRVK == MCSymbolRefExpr::VK_TLSGD; if (NeedsPadding && Is64BitsLP64) EmitAndCountInstruction(MCInstBuilder(X86::DATA16_PREFIX)); @@ -2164,6 +2185,8 @@ void X86AsmPrinter::emitInstruction(const MachineInstr *MI) { case X86::TLS_base_addr32: case X86::TLS_base_addr64: case X86::TLS_base_addrX32: + case X86::TLS_desc32: + case X86::TLS_desc64: return LowerTlsAddr(MCInstLowering, *MI); case X86::MOVPC32r: { diff --git a/llvm/test/CodeGen/X86/tls-desc.ll b/llvm/test/CodeGen/X86/tls-desc.ll new file mode 100644 index 000000000000..c73986e69e79 --- /dev/null +++ b/llvm/test/CodeGen/X86/tls-desc.ll @@ -0,0 +1,199 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4 +; RUN: llc < %s -mtriple=i686 --relocation-model=pic -enable-tlsdesc | FileCheck %s --check-prefix=X86 +; RUN: llc < %s -mtriple=x86_64-pc-linux-gnux32 --relocation-model=pic -enable-tlsdesc | FileCheck %s --check-prefix=X32 +; RUN: llc < %s -mtriple=x86_64 --relocation-model=pic -enable-tlsdesc | FileCheck %s --check-prefix=X64 + +@x = thread_local global i32 0, align 4 +@y = internal thread_local global i32 1, align 4 +@z = external hidden thread_local global i32, align 4 + +define ptr @f1() nounwind { +; X86-LABEL: f1: +; X86: # %bb.0: +; X86-NEXT: pushl %ebp +; X86-NEXT: pushl %ebx +; X86-NEXT: pushl %edi +; X86-NEXT: pushl %esi +; X86-NEXT: pushl %eax +; X86-NEXT: calll .L0$pb +; X86-NEXT: .L0$pb: +; X86-NEXT: popl %ebx +; X86-NEXT: .Ltmp0: +; X86-NEXT: addl $_GLOBAL_OFFSET_TABLE_+(.Ltmp0-.L0$pb), %ebx +; X86-NEXT: #APP +; X86-NEXT: #NO_APP +; X86-NEXT: movl %eax, (%esp) # 4-byte Spill +; X86-NEXT: leal x@tlsdesc(%ebx), %eax +; X86-NEXT: calll *x@tlscall(%eax) +; X86-NEXT: addl %gs:0, %eax +; X86-NEXT: movl (%esp), %ebx # 4-byte Reload +; X86-NEXT: #APP +; X86-NEXT: #NO_APP +; X86-NEXT: addl $4, %esp +; X86-NEXT: popl %esi +; X86-NEXT: popl %edi +; X86-NEXT: popl %ebx +; X86-NEXT: popl %ebp +; X86-NEXT: retl +; +; X32-LABEL: f1: +; X32: # %bb.0: +; X32-NEXT: pushq %rax +; X32-NEXT: #APP +; X32-NEXT: #NO_APP +; X32-NEXT: leal x@tlsdesc(%rip), %eax +; X32-NEXT: callq *x@tlscall(%eax) +; X32-NEXT: # kill: def $eax killed $eax def $rax +; X32-NEXT: addl %fs:0, %eax +; X32-NEXT: #APP +; X32-NEXT: #NO_APP +; X32-NEXT: popq %rcx +; X32-NEXT: retq +; +; X64-LABEL: f1: +; X64: # %bb.0: +; X64-NEXT: pushq %rax +; X64-NEXT: #APP +; X64-NEXT: #NO_APP +; X64-NEXT: leaq x@tlsdesc(%rip), %rax +; X64-NEXT: callq *x@tlscall(%rax) +; X64-NEXT: addq %fs:0, %rax +; X64-NEXT: #APP +; X64-NEXT: #NO_APP +; X64-NEXT: popq %rcx +; X64-NEXT: retq + %a = call { i32, i32, i32, i32, i32, i32 } asm sideeffect "", "=r,=r,=r,=r,=r,=r,~{dirflag},~{fpsr},~{flags}"() + %b = call ptr @llvm.threadlocal.address.p0(ptr @x) + %a.0 = extractvalue { i32, i32, i32, i32, i32, i32 } %a, 0 + %a.1 = extractvalue { i32, i32, i32, i32, i32, i32 } %a, 1 + %a.2 = extractvalue { i32, i32, i32, i32, i32, i32 } %a, 2 + %a.3 = extractvalue { i32, i32, i32, i32, i32, i32 } %a, 3 + %a.4 = extractvalue { i32, i32, i32, i32, i32, i32 } %a, 4 + %a.5 = extractvalue { i32, i32, i32, i32, i32, i32 } %a, 5 + call void asm sideeffect "", "r,r,r,r,r,r,~{dirflag},~{fpsr},~{flags}"(i32 %a.0, i32 %a.1, i32 %a.2, i32 %a.3, i32 %a.4, i32 %a.5) + ret ptr %b +} + +define i32 @f2() nounwind { +; X86-LABEL: f2: +; X86: # %bb.0: +; X86-NEXT: pushl %ebx +; X86-NEXT: calll .L1$pb +; X86-NEXT: .L1$pb: +; X86-NEXT: popl %ebx +; X86-NEXT: .Ltmp1: +; X86-NEXT: addl $_GLOBAL_OFFSET_TABLE_+(.Ltmp1-.L1$pb), %ebx +; X86-NEXT: movl %gs:0, %ecx +; X86-NEXT: leal x@tlsdesc(%ebx), %eax +; X86-NEXT: calll *x@tlscall(%eax) +; X86-NEXT: movl (%eax,%ecx), %eax +; X86-NEXT: popl %ebx +; X86-NEXT: retl +; +; X32-LABEL: f2: +; X32: # %bb.0: +; X32-NEXT: pushq %rax +; X32-NEXT: movl %fs:0, %ecx +; X32-NEXT: leal x@tlsdesc(%rip), %eax +; X32-NEXT: callq *x@tlscall(%eax) +; X32-NEXT: movl (%eax,%ecx), %eax +; X32-NEXT: popq %rcx +; X32-NEXT: retq +; +; X64-LABEL: f2: +; X64: # %bb.0: +; X64-NEXT: pushq %rax +; X64-NEXT: movq %fs:0, %rcx +; X64-NEXT: leaq x@tlsdesc(%rip), %rax +; X64-NEXT: callq *x@tlscall(%rax) +; X64-NEXT: movl (%rax,%rcx), %eax +; X64-NEXT: popq %rcx +; X64-NEXT: retq + %1 = tail call ptr @llvm.threadlocal.address.p0(ptr @x) + %2 = load i32, ptr %1 + ret i32 %2 +} + +define ptr @f3() nounwind { +; X86-LABEL: f3: +; X86: # %bb.0: +; X86-NEXT: pushl %ebx +; X86-NEXT: calll .L2$pb +; X86-NEXT: .L2$pb: +; X86-NEXT: popl %ebx +; X86-NEXT: .Ltmp2: +; X86-NEXT: addl $_GLOBAL_OFFSET_TABLE_+(.Ltmp2-.L2$pb), %ebx +; X86-NEXT: leal x@tlsdesc(%ebx), %eax +; X86-NEXT: calll *x@tlscall(%eax) +; X86-NEXT: addl %gs:0, %eax +; X86-NEXT: popl %ebx +; X86-NEXT: retl +; +; X32-LABEL: f3: +; X32: # %bb.0: +; X32-NEXT: pushq %rax +; X32-NEXT: leal x@tlsdesc(%rip), %eax +; X32-NEXT: callq *x@tlscall(%eax) +; X32-NEXT: # kill: def $eax killed $eax def $rax +; X32-NEXT: addl %fs:0, %eax +; X32-NEXT: popq %rcx +; X32-NEXT: retq +; +; X64-LABEL: f3: +; X64: # %bb.0: +; X64-NEXT: pushq %rax +; X64-NEXT: leaq x@tlsdesc(%rip), %rax +; X64-NEXT: callq *x@tlscall(%rax) +; X64-NEXT: addq %fs:0, %rax +; X64-NEXT: popq %rcx +; X64-NEXT: retq + %1 = tail call ptr @llvm.threadlocal.address.p0(ptr @x) + ret ptr %1 +} + +define i32 @f4() nounwind { +; X86-LABEL: f4: +; X86: # %bb.0: +; X86-NEXT: pushl %ebx +; X86-NEXT: calll .L3$pb +; X86-NEXT: .L3$pb: +; X86-NEXT: popl %ebx +; X86-NEXT: .Ltmp3: +; X86-NEXT: addl $_GLOBAL_OFFSET_TABLE_+(.Ltmp3-.L3$pb), %ebx +; X86-NEXT: movl %gs:0, %edx +; X86-NEXT: leal _TLS_MODULE_BASE_@tlsdesc(%ebx), %eax +; X86-NEXT: calll *_TLS_MODULE_BASE_@tlscall(%eax) +; X86-NEXT: movl y@DTPOFF(%eax,%edx), %ecx +; X86-NEXT: addl z@DTPOFF(%eax,%edx), %ecx +; X86-NEXT: movl %ecx, %eax +; X86-NEXT: popl %ebx +; X86-NEXT: retl +; +; X32-LABEL: f4: +; X32: # %bb.0: +; X32-NEXT: pushq %rax +; X32-NEXT: movl %fs:0, %edx +; X32-NEXT: leal _TLS_MODULE_BASE_@tlsdesc(%rip), %eax +; X32-NEXT: callq *_TLS_MODULE_BASE_@tlscall(%eax) +; X32-NEXT: movl y@DTPOFF(%eax,%edx), %ecx +; X32-NEXT: addl z@DTPOFF(%eax,%edx), %ecx +; X32-NEXT: movl %ecx, %eax +; X32-NEXT: popq %rcx +; X32-NEXT: retq +; +; X64-LABEL: f4: +; X64: # %bb.0: +; X64-NEXT: pushq %rax +; X64-NEXT: movq %fs:0, %rdx +; X64-NEXT: leaq _TLS_MODULE_BASE_@tlsdesc(%rip), %rax +; X64-NEXT: callq *_TLS_MODULE_BASE_@tlscall(%rax) +; X64-NEXT: movl y@DTPOFF(%rax,%rdx), %ecx +; X64-NEXT: addl z@DTPOFF(%rax,%rdx), %ecx +; X64-NEXT: movl %ecx, %eax +; X64-NEXT: popq %rcx +; X64-NEXT: retq + %1 = load i32, ptr @y, align 4 + %2 = load i32, ptr @z, align 4 + %3 = add nsw i32 %1, %2 + ret i32 %3 +} -- GitLab From 092999e70b349ac521cab2648152ababeb12873f Mon Sep 17 00:00:00 2001 From: Jay Foad Date: Fri, 15 Mar 2024 14:08:43 +0000 Subject: [PATCH 005/782] [AMDGPU] Update checks in new test after #85370 --- llvm/test/CodeGen/AMDGPU/itofp.i128.ll | 2186 +++++++++++------------- 1 file changed, 1001 insertions(+), 1185 deletions(-) diff --git a/llvm/test/CodeGen/AMDGPU/itofp.i128.ll b/llvm/test/CodeGen/AMDGPU/itofp.i128.ll index e4e8d52addf8..bfeb214c5af8 100644 --- a/llvm/test/CodeGen/AMDGPU/itofp.i128.ll +++ b/llvm/test/CodeGen/AMDGPU/itofp.i128.ll @@ -11,7 +11,7 @@ define float @sitofp_i128_to_f32(i128 %x) { ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] ; SDAG-NEXT: v_mov_b32_e32 v4, 0 ; SDAG-NEXT: s_and_saveexec_b64 s[6:7], vcc -; SDAG-NEXT: s_cbranch_execz .LBB0_16 +; SDAG-NEXT: s_cbranch_execz .LBB0_14 ; SDAG-NEXT: ; %bb.1: ; %itofp-if-end ; SDAG-NEXT: v_ashrrev_i32_e32 v5, 31, v3 ; SDAG-NEXT: v_xor_b32_e32 v0, v5, v0 @@ -32,112 +32,100 @@ define float @sitofp_i128_to_f32(i128 %x) { ; SDAG-NEXT: v_min_u32_e32 v6, v6, v7 ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] ; SDAG-NEXT: v_add_u32_e32 v6, 64, v6 -; SDAG-NEXT: v_cndmask_b32_e32 v9, v6, v2, vcc -; SDAG-NEXT: v_sub_u32_e32 v2, 0x7f, v9 -; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 25, v2 -; SDAG-NEXT: ; implicit-def: $vgpr6 +; SDAG-NEXT: v_cndmask_b32_e32 v7, v6, v2, vcc +; SDAG-NEXT: v_sub_u32_e32 v6, 0x80, v7 +; SDAG-NEXT: v_sub_u32_e32 v2, 0x7f, v7 +; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 25, v6 +; SDAG-NEXT: ; implicit-def: $vgpr8 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc ; SDAG-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; SDAG-NEXT: ; %bb.2: ; %itofp-if-else -; SDAG-NEXT: v_add_u32_e32 v4, 0xffffff98, v9 +; SDAG-NEXT: v_add_u32_e32 v4, 0xffffff98, v7 ; SDAG-NEXT: v_lshlrev_b64 v[0:1], v4, v[0:1] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v4 -; SDAG-NEXT: v_cndmask_b32_e32 v6, 0, v0, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v8, 0, v0, vcc +; SDAG-NEXT: ; implicit-def: $vgpr6 ; SDAG-NEXT: ; implicit-def: $vgpr0_vgpr1 -; SDAG-NEXT: ; implicit-def: $vgpr9 +; SDAG-NEXT: ; implicit-def: $vgpr7 ; SDAG-NEXT: ; implicit-def: $vgpr4_vgpr5 -; SDAG-NEXT: ; %bb.3: ; %Flow6 +; SDAG-NEXT: ; %bb.3: ; %Flow3 ; SDAG-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; SDAG-NEXT: s_cbranch_execz .LBB0_15 +; SDAG-NEXT: s_cbranch_execz .LBB0_13 ; SDAG-NEXT: ; %bb.4: ; %NodeBlock -; SDAG-NEXT: v_sub_u32_e32 v8, 0x80, v9 -; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 25, v8 -; SDAG-NEXT: s_mov_b64 s[10:11], 0 -; SDAG-NEXT: s_mov_b64 s[4:5], 0 +; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 25, v6 +; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc +; SDAG-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; SDAG-NEXT: s_cbranch_execz .LBB0_8 +; SDAG-NEXT: ; %bb.5: ; %LeafBlock +; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 26, v6 ; SDAG-NEXT: s_and_saveexec_b64 s[12:13], vcc -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: ; %bb.5: ; %LeafBlock1 -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 26, v8 -; SDAG-NEXT: s_and_b64 s[4:5], vcc, exec -; SDAG-NEXT: ; %bb.6: ; %Flow3 -; SDAG-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; SDAG-NEXT: ; %bb.7: ; %LeafBlock -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 25, v8 -; SDAG-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; SDAG-NEXT: s_and_b64 s[14:15], vcc, exec -; SDAG-NEXT: s_mov_b64 s[10:11], exec -; SDAG-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; SDAG-NEXT: ; %bb.8: ; %Flow4 -; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: v_mov_b32_e32 v7, v1 -; SDAG-NEXT: v_mov_b32_e32 v6, v0 -; SDAG-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: s_cbranch_execz .LBB0_10 -; SDAG-NEXT: ; %bb.9: ; %itofp-sw-default -; SDAG-NEXT: v_sub_u32_e32 v12, 0x66, v9 +; SDAG-NEXT: s_cbranch_execz .LBB0_7 +; SDAG-NEXT: ; %bb.6: ; %itofp-sw-default +; SDAG-NEXT: v_sub_u32_e32 v12, 0x66, v7 ; SDAG-NEXT: v_sub_u32_e32 v10, 64, v12 -; SDAG-NEXT: v_lshrrev_b64 v[6:7], v12, v[0:1] +; SDAG-NEXT: v_lshrrev_b64 v[8:9], v12, v[0:1] ; SDAG-NEXT: v_lshlrev_b64 v[10:11], v10, v[4:5] -; SDAG-NEXT: v_sub_u32_e32 v13, 38, v9 -; SDAG-NEXT: v_or_b32_e32 v11, v7, v11 -; SDAG-NEXT: v_or_b32_e32 v10, v6, v10 -; SDAG-NEXT: v_lshrrev_b64 v[6:7], v13, v[4:5] +; SDAG-NEXT: v_sub_u32_e32 v13, 38, v7 +; SDAG-NEXT: v_or_b32_e32 v11, v9, v11 +; SDAG-NEXT: v_or_b32_e32 v10, v8, v10 +; SDAG-NEXT: v_lshrrev_b64 v[8:9], v13, v[4:5] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v12 -; SDAG-NEXT: v_add_u32_e32 v14, 26, v9 -; SDAG-NEXT: v_cndmask_b32_e32 v7, v7, v11, vcc +; SDAG-NEXT: v_add_u32_e32 v14, 26, v7 +; SDAG-NEXT: v_cndmask_b32_e32 v9, v9, v11, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v12 -; SDAG-NEXT: v_cndmask_b32_e32 v6, v6, v10, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v8, v8, v10, vcc ; SDAG-NEXT: v_lshrrev_b64 v[10:11], v13, v[0:1] ; SDAG-NEXT: v_lshlrev_b64 v[12:13], v14, v[4:5] -; SDAG-NEXT: v_subrev_u32_e32 v9, 38, v9 -; SDAG-NEXT: v_cndmask_b32_e64 v15, v6, v0, s[4:5] -; SDAG-NEXT: v_or_b32_e32 v6, v13, v11 -; SDAG-NEXT: v_or_b32_e32 v11, v12, v10 -; SDAG-NEXT: v_lshlrev_b64 v[9:10], v9, v[0:1] +; SDAG-NEXT: v_subrev_u32_e32 v7, 38, v7 +; SDAG-NEXT: v_cndmask_b32_e64 v15, v8, v0, s[4:5] +; SDAG-NEXT: v_lshlrev_b64 v[7:8], v7, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v9, v9, v1, s[4:5] +; SDAG-NEXT: v_or_b32_e32 v11, v13, v11 +; SDAG-NEXT: v_or_b32_e32 v10, v12, v10 ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v14 -; SDAG-NEXT: v_cndmask_b32_e64 v7, v7, v1, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v6, v10, v6, vcc +; SDAG-NEXT: v_lshlrev_b64 v[0:1], v14, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e32 v8, v8, v11, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v14 -; SDAG-NEXT: v_cndmask_b32_e64 v10, v6, v5, s[4:5] -; SDAG-NEXT: v_lshlrev_b64 v[5:6], v14, v[0:1] -; SDAG-NEXT: v_cndmask_b32_e32 v9, v9, v11, vcc -; SDAG-NEXT: v_cndmask_b32_e64 v4, v9, v4, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v6, 0, v6, vcc -; SDAG-NEXT: v_cndmask_b32_e32 v9, 0, v5, vcc -; SDAG-NEXT: v_or_b32_e32 v5, v6, v10 -; SDAG-NEXT: v_or_b32_e32 v4, v9, v4 -; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] -; SDAG-NEXT: s_andn2_b64 s[10:11], s[10:11], exec -; SDAG-NEXT: v_cndmask_b32_e64 v4, 0, 1, vcc -; SDAG-NEXT: v_or_b32_e32 v6, v15, v4 -; SDAG-NEXT: .LBB0_10: ; %Flow5 +; SDAG-NEXT: v_cndmask_b32_e32 v7, v7, v10, vcc +; SDAG-NEXT: v_cndmask_b32_e64 v5, v8, v5, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e64 v4, v7, v4, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v0, 0, v0, vcc +; SDAG-NEXT: v_or_b32_e32 v1, v1, v5 +; SDAG-NEXT: v_or_b32_e32 v0, v0, v4 +; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; SDAG-NEXT: v_or_b32_e32 v8, v15, v0 +; SDAG-NEXT: v_mov_b32_e32 v0, v8 +; SDAG-NEXT: v_mov_b32_e32 v1, v9 +; SDAG-NEXT: .LBB0_7: ; %Flow1 ; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; SDAG-NEXT: ; %bb.11: ; %itofp-sw-bb -; SDAG-NEXT: v_lshlrev_b64 v[6:7], 1, v[0:1] -; SDAG-NEXT: ; %bb.12: ; %itofp-sw-epilog +; SDAG-NEXT: .LBB0_8: ; %Flow2 +; SDAG-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; SDAG-NEXT: ; %bb.9: ; %itofp-sw-bb +; SDAG-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; SDAG-NEXT: ; %bb.10: ; %itofp-sw-epilog ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: v_lshrrev_b32_e32 v0, 2, v6 -; SDAG-NEXT: v_and_or_b32 v0, v0, 1, v6 +; SDAG-NEXT: v_lshrrev_b32_e32 v4, 2, v0 +; SDAG-NEXT: v_and_or_b32 v0, v4, 1, v0 ; SDAG-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v7, vcc +; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc ; SDAG-NEXT: v_and_b32_e32 v4, 0x4000000, v0 ; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 0, v4 -; SDAG-NEXT: v_alignbit_b32 v6, v1, v0, 2 +; SDAG-NEXT: v_alignbit_b32 v8, v1, v0, 2 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc -; SDAG-NEXT: ; %bb.13: ; %itofp-if-then20 -; SDAG-NEXT: v_alignbit_b32 v6, v1, v0, 3 -; SDAG-NEXT: v_mov_b32_e32 v2, v8 -; SDAG-NEXT: ; %bb.14: ; %Flow +; SDAG-NEXT: ; %bb.11: ; %itofp-if-then20 +; SDAG-NEXT: v_alignbit_b32 v8, v1, v0, 3 +; SDAG-NEXT: v_mov_b32_e32 v2, v6 +; SDAG-NEXT: ; %bb.12: ; %Flow ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: .LBB0_15: ; %Flow7 +; SDAG-NEXT: .LBB0_13: ; %Flow4 ; SDAG-NEXT: s_or_b64 exec, exec, s[8:9] ; SDAG-NEXT: v_and_b32_e32 v0, 0x80000000, v3 ; SDAG-NEXT: v_lshl_add_u32 v1, v2, 23, 1.0 -; SDAG-NEXT: v_and_b32_e32 v2, 0x7fffff, v6 +; SDAG-NEXT: v_and_b32_e32 v2, 0x7fffff, v8 ; SDAG-NEXT: v_or3_b32 v4, v2, v0, v1 -; SDAG-NEXT: .LBB0_16: ; %Flow8 +; SDAG-NEXT: .LBB0_14: ; %Flow5 ; SDAG-NEXT: s_or_b64 exec, exec, s[6:7] ; SDAG-NEXT: v_mov_b32_e32 v0, v4 ; SDAG-NEXT: s_setpc_b64 s[30:31] @@ -151,144 +139,126 @@ define float @sitofp_i128_to_f32(i128 %x) { ; GISEL-NEXT: s_mov_b32 s4, 0 ; GISEL-NEXT: v_mov_b32_e32 v4, s4 ; GISEL-NEXT: s_and_saveexec_b64 s[6:7], vcc -; GISEL-NEXT: s_cbranch_execz .LBB0_16 +; GISEL-NEXT: s_cbranch_execz .LBB0_14 ; GISEL-NEXT: ; %bb.1: ; %itofp-if-end -; GISEL-NEXT: v_ashrrev_i32_e32 v8, 31, v3 -; GISEL-NEXT: v_xor_b32_e32 v0, v8, v0 -; GISEL-NEXT: v_xor_b32_e32 v1, v8, v1 -; GISEL-NEXT: v_sub_co_u32_e32 v0, vcc, v0, v8 -; GISEL-NEXT: v_xor_b32_e32 v2, v8, v2 -; GISEL-NEXT: v_subb_co_u32_e32 v1, vcc, v1, v8, vcc -; GISEL-NEXT: v_xor_b32_e32 v3, v8, v3 -; GISEL-NEXT: v_subb_co_u32_e32 v6, vcc, v2, v8, vcc -; GISEL-NEXT: v_subb_co_u32_e32 v7, vcc, v3, v8, vcc -; GISEL-NEXT: v_ffbh_u32_e32 v3, v0 -; GISEL-NEXT: v_ffbh_u32_e32 v2, v1 -; GISEL-NEXT: v_add_u32_e32 v3, 32, v3 -; GISEL-NEXT: v_ffbh_u32_e32 v4, v6 -; GISEL-NEXT: v_min_u32_e32 v2, v2, v3 -; GISEL-NEXT: v_ffbh_u32_e32 v3, v7 -; GISEL-NEXT: v_add_u32_e32 v4, 32, v4 -; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[6:7] -; GISEL-NEXT: v_add_u32_e32 v2, 64, v2 -; GISEL-NEXT: v_min_u32_e32 v3, v3, v4 -; GISEL-NEXT: v_cndmask_b32_e32 v11, v3, v2, vcc -; GISEL-NEXT: v_sub_u32_e32 v9, 0x7f, v11 -; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 24, v9 -; GISEL-NEXT: ; implicit-def: $vgpr2 +; GISEL-NEXT: v_ashrrev_i32_e32 v6, 31, v3 +; GISEL-NEXT: v_xor_b32_e32 v0, v6, v0 +; GISEL-NEXT: v_xor_b32_e32 v1, v6, v1 +; GISEL-NEXT: v_sub_co_u32_e32 v0, vcc, v0, v6 +; GISEL-NEXT: v_xor_b32_e32 v2, v6, v2 +; GISEL-NEXT: v_subb_co_u32_e32 v1, vcc, v1, v6, vcc +; GISEL-NEXT: v_xor_b32_e32 v3, v6, v3 +; GISEL-NEXT: v_subb_co_u32_e32 v2, vcc, v2, v6, vcc +; GISEL-NEXT: v_ffbh_u32_e32 v5, v0 +; GISEL-NEXT: v_subb_co_u32_e32 v3, vcc, v3, v6, vcc +; GISEL-NEXT: v_ffbh_u32_e32 v4, v1 +; GISEL-NEXT: v_add_u32_e32 v5, 32, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v7, v2 +; GISEL-NEXT: v_min_u32_e32 v4, v4, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v5, v3 +; GISEL-NEXT: v_add_u32_e32 v7, 32, v7 +; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[2:3] +; GISEL-NEXT: v_add_u32_e32 v4, 64, v4 +; GISEL-NEXT: v_min_u32_e32 v5, v5, v7 +; GISEL-NEXT: v_cndmask_b32_e32 v5, v5, v4, vcc +; GISEL-NEXT: v_sub_u32_e32 v8, 0x80, v5 +; GISEL-NEXT: v_sub_u32_e32 v7, 0x7f, v5 +; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 24, v8 +; GISEL-NEXT: ; implicit-def: $vgpr4 ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc ; GISEL-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; GISEL-NEXT: ; %bb.2: ; %itofp-if-else -; GISEL-NEXT: v_add_u32_e32 v2, 0xffffff98, v11 +; GISEL-NEXT: v_add_u32_e32 v2, 0xffffff98, v5 ; GISEL-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 -; GISEL-NEXT: v_cndmask_b32_e32 v2, 0, v0, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v0, vcc +; GISEL-NEXT: ; implicit-def: $vgpr8 ; GISEL-NEXT: ; implicit-def: $vgpr0 -; GISEL-NEXT: ; implicit-def: $vgpr11 -; GISEL-NEXT: ; implicit-def: $vgpr6 -; GISEL-NEXT: ; %bb.3: ; %Flow6 +; GISEL-NEXT: ; implicit-def: $vgpr5 +; GISEL-NEXT: ; implicit-def: $vgpr2 +; GISEL-NEXT: ; %bb.3: ; %Flow3 ; GISEL-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; GISEL-NEXT: s_cbranch_execz .LBB0_15 +; GISEL-NEXT: s_cbranch_execz .LBB0_13 ; GISEL-NEXT: ; %bb.4: ; %NodeBlock -; GISEL-NEXT: v_sub_u32_e32 v10, 0x80, v11 -; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 26, v10 -; GISEL-NEXT: s_mov_b64 s[10:11], 0 -; GISEL-NEXT: s_mov_b64 s[4:5], 0 +; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 26, v8 +; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc +; GISEL-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; GISEL-NEXT: s_cbranch_execz .LBB0_8 +; GISEL-NEXT: ; %bb.5: ; %LeafBlock +; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 26, v8 ; GISEL-NEXT: s_and_saveexec_b64 s[12:13], vcc -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: ; %bb.5: ; %LeafBlock1 -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 26, v10 -; GISEL-NEXT: s_andn2_b64 s[4:5], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.6: ; %Flow3 -; GISEL-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; GISEL-NEXT: ; %bb.7: ; %LeafBlock -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 25, v10 -; GISEL-NEXT: s_andn2_b64 s[10:11], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, -1 -; GISEL-NEXT: s_or_b64 s[10:11], s[10:11], s[14:15] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.8: ; %Flow4 -; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: v_mov_b32_e32 v5, v3 -; GISEL-NEXT: v_mov_b32_e32 v4, v2 -; GISEL-NEXT: v_mov_b32_e32 v3, v1 -; GISEL-NEXT: v_mov_b32_e32 v2, v0 -; GISEL-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: s_cbranch_execz .LBB0_10 -; GISEL-NEXT: ; %bb.9: ; %itofp-sw-default -; GISEL-NEXT: v_sub_u32_e32 v12, 0x66, v11 -; GISEL-NEXT: v_sub_u32_e32 v4, 64, v12 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v12, v[0:1] -; GISEL-NEXT: v_lshlrev_b64 v[4:5], v4, v[6:7] -; GISEL-NEXT: v_subrev_u32_e32 v13, 64, v12 -; GISEL-NEXT: v_or_b32_e32 v4, v2, v4 -; GISEL-NEXT: v_or_b32_e32 v5, v3, v5 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v13, v[6:7] -; GISEL-NEXT: v_add_u32_e32 v13, 26, v11 -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v12 -; GISEL-NEXT: v_sub_u32_e32 v11, 64, v13 -; GISEL-NEXT: v_cndmask_b32_e32 v2, v2, v4, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v3, v3, v5, vcc -; GISEL-NEXT: v_cmp_eq_u32_e32 vcc, 0, v12 -; GISEL-NEXT: v_lshrrev_b64 v[4:5], v13, -1 +; GISEL-NEXT: s_cbranch_execz .LBB0_7 +; GISEL-NEXT: ; %bb.6: ; %itofp-sw-default +; GISEL-NEXT: v_sub_u32_e32 v4, 0x66, v5 +; GISEL-NEXT: v_sub_u32_e32 v11, 64, v4 +; GISEL-NEXT: v_lshrrev_b64 v[9:10], v4, v[0:1] +; GISEL-NEXT: v_lshlrev_b64 v[11:12], v11, v[2:3] +; GISEL-NEXT: v_subrev_u32_e32 v13, 64, v4 +; GISEL-NEXT: v_or_b32_e32 v11, v9, v11 +; GISEL-NEXT: v_or_b32_e32 v12, v10, v12 +; GISEL-NEXT: v_lshrrev_b64 v[9:10], v13, v[2:3] +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v4 +; GISEL-NEXT: v_add_u32_e32 v5, 26, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v9, v9, v11, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v10, v10, v12, vcc +; GISEL-NEXT: v_cmp_eq_u32_e32 vcc, 0, v4 +; GISEL-NEXT: v_sub_u32_e32 v11, 64, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v13, v9, v0, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v4, v10, v1, vcc +; GISEL-NEXT: v_lshrrev_b64 v[9:10], v5, -1 ; GISEL-NEXT: v_lshlrev_b64 v[11:12], v11, -1 -; GISEL-NEXT: v_subrev_u32_e32 v14, 64, v13 -; GISEL-NEXT: v_or_b32_e32 v15, v4, v11 -; GISEL-NEXT: v_or_b32_e32 v16, v5, v12 +; GISEL-NEXT: v_subrev_u32_e32 v14, 64, v5 +; GISEL-NEXT: v_or_b32_e32 v15, v9, v11 +; GISEL-NEXT: v_or_b32_e32 v16, v10, v12 ; GISEL-NEXT: v_lshrrev_b64 v[11:12], v14, -1 -; GISEL-NEXT: v_cndmask_b32_e32 v2, v2, v0, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v3, v3, v1, vcc -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v13 +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v5 ; GISEL-NEXT: v_cndmask_b32_e32 v11, v11, v15, vcc ; GISEL-NEXT: v_cndmask_b32_e32 v12, v12, v16, vcc -; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v13 -; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v5, 0, v5, vcc -; GISEL-NEXT: v_cndmask_b32_e64 v11, v11, -1, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e64 v12, v12, -1, s[4:5] -; GISEL-NEXT: v_and_b32_e32 v4, v4, v6 -; GISEL-NEXT: v_and_b32_e32 v5, v5, v7 -; GISEL-NEXT: v_and_or_b32 v4, v11, v0, v4 -; GISEL-NEXT: v_and_or_b32 v5, v12, v1, v5 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[10:11], exec -; GISEL-NEXT: v_cndmask_b32_e64 v4, 0, 1, vcc -; GISEL-NEXT: s_and_b64 s[10:11], exec, 0 -; GISEL-NEXT: v_or_b32_e32 v2, v2, v4 -; GISEL-NEXT: s_or_b64 s[10:11], s[4:5], s[10:11] -; GISEL-NEXT: .LBB0_10: ; %Flow5 +; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v9, 0, v9, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v10, 0, v10, vcc +; GISEL-NEXT: v_cndmask_b32_e64 v5, v11, -1, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e64 v11, v12, -1, s[4:5] +; GISEL-NEXT: v_and_b32_e32 v2, v9, v2 +; GISEL-NEXT: v_and_b32_e32 v3, v10, v3 +; GISEL-NEXT: v_and_or_b32 v0, v5, v0, v2 +; GISEL-NEXT: v_and_or_b32 v1, v11, v1, v3 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; GISEL-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; GISEL-NEXT: v_or_b32_e32 v3, v13, v0 +; GISEL-NEXT: v_mov_b32_e32 v0, v3 +; GISEL-NEXT: v_mov_b32_e32 v1, v4 +; GISEL-NEXT: v_mov_b32_e32 v2, v5 +; GISEL-NEXT: v_mov_b32_e32 v3, v6 +; GISEL-NEXT: .LBB0_7: ; %Flow1 ; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; GISEL-NEXT: ; %bb.11: ; %itofp-sw-bb -; GISEL-NEXT: v_lshlrev_b64 v[2:3], 1, v[0:1] -; GISEL-NEXT: ; %bb.12: ; %itofp-sw-epilog +; GISEL-NEXT: .LBB0_8: ; %Flow2 +; GISEL-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; GISEL-NEXT: ; %bb.9: ; %itofp-sw-bb +; GISEL-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; GISEL-NEXT: ; %bb.10: ; %itofp-sw-epilog ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: v_bfe_u32 v0, v2, 2, 1 -; GISEL-NEXT: v_or_b32_e32 v0, v2, v0 +; GISEL-NEXT: v_bfe_u32 v2, v0, 2, 1 +; GISEL-NEXT: v_or_b32_e32 v0, v0, v2 ; GISEL-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v3, vcc -; GISEL-NEXT: v_lshrrev_b64 v[2:3], 2, v[0:1] -; GISEL-NEXT: v_and_b32_e32 v3, 0x4000000, v0 -; GISEL-NEXT: v_mov_b32_e32 v4, 0 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[3:4] +; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc +; GISEL-NEXT: v_and_b32_e32 v2, 0x4000000, v0 +; GISEL-NEXT: v_mov_b32_e32 v3, 0 +; GISEL-NEXT: v_lshrrev_b64 v[4:5], 2, v[0:1] +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc -; GISEL-NEXT: ; %bb.13: ; %itofp-if-then20 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], 3, v[0:1] -; GISEL-NEXT: v_mov_b32_e32 v9, v10 -; GISEL-NEXT: ; %bb.14: ; %Flow +; GISEL-NEXT: ; %bb.11: ; %itofp-if-then20 +; GISEL-NEXT: v_lshrrev_b64 v[4:5], 3, v[0:1] +; GISEL-NEXT: v_mov_b32_e32 v7, v8 +; GISEL-NEXT: ; %bb.12: ; %Flow ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: .LBB0_15: ; %Flow7 +; GISEL-NEXT: .LBB0_13: ; %Flow4 ; GISEL-NEXT: s_or_b64 exec, exec, s[8:9] -; GISEL-NEXT: v_and_b32_e32 v0, 0x80000000, v8 -; GISEL-NEXT: v_lshl_add_u32 v1, v9, 23, 1.0 -; GISEL-NEXT: v_and_b32_e32 v2, 0x7fffff, v2 +; GISEL-NEXT: v_and_b32_e32 v0, 0x80000000, v6 +; GISEL-NEXT: v_lshl_add_u32 v1, v7, 23, 1.0 +; GISEL-NEXT: v_and_b32_e32 v2, 0x7fffff, v4 ; GISEL-NEXT: v_or3_b32 v4, v2, v0, v1 -; GISEL-NEXT: .LBB0_16: ; %Flow8 +; GISEL-NEXT: .LBB0_14: ; %Flow5 ; GISEL-NEXT: s_or_b64 exec, exec, s[6:7] ; GISEL-NEXT: v_mov_b32_e32 v0, v4 ; GISEL-NEXT: s_setpc_b64 s[30:31] @@ -305,7 +275,7 @@ define float @uitofp_i128_to_f32(i128 %x) { ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] ; SDAG-NEXT: v_mov_b32_e32 v4, 0 ; SDAG-NEXT: s_and_saveexec_b64 s[6:7], vcc -; SDAG-NEXT: s_cbranch_execz .LBB1_16 +; SDAG-NEXT: s_cbranch_execz .LBB1_14 ; SDAG-NEXT: ; %bb.1: ; %itofp-if-end ; SDAG-NEXT: v_ffbh_u32_e32 v4, v2 ; SDAG-NEXT: v_add_u32_e32 v4, 32, v4 @@ -317,113 +287,99 @@ define float @uitofp_i128_to_f32(i128 %x) { ; SDAG-NEXT: v_min_u32_e32 v5, v5, v6 ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] ; SDAG-NEXT: v_add_u32_e32 v5, 64, v5 -; SDAG-NEXT: v_cndmask_b32_e32 v8, v5, v4, vcc -; SDAG-NEXT: v_sub_u32_e32 v6, 0x7f, v8 -; SDAG-NEXT: v_mov_b32_e32 v5, v1 -; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 25, v6 -; SDAG-NEXT: v_mov_b32_e32 v4, v0 -; SDAG-NEXT: ; implicit-def: $vgpr9 +; SDAG-NEXT: v_cndmask_b32_e32 v6, v5, v4, vcc +; SDAG-NEXT: v_sub_u32_e32 v5, 0x80, v6 +; SDAG-NEXT: v_sub_u32_e32 v4, 0x7f, v6 +; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 25, v5 +; SDAG-NEXT: ; implicit-def: $vgpr7 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc ; SDAG-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; SDAG-NEXT: ; %bb.2: ; %itofp-if-else -; SDAG-NEXT: v_add_u32_e32 v2, 0xffffff98, v8 +; SDAG-NEXT: v_add_u32_e32 v2, 0xffffff98, v6 ; SDAG-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 -; SDAG-NEXT: v_cndmask_b32_e32 v9, 0, v0, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v7, 0, v0, vcc +; SDAG-NEXT: ; implicit-def: $vgpr5 ; SDAG-NEXT: ; implicit-def: $vgpr0_vgpr1 -; SDAG-NEXT: ; implicit-def: $vgpr8 +; SDAG-NEXT: ; implicit-def: $vgpr6 ; SDAG-NEXT: ; implicit-def: $vgpr2_vgpr3 -; SDAG-NEXT: ; implicit-def: $vgpr4_vgpr5 -; SDAG-NEXT: ; %bb.3: ; %Flow6 +; SDAG-NEXT: ; %bb.3: ; %Flow3 ; SDAG-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; SDAG-NEXT: s_cbranch_execz .LBB1_15 +; SDAG-NEXT: s_cbranch_execz .LBB1_13 ; SDAG-NEXT: ; %bb.4: ; %NodeBlock -; SDAG-NEXT: v_sub_u32_e32 v7, 0x80, v8 -; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 25, v7 -; SDAG-NEXT: s_mov_b64 s[10:11], 0 -; SDAG-NEXT: s_mov_b64 s[4:5], 0 +; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 25, v5 +; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc +; SDAG-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; SDAG-NEXT: s_cbranch_execz .LBB1_8 +; SDAG-NEXT: ; %bb.5: ; %LeafBlock +; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 26, v5 ; SDAG-NEXT: s_and_saveexec_b64 s[12:13], vcc -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: ; %bb.5: ; %LeafBlock1 -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 26, v7 -; SDAG-NEXT: s_and_b64 s[4:5], vcc, exec -; SDAG-NEXT: ; %bb.6: ; %Flow3 -; SDAG-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; SDAG-NEXT: ; %bb.7: ; %LeafBlock -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 25, v7 -; SDAG-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; SDAG-NEXT: s_and_b64 s[14:15], vcc, exec -; SDAG-NEXT: s_mov_b64 s[10:11], exec -; SDAG-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; SDAG-NEXT: ; implicit-def: $vgpr4_vgpr5 -; SDAG-NEXT: ; %bb.8: ; %Flow4 -; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: s_cbranch_execz .LBB1_10 -; SDAG-NEXT: ; %bb.9: ; %itofp-sw-default -; SDAG-NEXT: v_sub_u32_e32 v11, 0x66, v8 +; SDAG-NEXT: s_cbranch_execz .LBB1_7 +; SDAG-NEXT: ; %bb.6: ; %itofp-sw-default +; SDAG-NEXT: v_sub_u32_e32 v11, 0x66, v6 ; SDAG-NEXT: v_sub_u32_e32 v9, 64, v11 -; SDAG-NEXT: v_lshrrev_b64 v[4:5], v11, v[0:1] +; SDAG-NEXT: v_lshrrev_b64 v[7:8], v11, v[0:1] ; SDAG-NEXT: v_lshlrev_b64 v[9:10], v9, v[2:3] -; SDAG-NEXT: v_sub_u32_e32 v12, 38, v8 -; SDAG-NEXT: v_or_b32_e32 v10, v5, v10 -; SDAG-NEXT: v_or_b32_e32 v9, v4, v9 -; SDAG-NEXT: v_lshrrev_b64 v[4:5], v12, v[2:3] +; SDAG-NEXT: v_sub_u32_e32 v12, 38, v6 +; SDAG-NEXT: v_or_b32_e32 v10, v8, v10 +; SDAG-NEXT: v_or_b32_e32 v9, v7, v9 +; SDAG-NEXT: v_lshrrev_b64 v[7:8], v12, v[2:3] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v11 -; SDAG-NEXT: v_add_u32_e32 v13, 26, v8 -; SDAG-NEXT: v_cndmask_b32_e32 v5, v5, v10, vcc +; SDAG-NEXT: v_add_u32_e32 v13, 26, v6 +; SDAG-NEXT: v_cndmask_b32_e32 v8, v8, v10, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v11 -; SDAG-NEXT: v_cndmask_b32_e32 v4, v4, v9, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v7, v7, v9, vcc ; SDAG-NEXT: v_lshrrev_b64 v[9:10], v12, v[0:1] ; SDAG-NEXT: v_lshlrev_b64 v[11:12], v13, v[2:3] -; SDAG-NEXT: v_subrev_u32_e32 v8, 38, v8 -; SDAG-NEXT: v_cndmask_b32_e64 v14, v4, v0, s[4:5] -; SDAG-NEXT: v_or_b32_e32 v4, v12, v10 -; SDAG-NEXT: v_or_b32_e32 v10, v11, v9 -; SDAG-NEXT: v_lshlrev_b64 v[8:9], v8, v[0:1] +; SDAG-NEXT: v_subrev_u32_e32 v6, 38, v6 +; SDAG-NEXT: v_cndmask_b32_e64 v14, v7, v0, s[4:5] +; SDAG-NEXT: v_lshlrev_b64 v[6:7], v6, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v8, v8, v1, s[4:5] +; SDAG-NEXT: v_or_b32_e32 v10, v12, v10 +; SDAG-NEXT: v_or_b32_e32 v9, v11, v9 ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v13 -; SDAG-NEXT: v_cndmask_b32_e64 v5, v5, v1, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v4, v9, v4, vcc +; SDAG-NEXT: v_lshlrev_b64 v[0:1], v13, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e32 v7, v7, v10, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v13 -; SDAG-NEXT: v_cndmask_b32_e64 v9, v4, v3, s[4:5] -; SDAG-NEXT: v_lshlrev_b64 v[3:4], v13, v[0:1] -; SDAG-NEXT: v_cndmask_b32_e32 v8, v8, v10, vcc -; SDAG-NEXT: v_cndmask_b32_e64 v2, v8, v2, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc -; SDAG-NEXT: v_cndmask_b32_e32 v8, 0, v3, vcc -; SDAG-NEXT: v_or_b32_e32 v3, v4, v9 -; SDAG-NEXT: v_or_b32_e32 v2, v8, v2 -; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] -; SDAG-NEXT: s_andn2_b64 s[10:11], s[10:11], exec -; SDAG-NEXT: v_cndmask_b32_e64 v2, 0, 1, vcc -; SDAG-NEXT: v_or_b32_e32 v4, v14, v2 -; SDAG-NEXT: .LBB1_10: ; %Flow5 +; SDAG-NEXT: v_cndmask_b32_e32 v6, v6, v9, vcc +; SDAG-NEXT: v_cndmask_b32_e64 v3, v7, v3, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e64 v2, v6, v2, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v0, 0, v0, vcc +; SDAG-NEXT: v_or_b32_e32 v1, v1, v3 +; SDAG-NEXT: v_or_b32_e32 v0, v0, v2 +; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; SDAG-NEXT: v_or_b32_e32 v7, v14, v0 +; SDAG-NEXT: v_mov_b32_e32 v0, v7 +; SDAG-NEXT: v_mov_b32_e32 v1, v8 +; SDAG-NEXT: .LBB1_7: ; %Flow1 ; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; SDAG-NEXT: ; %bb.11: ; %itofp-sw-bb -; SDAG-NEXT: v_lshlrev_b64 v[4:5], 1, v[0:1] -; SDAG-NEXT: ; %bb.12: ; %itofp-sw-epilog +; SDAG-NEXT: .LBB1_8: ; %Flow2 +; SDAG-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; SDAG-NEXT: ; %bb.9: ; %itofp-sw-bb +; SDAG-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; SDAG-NEXT: ; %bb.10: ; %itofp-sw-epilog ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: v_lshrrev_b32_e32 v0, 2, v4 -; SDAG-NEXT: v_and_or_b32 v0, v0, 1, v4 +; SDAG-NEXT: v_lshrrev_b32_e32 v2, 2, v0 +; SDAG-NEXT: v_and_or_b32 v0, v2, 1, v0 ; SDAG-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v5, vcc +; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc ; SDAG-NEXT: v_and_b32_e32 v2, 0x4000000, v0 ; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 0, v2 -; SDAG-NEXT: v_alignbit_b32 v9, v1, v0, 2 +; SDAG-NEXT: v_alignbit_b32 v7, v1, v0, 2 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc -; SDAG-NEXT: ; %bb.13: ; %itofp-if-then20 -; SDAG-NEXT: v_alignbit_b32 v9, v1, v0, 3 -; SDAG-NEXT: v_mov_b32_e32 v6, v7 -; SDAG-NEXT: ; %bb.14: ; %Flow +; SDAG-NEXT: ; %bb.11: ; %itofp-if-then20 +; SDAG-NEXT: v_alignbit_b32 v7, v1, v0, 3 +; SDAG-NEXT: v_mov_b32_e32 v4, v5 +; SDAG-NEXT: ; %bb.12: ; %Flow ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: .LBB1_15: ; %Flow7 +; SDAG-NEXT: .LBB1_13: ; %Flow4 ; SDAG-NEXT: s_or_b64 exec, exec, s[8:9] -; SDAG-NEXT: v_and_b32_e32 v0, 0x7fffff, v9 -; SDAG-NEXT: v_lshl_or_b32 v0, v6, 23, v0 +; SDAG-NEXT: v_and_b32_e32 v0, 0x7fffff, v7 +; SDAG-NEXT: v_lshl_or_b32 v0, v4, 23, v0 ; SDAG-NEXT: v_add_u32_e32 v4, 1.0, v0 -; SDAG-NEXT: .LBB1_16: ; %Flow8 +; SDAG-NEXT: .LBB1_14: ; %Flow5 ; SDAG-NEXT: s_or_b64 exec, exec, s[6:7] ; SDAG-NEXT: v_mov_b32_e32 v0, v4 ; SDAG-NEXT: s_setpc_b64 s[30:31] @@ -431,144 +387,124 @@ define float @uitofp_i128_to_f32(i128 %x) { ; GISEL-LABEL: uitofp_i128_to_f32: ; GISEL: ; %bb.0: ; %itofp-entry ; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) -; GISEL-NEXT: v_mov_b32_e32 v4, v2 -; GISEL-NEXT: v_mov_b32_e32 v5, v3 -; GISEL-NEXT: v_or_b32_e32 v2, v0, v4 -; GISEL-NEXT: v_or_b32_e32 v3, v1, v5 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] +; GISEL-NEXT: v_or_b32_e32 v4, v0, v2 +; GISEL-NEXT: v_or_b32_e32 v5, v1, v3 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] ; GISEL-NEXT: s_mov_b32 s4, 0 -; GISEL-NEXT: v_mov_b32_e32 v2, s4 +; GISEL-NEXT: v_mov_b32_e32 v4, s4 ; GISEL-NEXT: s_and_saveexec_b64 s[6:7], vcc -; GISEL-NEXT: s_cbranch_execz .LBB1_16 +; GISEL-NEXT: s_cbranch_execz .LBB1_14 ; GISEL-NEXT: ; %bb.1: ; %itofp-if-end -; GISEL-NEXT: v_ffbh_u32_e32 v3, v0 -; GISEL-NEXT: v_ffbh_u32_e32 v2, v1 -; GISEL-NEXT: v_add_u32_e32 v3, 32, v3 -; GISEL-NEXT: v_ffbh_u32_e32 v6, v4 -; GISEL-NEXT: v_min_u32_e32 v2, v2, v3 -; GISEL-NEXT: v_ffbh_u32_e32 v3, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v5, v0 +; GISEL-NEXT: v_ffbh_u32_e32 v4, v1 +; GISEL-NEXT: v_add_u32_e32 v5, 32, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v6, v2 +; GISEL-NEXT: v_min_u32_e32 v4, v4, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v5, v3 ; GISEL-NEXT: v_add_u32_e32 v6, 32, v6 -; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[4:5] -; GISEL-NEXT: v_add_u32_e32 v2, 64, v2 -; GISEL-NEXT: v_min_u32_e32 v3, v3, v6 -; GISEL-NEXT: v_cndmask_b32_e32 v12, v3, v2, vcc -; GISEL-NEXT: v_sub_u32_e32 v10, 0x7f, v12 -; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 24, v10 -; GISEL-NEXT: ; implicit-def: $vgpr2 +; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[2:3] +; GISEL-NEXT: v_add_u32_e32 v4, 64, v4 +; GISEL-NEXT: v_min_u32_e32 v5, v5, v6 +; GISEL-NEXT: v_cndmask_b32_e32 v5, v5, v4, vcc +; GISEL-NEXT: v_sub_u32_e32 v7, 0x80, v5 +; GISEL-NEXT: v_sub_u32_e32 v6, 0x7f, v5 +; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 24, v7 +; GISEL-NEXT: ; implicit-def: $vgpr4 ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc ; GISEL-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; GISEL-NEXT: ; %bb.2: ; %itofp-if-else -; GISEL-NEXT: v_add_u32_e32 v2, 0xffffff98, v12 +; GISEL-NEXT: v_add_u32_e32 v2, 0xffffff98, v5 ; GISEL-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 -; GISEL-NEXT: v_cndmask_b32_e32 v2, 0, v0, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v0, vcc +; GISEL-NEXT: ; implicit-def: $vgpr7 ; GISEL-NEXT: ; implicit-def: $vgpr0 -; GISEL-NEXT: ; implicit-def: $vgpr12 -; GISEL-NEXT: ; implicit-def: $vgpr4 -; GISEL-NEXT: ; %bb.3: ; %Flow6 +; GISEL-NEXT: ; implicit-def: $vgpr5 +; GISEL-NEXT: ; implicit-def: $vgpr2 +; GISEL-NEXT: ; %bb.3: ; %Flow3 ; GISEL-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; GISEL-NEXT: s_cbranch_execz .LBB1_15 +; GISEL-NEXT: s_cbranch_execz .LBB1_13 ; GISEL-NEXT: ; %bb.4: ; %NodeBlock -; GISEL-NEXT: v_sub_u32_e32 v11, 0x80, v12 -; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 26, v11 -; GISEL-NEXT: s_mov_b64 s[10:11], 0 -; GISEL-NEXT: s_mov_b64 s[4:5], 0 +; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 26, v7 +; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc +; GISEL-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; GISEL-NEXT: s_cbranch_execz .LBB1_8 +; GISEL-NEXT: ; %bb.5: ; %LeafBlock +; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 26, v7 ; GISEL-NEXT: s_and_saveexec_b64 s[12:13], vcc -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: ; %bb.5: ; %LeafBlock1 -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 26, v11 -; GISEL-NEXT: s_andn2_b64 s[4:5], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.6: ; %Flow3 -; GISEL-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; GISEL-NEXT: ; %bb.7: ; %LeafBlock -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 25, v11 -; GISEL-NEXT: s_andn2_b64 s[10:11], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, -1 -; GISEL-NEXT: s_or_b64 s[10:11], s[10:11], s[14:15] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.8: ; %Flow4 -; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: v_mov_b32_e32 v9, v3 -; GISEL-NEXT: v_mov_b32_e32 v7, v1 -; GISEL-NEXT: v_mov_b32_e32 v6, v0 -; GISEL-NEXT: v_mov_b32_e32 v8, v2 -; GISEL-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: s_cbranch_execz .LBB1_10 -; GISEL-NEXT: ; %bb.9: ; %itofp-sw-default -; GISEL-NEXT: v_sub_u32_e32 v8, 0x66, v12 -; GISEL-NEXT: v_sub_u32_e32 v6, 64, v8 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v8, v[0:1] -; GISEL-NEXT: v_lshlrev_b64 v[6:7], v6, v[4:5] -; GISEL-NEXT: v_subrev_u32_e32 v9, 64, v8 -; GISEL-NEXT: v_or_b32_e32 v6, v2, v6 -; GISEL-NEXT: v_or_b32_e32 v7, v3, v7 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v9, v[4:5] -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v8 -; GISEL-NEXT: v_add_u32_e32 v12, 26, v12 -; GISEL-NEXT: v_cndmask_b32_e32 v2, v2, v6, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v3, v3, v7, vcc -; GISEL-NEXT: v_cmp_eq_u32_e32 vcc, 0, v8 -; GISEL-NEXT: v_sub_u32_e32 v8, 64, v12 -; GISEL-NEXT: v_cndmask_b32_e32 v6, v2, v0, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v7, v3, v1, vcc -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v12, -1 -; GISEL-NEXT: v_lshlrev_b64 v[8:9], v8, -1 -; GISEL-NEXT: v_subrev_u32_e32 v13, 64, v12 -; GISEL-NEXT: v_or_b32_e32 v14, v2, v8 -; GISEL-NEXT: v_or_b32_e32 v15, v3, v9 -; GISEL-NEXT: v_lshrrev_b64 v[8:9], v13, -1 -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v12 -; GISEL-NEXT: v_cndmask_b32_e32 v8, v8, v14, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v9, v9, v15, vcc -; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v12 -; GISEL-NEXT: v_cndmask_b32_e32 v2, 0, v2, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc -; GISEL-NEXT: v_cndmask_b32_e64 v8, v8, -1, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e64 v9, v9, -1, s[4:5] -; GISEL-NEXT: v_and_b32_e32 v2, v2, v4 -; GISEL-NEXT: v_and_b32_e32 v3, v3, v5 -; GISEL-NEXT: v_and_or_b32 v2, v8, v0, v2 -; GISEL-NEXT: v_and_or_b32 v3, v9, v1, v3 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[10:11], exec -; GISEL-NEXT: v_cndmask_b32_e64 v2, 0, 1, vcc -; GISEL-NEXT: s_and_b64 s[10:11], exec, 0 -; GISEL-NEXT: v_or_b32_e32 v6, v6, v2 -; GISEL-NEXT: s_or_b64 s[10:11], s[4:5], s[10:11] -; GISEL-NEXT: .LBB1_10: ; %Flow5 +; GISEL-NEXT: s_cbranch_execz .LBB1_7 +; GISEL-NEXT: ; %bb.6: ; %itofp-sw-default +; GISEL-NEXT: v_sub_u32_e32 v4, 0x66, v5 +; GISEL-NEXT: v_sub_u32_e32 v10, 64, v4 +; GISEL-NEXT: v_lshrrev_b64 v[8:9], v4, v[0:1] +; GISEL-NEXT: v_lshlrev_b64 v[10:11], v10, v[2:3] +; GISEL-NEXT: v_subrev_u32_e32 v12, 64, v4 +; GISEL-NEXT: v_or_b32_e32 v10, v8, v10 +; GISEL-NEXT: v_or_b32_e32 v11, v9, v11 +; GISEL-NEXT: v_lshrrev_b64 v[8:9], v12, v[2:3] +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v4 +; GISEL-NEXT: v_add_u32_e32 v5, 26, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v8, v8, v10, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v9, v9, v11, vcc +; GISEL-NEXT: v_cmp_eq_u32_e32 vcc, 0, v4 +; GISEL-NEXT: v_sub_u32_e32 v10, 64, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v12, v8, v0, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v4, v9, v1, vcc +; GISEL-NEXT: v_lshrrev_b64 v[8:9], v5, -1 +; GISEL-NEXT: v_lshlrev_b64 v[10:11], v10, -1 +; GISEL-NEXT: v_subrev_u32_e32 v13, 64, v5 +; GISEL-NEXT: v_or_b32_e32 v14, v8, v10 +; GISEL-NEXT: v_or_b32_e32 v15, v9, v11 +; GISEL-NEXT: v_lshrrev_b64 v[10:11], v13, -1 +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v10, v10, v14, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v11, v11, v15, vcc +; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v8, 0, v8, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v9, 0, v9, vcc +; GISEL-NEXT: v_cndmask_b32_e64 v5, v10, -1, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e64 v10, v11, -1, s[4:5] +; GISEL-NEXT: v_and_b32_e32 v2, v8, v2 +; GISEL-NEXT: v_and_b32_e32 v3, v9, v3 +; GISEL-NEXT: v_and_or_b32 v0, v5, v0, v2 +; GISEL-NEXT: v_and_or_b32 v1, v10, v1, v3 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; GISEL-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; GISEL-NEXT: v_or_b32_e32 v3, v12, v0 +; GISEL-NEXT: v_mov_b32_e32 v0, v3 +; GISEL-NEXT: v_mov_b32_e32 v1, v4 +; GISEL-NEXT: v_mov_b32_e32 v2, v5 +; GISEL-NEXT: v_mov_b32_e32 v3, v6 +; GISEL-NEXT: .LBB1_7: ; %Flow1 ; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; GISEL-NEXT: ; %bb.11: ; %itofp-sw-bb -; GISEL-NEXT: v_lshlrev_b64 v[6:7], 1, v[0:1] -; GISEL-NEXT: ; %bb.12: ; %itofp-sw-epilog +; GISEL-NEXT: .LBB1_8: ; %Flow2 +; GISEL-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; GISEL-NEXT: ; %bb.9: ; %itofp-sw-bb +; GISEL-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; GISEL-NEXT: ; %bb.10: ; %itofp-sw-epilog ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: v_bfe_u32 v0, v6, 2, 1 -; GISEL-NEXT: v_or_b32_e32 v0, v6, v0 +; GISEL-NEXT: v_bfe_u32 v2, v0, 2, 1 +; GISEL-NEXT: v_or_b32_e32 v0, v0, v2 ; GISEL-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v7, vcc -; GISEL-NEXT: v_lshrrev_b64 v[2:3], 2, v[0:1] -; GISEL-NEXT: v_and_b32_e32 v3, 0x4000000, v0 -; GISEL-NEXT: v_mov_b32_e32 v4, 0 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[3:4] +; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc +; GISEL-NEXT: v_and_b32_e32 v2, 0x4000000, v0 +; GISEL-NEXT: v_mov_b32_e32 v3, 0 +; GISEL-NEXT: v_lshrrev_b64 v[4:5], 2, v[0:1] +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc -; GISEL-NEXT: ; %bb.13: ; %itofp-if-then20 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], 3, v[0:1] -; GISEL-NEXT: v_mov_b32_e32 v10, v11 -; GISEL-NEXT: ; %bb.14: ; %Flow +; GISEL-NEXT: ; %bb.11: ; %itofp-if-then20 +; GISEL-NEXT: v_lshrrev_b64 v[4:5], 3, v[0:1] +; GISEL-NEXT: v_mov_b32_e32 v6, v7 +; GISEL-NEXT: ; %bb.12: ; %Flow ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: .LBB1_15: ; %Flow7 +; GISEL-NEXT: .LBB1_13: ; %Flow4 ; GISEL-NEXT: s_or_b64 exec, exec, s[8:9] -; GISEL-NEXT: v_lshl_add_u32 v0, v10, 23, 1.0 +; GISEL-NEXT: v_lshl_add_u32 v0, v6, 23, 1.0 ; GISEL-NEXT: v_mov_b32_e32 v1, 0x7fffff -; GISEL-NEXT: v_and_or_b32 v2, v2, v1, v0 -; GISEL-NEXT: .LBB1_16: ; %Flow8 +; GISEL-NEXT: v_and_or_b32 v4, v4, v1, v0 +; GISEL-NEXT: .LBB1_14: ; %Flow5 ; GISEL-NEXT: s_or_b64 exec, exec, s[6:7] -; GISEL-NEXT: v_mov_b32_e32 v0, v2 +; GISEL-NEXT: v_mov_b32_e32 v0, v4 ; GISEL-NEXT: s_setpc_b64 s[30:31] %cvt = uitofp i128 %x to float ret float %cvt @@ -586,7 +522,7 @@ define double @sitofp_i128_to_f64(i128 %x) { ; SDAG-NEXT: v_mov_b32_e32 v0, 0 ; SDAG-NEXT: v_mov_b32_e32 v1, 0 ; SDAG-NEXT: s_and_saveexec_b64 s[6:7], vcc -; SDAG-NEXT: s_cbranch_execz .LBB2_16 +; SDAG-NEXT: s_cbranch_execz .LBB2_14 ; SDAG-NEXT: ; %bb.1: ; %itofp-if-end ; SDAG-NEXT: v_ashrrev_i32_e32 v0, 31, v3 ; SDAG-NEXT: v_xor_b32_e32 v4, v0, v4 @@ -607,128 +543,116 @@ define double @sitofp_i128_to_f64(i128 %x) { ; SDAG-NEXT: v_min_u32_e32 v1, v1, v2 ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[6:7] ; SDAG-NEXT: v_add_u32_e32 v1, 64, v1 -; SDAG-NEXT: v_cndmask_b32_e32 v11, v1, v0, vcc -; SDAG-NEXT: v_sub_u32_e32 v2, 0x7f, v11 -; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 54, v2 -; SDAG-NEXT: ; implicit-def: $vgpr8 +; SDAG-NEXT: v_cndmask_b32_e32 v9, v1, v0, vcc +; SDAG-NEXT: v_sub_u32_e32 v8, 0x80, v9 +; SDAG-NEXT: v_sub_u32_e32 v2, 0x7f, v9 +; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 54, v8 +; SDAG-NEXT: ; implicit-def: $vgpr10 ; SDAG-NEXT: ; implicit-def: $vgpr0_vgpr1 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc ; SDAG-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; SDAG-NEXT: ; %bb.2: ; %itofp-if-else -; SDAG-NEXT: v_add_u32_e32 v6, 0xffffffb5, v11 +; SDAG-NEXT: v_add_u32_e32 v6, 0xffffffb5, v9 ; SDAG-NEXT: v_lshlrev_b64 v[0:1], v6, v[4:5] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v6 -; SDAG-NEXT: v_cndmask_b32_e32 v8, 0, v1, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v10, 0, v1, vcc ; SDAG-NEXT: v_cndmask_b32_e32 v0, 0, v0, vcc +; SDAG-NEXT: ; implicit-def: $vgpr8 ; SDAG-NEXT: ; implicit-def: $vgpr6_vgpr7 ; SDAG-NEXT: ; implicit-def: $vgpr4_vgpr5 -; SDAG-NEXT: ; implicit-def: $vgpr11 -; SDAG-NEXT: ; %bb.3: ; %Flow6 +; SDAG-NEXT: ; implicit-def: $vgpr9 +; SDAG-NEXT: ; %bb.3: ; %Flow3 ; SDAG-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; SDAG-NEXT: s_cbranch_execz .LBB2_15 +; SDAG-NEXT: s_cbranch_execz .LBB2_13 ; SDAG-NEXT: ; %bb.4: ; %NodeBlock -; SDAG-NEXT: v_sub_u32_e32 v10, 0x80, v11 -; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 54, v10 -; SDAG-NEXT: s_mov_b64 s[10:11], 0 -; SDAG-NEXT: s_mov_b64 s[4:5], 0 +; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 54, v8 +; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc +; SDAG-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; SDAG-NEXT: s_cbranch_execz .LBB2_8 +; SDAG-NEXT: ; %bb.5: ; %LeafBlock +; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 55, v8 ; SDAG-NEXT: s_and_saveexec_b64 s[12:13], vcc -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: ; %bb.5: ; %LeafBlock1 -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 55, v10 -; SDAG-NEXT: s_and_b64 s[4:5], vcc, exec -; SDAG-NEXT: ; %bb.6: ; %Flow3 -; SDAG-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; SDAG-NEXT: ; %bb.7: ; %LeafBlock -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 54, v10 -; SDAG-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; SDAG-NEXT: s_and_b64 s[14:15], vcc, exec -; SDAG-NEXT: s_mov_b64 s[10:11], exec -; SDAG-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; SDAG-NEXT: ; %bb.8: ; %Flow4 -; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: v_mov_b32_e32 v0, v4 -; SDAG-NEXT: v_mov_b32_e32 v9, v7 -; SDAG-NEXT: v_mov_b32_e32 v1, v5 -; SDAG-NEXT: v_mov_b32_e32 v8, v6 -; SDAG-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: s_cbranch_execz .LBB2_10 -; SDAG-NEXT: ; %bb.9: ; %itofp-sw-default -; SDAG-NEXT: v_sub_u32_e32 v12, 0x49, v11 -; SDAG-NEXT: v_sub_u32_e32 v8, 64, v12 +; SDAG-NEXT: s_cbranch_execz .LBB2_7 +; SDAG-NEXT: ; %bb.6: ; %itofp-sw-default +; SDAG-NEXT: v_sub_u32_e32 v12, 0x49, v9 +; SDAG-NEXT: v_sub_u32_e32 v10, 64, v12 ; SDAG-NEXT: v_lshrrev_b64 v[0:1], v12, v[4:5] -; SDAG-NEXT: v_lshlrev_b64 v[8:9], v8, v[6:7] -; SDAG-NEXT: v_sub_u32_e32 v13, 9, v11 -; SDAG-NEXT: v_or_b32_e32 v9, v1, v9 -; SDAG-NEXT: v_or_b32_e32 v8, v0, v8 +; SDAG-NEXT: v_lshlrev_b64 v[10:11], v10, v[6:7] +; SDAG-NEXT: v_sub_u32_e32 v13, 9, v9 +; SDAG-NEXT: v_or_b32_e32 v11, v1, v11 +; SDAG-NEXT: v_or_b32_e32 v10, v0, v10 ; SDAG-NEXT: v_lshrrev_b64 v[0:1], v13, v[6:7] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v12 -; SDAG-NEXT: v_cndmask_b32_e32 v1, v1, v9, vcc -; SDAG-NEXT: v_cndmask_b32_e32 v0, v0, v8, vcc -; SDAG-NEXT: v_lshrrev_b64 v[8:9], v12, v[6:7] +; SDAG-NEXT: v_add_u32_e32 v16, 55, v9 +; SDAG-NEXT: v_cndmask_b32_e32 v1, v1, v11, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v12 -; SDAG-NEXT: v_cndmask_b32_e32 v9, 0, v9, vcc -; SDAG-NEXT: v_add_u32_e32 v9, 55, v11 +; SDAG-NEXT: v_cndmask_b32_e32 v0, v0, v10, vcc +; SDAG-NEXT: v_lshrrev_b64 v[10:11], v12, v[6:7] ; SDAG-NEXT: v_lshrrev_b64 v[12:13], v13, v[4:5] -; SDAG-NEXT: v_lshlrev_b64 v[14:15], v9, v[6:7] -; SDAG-NEXT: v_add_u32_e32 v11, -9, v11 +; SDAG-NEXT: v_lshlrev_b64 v[14:15], v16, v[6:7] +; SDAG-NEXT: v_add_u32_e32 v9, -9, v9 +; SDAG-NEXT: v_or_b32_e32 v15, v15, v13 ; SDAG-NEXT: v_or_b32_e32 v14, v14, v12 -; SDAG-NEXT: v_lshlrev_b64 v[11:12], v11, v[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v8, 0, v8, vcc -; SDAG-NEXT: v_or_b32_e32 v13, v15, v13 -; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v9 +; SDAG-NEXT: v_lshlrev_b64 v[12:13], v9, v[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v11, 0, v11, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v10, 0, v10, vcc +; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v16 ; SDAG-NEXT: v_cndmask_b32_e64 v1, v1, v5, s[4:5] ; SDAG-NEXT: v_cndmask_b32_e64 v0, v0, v4, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v12, v12, v13, vcc -; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v9 -; SDAG-NEXT: v_cndmask_b32_e64 v13, v12, v7, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v14, v11, v14, vcc -; SDAG-NEXT: v_lshlrev_b64 v[11:12], v9, v[4:5] -; SDAG-NEXT: v_cndmask_b32_e64 v9, v14, v6, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v12, 0, v12, vcc -; SDAG-NEXT: v_cndmask_b32_e32 v11, 0, v11, vcc -; SDAG-NEXT: v_or_b32_e32 v12, v12, v13 -; SDAG-NEXT: v_or_b32_e32 v11, v11, v9 -; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[11:12] -; SDAG-NEXT: s_andn2_b64 s[10:11], s[10:11], exec -; SDAG-NEXT: v_cndmask_b32_e64 v9, 0, 1, vcc -; SDAG-NEXT: v_or_b32_e32 v0, v0, v9 -; SDAG-NEXT: .LBB2_10: ; %Flow5 +; SDAG-NEXT: v_cndmask_b32_e32 v9, v13, v15, vcc +; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v16 +; SDAG-NEXT: v_lshlrev_b64 v[4:5], v16, v[4:5] +; SDAG-NEXT: v_cndmask_b32_e64 v7, v9, v7, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v9, v12, v14, vcc +; SDAG-NEXT: v_cndmask_b32_e64 v6, v9, v6, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v5, 0, v5, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc +; SDAG-NEXT: v_or_b32_e32 v5, v5, v7 +; SDAG-NEXT: v_or_b32_e32 v4, v4, v6 +; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] +; SDAG-NEXT: v_mov_b32_e32 v6, v10 +; SDAG-NEXT: v_cndmask_b32_e64 v4, 0, 1, vcc +; SDAG-NEXT: v_or_b32_e32 v0, v0, v4 +; SDAG-NEXT: v_mov_b32_e32 v5, v1 +; SDAG-NEXT: v_mov_b32_e32 v4, v0 +; SDAG-NEXT: v_mov_b32_e32 v7, v11 +; SDAG-NEXT: .LBB2_7: ; %Flow1 ; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; SDAG-NEXT: ; %bb.11: ; %itofp-sw-bb -; SDAG-NEXT: v_lshlrev_b64 v[8:9], 1, v[6:7] -; SDAG-NEXT: v_lshrrev_b32_e32 v6, 31, v5 -; SDAG-NEXT: v_lshlrev_b64 v[0:1], 1, v[4:5] -; SDAG-NEXT: v_or_b32_e32 v8, v8, v6 -; SDAG-NEXT: ; %bb.12: ; %itofp-sw-epilog +; SDAG-NEXT: .LBB2_8: ; %Flow2 +; SDAG-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; SDAG-NEXT: ; %bb.9: ; %itofp-sw-bb +; SDAG-NEXT: v_lshlrev_b64 v[6:7], 1, v[6:7] +; SDAG-NEXT: v_lshrrev_b32_e32 v0, 31, v5 +; SDAG-NEXT: v_lshlrev_b64 v[4:5], 1, v[4:5] +; SDAG-NEXT: v_or_b32_e32 v6, v6, v0 +; SDAG-NEXT: ; %bb.10: ; %itofp-sw-epilog ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: v_lshrrev_b32_e32 v4, 2, v0 -; SDAG-NEXT: v_and_or_b32 v0, v4, 1, v0 +; SDAG-NEXT: v_lshrrev_b32_e32 v0, 2, v4 +; SDAG-NEXT: v_and_or_b32 v0, v0, 1, v4 ; SDAG-NEXT: v_add_co_u32_e32 v4, vcc, 1, v0 -; SDAG-NEXT: v_addc_co_u32_e32 v5, vcc, 0, v1, vcc -; SDAG-NEXT: v_addc_co_u32_e32 v6, vcc, 0, v8, vcc +; SDAG-NEXT: v_addc_co_u32_e32 v5, vcc, 0, v5, vcc +; SDAG-NEXT: v_addc_co_u32_e32 v6, vcc, 0, v6, vcc ; SDAG-NEXT: v_lshrrev_b64 v[0:1], 2, v[4:5] ; SDAG-NEXT: v_lshlrev_b32_e32 v7, 30, v6 -; SDAG-NEXT: v_or_b32_e32 v8, v1, v7 +; SDAG-NEXT: v_or_b32_e32 v10, v1, v7 ; SDAG-NEXT: v_and_b32_e32 v1, 0x800000, v5 ; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 0, v1 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc -; SDAG-NEXT: ; %bb.13: ; %itofp-if-then20 +; SDAG-NEXT: ; %bb.11: ; %itofp-if-then20 ; SDAG-NEXT: v_lshrrev_b64 v[0:1], 3, v[4:5] ; SDAG-NEXT: v_lshlrev_b32_e32 v2, 29, v6 -; SDAG-NEXT: v_or_b32_e32 v8, v1, v2 -; SDAG-NEXT: v_mov_b32_e32 v2, v10 -; SDAG-NEXT: ; %bb.14: ; %Flow +; SDAG-NEXT: v_or_b32_e32 v10, v1, v2 +; SDAG-NEXT: v_mov_b32_e32 v2, v8 +; SDAG-NEXT: ; %bb.12: ; %Flow ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: .LBB2_15: ; %Flow7 +; SDAG-NEXT: .LBB2_13: ; %Flow4 ; SDAG-NEXT: s_or_b64 exec, exec, s[8:9] ; SDAG-NEXT: v_and_b32_e32 v1, 0x80000000, v3 ; SDAG-NEXT: v_mov_b32_e32 v3, 0x3ff00000 ; SDAG-NEXT: v_lshl_add_u32 v2, v2, 20, v3 -; SDAG-NEXT: v_and_b32_e32 v3, 0xfffff, v8 +; SDAG-NEXT: v_and_b32_e32 v3, 0xfffff, v10 ; SDAG-NEXT: v_or3_b32 v1, v3, v1, v2 -; SDAG-NEXT: .LBB2_16: ; %Flow8 +; SDAG-NEXT: .LBB2_14: ; %Flow5 ; SDAG-NEXT: s_or_b64 exec, exec, s[6:7] ; SDAG-NEXT: s_setpc_b64 s[30:31] ; @@ -744,156 +668,142 @@ define double @sitofp_i128_to_f64(i128 %x) { ; GISEL-NEXT: v_mov_b32_e32 v0, s4 ; GISEL-NEXT: v_mov_b32_e32 v1, s5 ; GISEL-NEXT: s_and_saveexec_b64 s[6:7], vcc -; GISEL-NEXT: s_cbranch_execz .LBB2_16 +; GISEL-NEXT: s_cbranch_execz .LBB2_14 ; GISEL-NEXT: ; %bb.1: ; %itofp-if-end -; GISEL-NEXT: v_ashrrev_i32_e32 v8, 31, v3 -; GISEL-NEXT: v_xor_b32_e32 v0, v8, v4 -; GISEL-NEXT: v_xor_b32_e32 v1, v8, v5 -; GISEL-NEXT: v_sub_co_u32_e32 v0, vcc, v0, v8 -; GISEL-NEXT: v_xor_b32_e32 v2, v8, v2 -; GISEL-NEXT: v_subb_co_u32_e32 v1, vcc, v1, v8, vcc -; GISEL-NEXT: v_xor_b32_e32 v3, v8, v3 -; GISEL-NEXT: v_subb_co_u32_e32 v2, vcc, v2, v8, vcc +; GISEL-NEXT: v_ashrrev_i32_e32 v6, 31, v3 +; GISEL-NEXT: v_xor_b32_e32 v0, v6, v4 +; GISEL-NEXT: v_xor_b32_e32 v1, v6, v5 +; GISEL-NEXT: v_sub_co_u32_e32 v0, vcc, v0, v6 +; GISEL-NEXT: v_xor_b32_e32 v2, v6, v2 +; GISEL-NEXT: v_subb_co_u32_e32 v1, vcc, v1, v6, vcc +; GISEL-NEXT: v_xor_b32_e32 v3, v6, v3 +; GISEL-NEXT: v_subb_co_u32_e32 v2, vcc, v2, v6, vcc ; GISEL-NEXT: v_ffbh_u32_e32 v5, v0 -; GISEL-NEXT: v_subb_co_u32_e32 v3, vcc, v3, v8, vcc +; GISEL-NEXT: v_subb_co_u32_e32 v3, vcc, v3, v6, vcc ; GISEL-NEXT: v_ffbh_u32_e32 v4, v1 ; GISEL-NEXT: v_add_u32_e32 v5, 32, v5 -; GISEL-NEXT: v_ffbh_u32_e32 v6, v2 +; GISEL-NEXT: v_ffbh_u32_e32 v7, v2 ; GISEL-NEXT: v_min_u32_e32 v4, v4, v5 ; GISEL-NEXT: v_ffbh_u32_e32 v5, v3 -; GISEL-NEXT: v_add_u32_e32 v6, 32, v6 +; GISEL-NEXT: v_add_u32_e32 v7, 32, v7 ; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[2:3] ; GISEL-NEXT: v_add_u32_e32 v4, 64, v4 -; GISEL-NEXT: v_min_u32_e32 v5, v5, v6 -; GISEL-NEXT: v_cndmask_b32_e32 v11, v5, v4, vcc -; GISEL-NEXT: v_sub_u32_e32 v9, 0x7f, v11 -; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 53, v9 -; GISEL-NEXT: ; implicit-def: $vgpr6 +; GISEL-NEXT: v_min_u32_e32 v5, v5, v7 +; GISEL-NEXT: v_cndmask_b32_e32 v9, v5, v4, vcc +; GISEL-NEXT: v_sub_u32_e32 v8, 0x80, v9 +; GISEL-NEXT: v_sub_u32_e32 v7, 0x7f, v9 +; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 53, v8 +; GISEL-NEXT: ; implicit-def: $vgpr10 ; GISEL-NEXT: ; implicit-def: $vgpr4_vgpr5 ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc ; GISEL-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; GISEL-NEXT: ; %bb.2: ; %itofp-if-else -; GISEL-NEXT: v_add_u32_e32 v2, 0xffffffb5, v11 +; GISEL-NEXT: v_add_u32_e32 v2, 0xffffffb5, v9 ; GISEL-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 ; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v0, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v6, 0, v1, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v10, 0, v1, vcc +; GISEL-NEXT: ; implicit-def: $vgpr8 ; GISEL-NEXT: ; implicit-def: $vgpr0 -; GISEL-NEXT: ; implicit-def: $vgpr11 -; GISEL-NEXT: ; %bb.3: ; %Flow6 +; GISEL-NEXT: ; implicit-def: $vgpr9 +; GISEL-NEXT: ; %bb.3: ; %Flow3 ; GISEL-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; GISEL-NEXT: s_cbranch_execz .LBB2_15 +; GISEL-NEXT: s_cbranch_execz .LBB2_13 ; GISEL-NEXT: ; %bb.4: ; %NodeBlock -; GISEL-NEXT: v_sub_u32_e32 v10, 0x80, v11 -; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 55, v10 -; GISEL-NEXT: s_mov_b64 s[10:11], 0 -; GISEL-NEXT: s_mov_b64 s[4:5], 0 +; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 55, v8 +; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc +; GISEL-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; GISEL-NEXT: s_cbranch_execz .LBB2_8 +; GISEL-NEXT: ; %bb.5: ; %LeafBlock +; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 55, v8 ; GISEL-NEXT: s_and_saveexec_b64 s[12:13], vcc -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: ; %bb.5: ; %LeafBlock1 -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 55, v10 -; GISEL-NEXT: s_andn2_b64 s[4:5], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.6: ; %Flow3 -; GISEL-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; GISEL-NEXT: ; %bb.7: ; %LeafBlock -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 54, v10 -; GISEL-NEXT: s_andn2_b64 s[10:11], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, -1 -; GISEL-NEXT: s_or_b64 s[10:11], s[10:11], s[14:15] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.8: ; %Flow4 -; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: v_mov_b32_e32 v7, v3 -; GISEL-NEXT: v_mov_b32_e32 v6, v2 -; GISEL-NEXT: v_mov_b32_e32 v5, v1 -; GISEL-NEXT: v_mov_b32_e32 v4, v0 -; GISEL-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: s_cbranch_execz .LBB2_10 -; GISEL-NEXT: ; %bb.9: ; %itofp-sw-default -; GISEL-NEXT: v_sub_u32_e32 v14, 0x49, v11 -; GISEL-NEXT: v_sub_u32_e32 v6, 64, v14 +; GISEL-NEXT: s_cbranch_execz .LBB2_7 +; GISEL-NEXT: ; %bb.6: ; %itofp-sw-default +; GISEL-NEXT: v_sub_u32_e32 v14, 0x49, v9 +; GISEL-NEXT: v_sub_u32_e32 v10, 64, v14 ; GISEL-NEXT: v_lshrrev_b64 v[4:5], v14, v[0:1] -; GISEL-NEXT: v_lshlrev_b64 v[6:7], v6, v[2:3] +; GISEL-NEXT: v_lshlrev_b64 v[10:11], v10, v[2:3] ; GISEL-NEXT: v_subrev_u32_e32 v15, 64, v14 -; GISEL-NEXT: v_or_b32_e32 v6, v4, v6 -; GISEL-NEXT: v_or_b32_e32 v7, v5, v7 +; GISEL-NEXT: v_or_b32_e32 v10, v4, v10 +; GISEL-NEXT: v_or_b32_e32 v11, v5, v11 ; GISEL-NEXT: v_lshrrev_b64 v[4:5], v15, v[2:3] -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v14 ; GISEL-NEXT: v_lshrrev_b64 v[12:13], v14, v[2:3] -; GISEL-NEXT: v_cndmask_b32_e32 v5, v5, v7, vcc -; GISEL-NEXT: v_add_u32_e32 v7, 55, v11 -; GISEL-NEXT: v_sub_u32_e32 v13, 64, v7 -; GISEL-NEXT: v_cndmask_b32_e32 v4, v4, v6, vcc +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v14 ; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v14 -; GISEL-NEXT: v_cndmask_b32_e32 v6, 0, v12, vcc -; GISEL-NEXT: v_lshrrev_b64 v[11:12], v7, -1 -; GISEL-NEXT: v_lshlrev_b64 v[13:14], v13, -1 -; GISEL-NEXT: v_subrev_u32_e32 v15, 64, v7 -; GISEL-NEXT: v_or_b32_e32 v16, v11, v13 -; GISEL-NEXT: v_or_b32_e32 v17, v12, v14 -; GISEL-NEXT: v_lshrrev_b64 v[13:14], v15, -1 -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v7 -; GISEL-NEXT: v_cndmask_b32_e64 v4, v4, v0, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e64 v5, v5, v1, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e32 v13, v13, v16, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v14, v14, v17, vcc -; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v7 -; GISEL-NEXT: v_cndmask_b32_e32 v11, 0, v11, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v12, 0, v12, vcc -; GISEL-NEXT: v_cndmask_b32_e64 v7, v13, -1, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e64 v13, v14, -1, s[4:5] -; GISEL-NEXT: v_and_b32_e32 v11, v11, v2 -; GISEL-NEXT: v_and_b32_e32 v12, v12, v3 -; GISEL-NEXT: v_and_or_b32 v11, v7, v0, v11 -; GISEL-NEXT: v_and_or_b32 v12, v13, v1, v12 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[11:12] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[10:11], exec -; GISEL-NEXT: v_cndmask_b32_e64 v7, 0, 1, vcc -; GISEL-NEXT: s_and_b64 s[10:11], exec, 0 -; GISEL-NEXT: v_or_b32_e32 v4, v4, v7 -; GISEL-NEXT: s_or_b64 s[10:11], s[4:5], s[10:11] -; GISEL-NEXT: .LBB2_10: ; %Flow5 +; GISEL-NEXT: v_add_u32_e32 v14, 55, v9 +; GISEL-NEXT: v_cndmask_b32_e32 v4, v4, v10, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v5, v5, v11, vcc +; GISEL-NEXT: v_sub_u32_e32 v11, 64, v14 +; GISEL-NEXT: v_cndmask_b32_e64 v13, v4, v0, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e64 v4, v5, v1, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e32 v5, 0, v12, vcc +; GISEL-NEXT: v_lshrrev_b64 v[9:10], v14, -1 +; GISEL-NEXT: v_lshlrev_b64 v[11:12], v11, -1 +; GISEL-NEXT: v_subrev_u32_e32 v15, 64, v14 +; GISEL-NEXT: v_or_b32_e32 v16, v9, v11 +; GISEL-NEXT: v_or_b32_e32 v17, v10, v12 +; GISEL-NEXT: v_lshrrev_b64 v[11:12], v15, -1 +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v14 +; GISEL-NEXT: v_cndmask_b32_e32 v11, v11, v16, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v12, v12, v17, vcc +; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v14 +; GISEL-NEXT: v_cndmask_b32_e32 v9, 0, v9, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v10, 0, v10, vcc +; GISEL-NEXT: v_cndmask_b32_e64 v11, v11, -1, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e64 v12, v12, -1, s[4:5] +; GISEL-NEXT: v_and_b32_e32 v2, v9, v2 +; GISEL-NEXT: v_and_b32_e32 v3, v10, v3 +; GISEL-NEXT: v_and_or_b32 v0, v11, v0, v2 +; GISEL-NEXT: v_and_or_b32 v1, v12, v1, v3 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; GISEL-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; GISEL-NEXT: v_or_b32_e32 v3, v13, v0 +; GISEL-NEXT: v_mov_b32_e32 v0, v3 +; GISEL-NEXT: v_mov_b32_e32 v1, v4 +; GISEL-NEXT: v_mov_b32_e32 v2, v5 +; GISEL-NEXT: v_mov_b32_e32 v3, v6 +; GISEL-NEXT: .LBB2_7: ; %Flow1 ; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; GISEL-NEXT: ; %bb.11: ; %itofp-sw-bb +; GISEL-NEXT: .LBB2_8: ; %Flow2 +; GISEL-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; GISEL-NEXT: ; %bb.9: ; %itofp-sw-bb +; GISEL-NEXT: v_lshlrev_b64 v[9:10], 1, v[0:1] ; GISEL-NEXT: v_lshlrev_b64 v[2:3], 1, v[2:3] -; GISEL-NEXT: v_lshlrev_b64 v[4:5], 1, v[0:1] ; GISEL-NEXT: v_lshrrev_b32_e32 v0, 31, v1 -; GISEL-NEXT: v_or_b32_e32 v6, v2, v0 -; GISEL-NEXT: ; %bb.12: ; %itofp-sw-epilog +; GISEL-NEXT: v_or_b32_e32 v11, v2, v0 +; GISEL-NEXT: v_mov_b32_e32 v0, v9 +; GISEL-NEXT: v_mov_b32_e32 v1, v10 +; GISEL-NEXT: v_mov_b32_e32 v2, v11 +; GISEL-NEXT: v_mov_b32_e32 v3, v12 +; GISEL-NEXT: ; %bb.10: ; %itofp-sw-epilog ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: v_bfe_u32 v0, v4, 2, 1 -; GISEL-NEXT: v_or_b32_e32 v0, v4, v0 +; GISEL-NEXT: v_bfe_u32 v3, v0, 2, 1 +; GISEL-NEXT: v_or_b32_e32 v0, v0, v3 ; GISEL-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v5, vcc -; GISEL-NEXT: v_addc_co_u32_e32 v2, vcc, 0, v6, vcc +; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc +; GISEL-NEXT: v_addc_co_u32_e32 v2, vcc, 0, v2, vcc ; GISEL-NEXT: v_lshrrev_b64 v[4:5], 2, v[0:1] -; GISEL-NEXT: v_mov_b32_e32 v6, 0 -; GISEL-NEXT: v_and_b32_e32 v7, 0x800000, v1 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[6:7] -; GISEL-NEXT: v_lshl_or_b32 v6, v2, 30, v5 +; GISEL-NEXT: v_mov_b32_e32 v9, 0 +; GISEL-NEXT: v_and_b32_e32 v10, 0x800000, v1 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[9:10] +; GISEL-NEXT: v_lshl_or_b32 v10, v2, 30, v5 ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc -; GISEL-NEXT: ; %bb.13: ; %itofp-if-then20 +; GISEL-NEXT: ; %bb.11: ; %itofp-if-then20 ; GISEL-NEXT: v_lshrrev_b64 v[4:5], 3, v[0:1] -; GISEL-NEXT: v_mov_b32_e32 v9, v10 -; GISEL-NEXT: v_lshl_or_b32 v6, v2, 29, v5 -; GISEL-NEXT: ; %bb.14: ; %Flow +; GISEL-NEXT: v_mov_b32_e32 v7, v8 +; GISEL-NEXT: v_lshl_or_b32 v10, v2, 29, v5 +; GISEL-NEXT: ; %bb.12: ; %Flow ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: .LBB2_15: ; %Flow7 +; GISEL-NEXT: .LBB2_13: ; %Flow4 ; GISEL-NEXT: s_or_b64 exec, exec, s[8:9] -; GISEL-NEXT: v_and_b32_e32 v0, 0x80000000, v8 +; GISEL-NEXT: v_and_b32_e32 v0, 0x80000000, v6 ; GISEL-NEXT: v_mov_b32_e32 v1, 0x3ff00000 ; GISEL-NEXT: v_mov_b32_e32 v2, 0xfffff -; GISEL-NEXT: v_lshl_add_u32 v1, v9, 20, v1 -; GISEL-NEXT: v_and_or_b32 v2, v6, v2, v0 +; GISEL-NEXT: v_lshl_add_u32 v1, v7, 20, v1 +; GISEL-NEXT: v_and_or_b32 v2, v10, v2, v0 ; GISEL-NEXT: v_and_or_b32 v0, v4, -1, 0 ; GISEL-NEXT: v_or3_b32 v1, v2, v1, 0 -; GISEL-NEXT: .LBB2_16: ; %Flow8 +; GISEL-NEXT: .LBB2_14: ; %Flow5 ; GISEL-NEXT: s_or_b64 exec, exec, s[6:7] ; GISEL-NEXT: s_setpc_b64 s[30:31] %cvt = sitofp i128 %x to double @@ -910,7 +820,7 @@ define double @uitofp_i128_to_f64(i128 %x) { ; SDAG-NEXT: v_mov_b32_e32 v4, 0 ; SDAG-NEXT: v_mov_b32_e32 v5, 0 ; SDAG-NEXT: s_and_saveexec_b64 s[6:7], vcc -; SDAG-NEXT: s_cbranch_execz .LBB3_16 +; SDAG-NEXT: s_cbranch_execz .LBB3_14 ; SDAG-NEXT: ; %bb.1: ; %itofp-if-end ; SDAG-NEXT: v_ffbh_u32_e32 v4, v2 ; SDAG-NEXT: v_add_u32_e32 v4, 32, v4 @@ -922,128 +832,112 @@ define double @uitofp_i128_to_f64(i128 %x) { ; SDAG-NEXT: v_min_u32_e32 v5, v5, v6 ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] ; SDAG-NEXT: v_add_u32_e32 v5, 64, v5 -; SDAG-NEXT: v_cndmask_b32_e32 v11, v5, v4, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v8, v5, v4, vcc +; SDAG-NEXT: v_sub_u32_e32 v7, 0x80, v8 +; SDAG-NEXT: v_sub_u32_e32 v6, 0x7f, v8 +; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 54, v7 +; SDAG-NEXT: ; implicit-def: $vgpr9 ; SDAG-NEXT: ; implicit-def: $vgpr4_vgpr5 -; SDAG-NEXT: v_sub_u32_e32 v9, 0x7f, v11 -; SDAG-NEXT: v_mov_b32_e32 v6, v1 -; SDAG-NEXT: v_mov_b32_e32 v8, v3 -; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 54, v9 -; SDAG-NEXT: v_mov_b32_e32 v5, v0 -; SDAG-NEXT: v_mov_b32_e32 v7, v2 -; SDAG-NEXT: ; implicit-def: $vgpr12 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc ; SDAG-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; SDAG-NEXT: ; %bb.2: ; %itofp-if-else -; SDAG-NEXT: v_add_u32_e32 v2, 0xffffffb5, v11 +; SDAG-NEXT: v_add_u32_e32 v2, 0xffffffb5, v8 ; SDAG-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 -; SDAG-NEXT: v_cndmask_b32_e32 v12, 0, v1, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v9, 0, v1, vcc ; SDAG-NEXT: v_cndmask_b32_e32 v4, 0, v0, vcc +; SDAG-NEXT: ; implicit-def: $vgpr7 ; SDAG-NEXT: ; implicit-def: $vgpr2_vgpr3 ; SDAG-NEXT: ; implicit-def: $vgpr0_vgpr1 -; SDAG-NEXT: ; implicit-def: $vgpr11 -; SDAG-NEXT: ; implicit-def: $vgpr5_vgpr6 -; SDAG-NEXT: ; implicit-def: $vgpr7_vgpr8 -; SDAG-NEXT: ; %bb.3: ; %Flow6 +; SDAG-NEXT: ; implicit-def: $vgpr8 +; SDAG-NEXT: ; %bb.3: ; %Flow3 ; SDAG-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; SDAG-NEXT: s_cbranch_execz .LBB3_15 +; SDAG-NEXT: s_cbranch_execz .LBB3_13 ; SDAG-NEXT: ; %bb.4: ; %NodeBlock -; SDAG-NEXT: v_sub_u32_e32 v10, 0x80, v11 -; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 54, v10 -; SDAG-NEXT: s_mov_b64 s[10:11], 0 -; SDAG-NEXT: s_mov_b64 s[4:5], 0 +; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 54, v7 +; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc +; SDAG-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; SDAG-NEXT: s_cbranch_execz .LBB3_8 +; SDAG-NEXT: ; %bb.5: ; %LeafBlock +; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 55, v7 ; SDAG-NEXT: s_and_saveexec_b64 s[12:13], vcc -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: ; %bb.5: ; %LeafBlock1 -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 55, v10 -; SDAG-NEXT: s_and_b64 s[4:5], vcc, exec -; SDAG-NEXT: ; %bb.6: ; %Flow3 -; SDAG-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; SDAG-NEXT: ; %bb.7: ; %LeafBlock -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 54, v10 -; SDAG-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; SDAG-NEXT: s_and_b64 s[14:15], vcc, exec -; SDAG-NEXT: s_mov_b64 s[10:11], exec -; SDAG-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; SDAG-NEXT: ; implicit-def: $vgpr5_vgpr6 -; SDAG-NEXT: ; implicit-def: $vgpr7_vgpr8 -; SDAG-NEXT: ; %bb.8: ; %Flow4 -; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: s_cbranch_execz .LBB3_10 -; SDAG-NEXT: ; %bb.9: ; %itofp-sw-default -; SDAG-NEXT: v_sub_u32_e32 v8, 0x49, v11 -; SDAG-NEXT: v_sub_u32_e32 v6, 64, v8 -; SDAG-NEXT: v_lshrrev_b64 v[4:5], v8, v[0:1] -; SDAG-NEXT: v_lshlrev_b64 v[6:7], v6, v[2:3] -; SDAG-NEXT: v_sub_u32_e32 v13, 9, v11 -; SDAG-NEXT: v_or_b32_e32 v7, v5, v7 -; SDAG-NEXT: v_or_b32_e32 v12, v4, v6 -; SDAG-NEXT: v_lshrrev_b64 v[4:5], v13, v[2:3] -; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v8 -; SDAG-NEXT: v_cndmask_b32_e32 v5, v5, v7, vcc -; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v8 -; SDAG-NEXT: v_cndmask_b32_e64 v6, v5, v1, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v7, v4, v12, vcc -; SDAG-NEXT: v_lshrrev_b64 v[4:5], v8, v[2:3] -; SDAG-NEXT: v_add_u32_e32 v16, 55, v11 -; SDAG-NEXT: v_cndmask_b32_e64 v8, v7, v0, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v5, 0, v5, vcc -; SDAG-NEXT: v_lshrrev_b64 v[12:13], v13, v[0:1] -; SDAG-NEXT: v_lshlrev_b64 v[14:15], v16, v[2:3] -; SDAG-NEXT: v_cndmask_b32_e32 v7, 0, v4, vcc -; SDAG-NEXT: v_add_u32_e32 v4, -9, v11 -; SDAG-NEXT: v_lshlrev_b64 v[4:5], v4, v[0:1] -; SDAG-NEXT: v_or_b32_e32 v13, v15, v13 -; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v16 -; SDAG-NEXT: v_or_b32_e32 v12, v14, v12 -; SDAG-NEXT: v_cndmask_b32_e32 v5, v5, v13, vcc -; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v16 -; SDAG-NEXT: v_cndmask_b32_e64 v11, v5, v3, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v12, v4, v12, vcc -; SDAG-NEXT: v_lshlrev_b64 v[4:5], v16, v[0:1] -; SDAG-NEXT: v_cndmask_b32_e64 v12, v12, v2, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v5, 0, v5, vcc -; SDAG-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc -; SDAG-NEXT: v_or_b32_e32 v5, v5, v11 -; SDAG-NEXT: v_or_b32_e32 v4, v4, v12 -; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] -; SDAG-NEXT: s_andn2_b64 s[10:11], s[10:11], exec -; SDAG-NEXT: v_cndmask_b32_e64 v4, 0, 1, vcc -; SDAG-NEXT: v_or_b32_e32 v5, v8, v4 -; SDAG-NEXT: .LBB3_10: ; %Flow5 +; SDAG-NEXT: s_cbranch_execz .LBB3_7 +; SDAG-NEXT: ; %bb.6: ; %itofp-sw-default +; SDAG-NEXT: v_sub_u32_e32 v11, 0x49, v8 +; SDAG-NEXT: v_sub_u32_e32 v9, 64, v11 +; SDAG-NEXT: v_lshrrev_b64 v[4:5], v11, v[0:1] +; SDAG-NEXT: v_lshlrev_b64 v[9:10], v9, v[2:3] +; SDAG-NEXT: v_sub_u32_e32 v12, 9, v8 +; SDAG-NEXT: v_or_b32_e32 v10, v5, v10 +; SDAG-NEXT: v_or_b32_e32 v9, v4, v9 +; SDAG-NEXT: v_lshrrev_b64 v[4:5], v12, v[2:3] +; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v11 +; SDAG-NEXT: v_add_u32_e32 v15, 55, v8 +; SDAG-NEXT: v_cndmask_b32_e32 v5, v5, v10, vcc +; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v11 +; SDAG-NEXT: v_cndmask_b32_e32 v4, v4, v9, vcc +; SDAG-NEXT: v_lshrrev_b64 v[9:10], v11, v[2:3] +; SDAG-NEXT: v_lshrrev_b64 v[11:12], v12, v[0:1] +; SDAG-NEXT: v_lshlrev_b64 v[13:14], v15, v[2:3] +; SDAG-NEXT: v_add_u32_e32 v8, -9, v8 +; SDAG-NEXT: v_or_b32_e32 v14, v14, v12 +; SDAG-NEXT: v_or_b32_e32 v13, v13, v11 +; SDAG-NEXT: v_lshlrev_b64 v[11:12], v8, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e32 v10, 0, v10, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v9, 0, v9, vcc +; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v15 +; SDAG-NEXT: v_cndmask_b32_e64 v5, v5, v1, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e64 v4, v4, v0, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v8, v12, v14, vcc +; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v15 +; SDAG-NEXT: v_lshlrev_b64 v[0:1], v15, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v3, v8, v3, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v8, v11, v13, vcc +; SDAG-NEXT: v_cndmask_b32_e64 v2, v8, v2, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v0, 0, v0, vcc +; SDAG-NEXT: v_or_b32_e32 v1, v1, v3 +; SDAG-NEXT: v_or_b32_e32 v0, v0, v2 +; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; SDAG-NEXT: v_mov_b32_e32 v2, v9 +; SDAG-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; SDAG-NEXT: v_or_b32_e32 v4, v4, v0 +; SDAG-NEXT: v_mov_b32_e32 v0, v4 +; SDAG-NEXT: v_mov_b32_e32 v1, v5 +; SDAG-NEXT: v_mov_b32_e32 v3, v10 +; SDAG-NEXT: .LBB3_7: ; %Flow1 ; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; SDAG-NEXT: ; %bb.11: ; %itofp-sw-bb -; SDAG-NEXT: v_lshlrev_b64 v[7:8], 1, v[2:3] -; SDAG-NEXT: v_lshrrev_b32_e32 v2, 31, v1 -; SDAG-NEXT: v_lshlrev_b64 v[5:6], 1, v[0:1] -; SDAG-NEXT: v_or_b32_e32 v7, v7, v2 -; SDAG-NEXT: ; %bb.12: ; %itofp-sw-epilog +; SDAG-NEXT: .LBB3_8: ; %Flow2 +; SDAG-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; SDAG-NEXT: ; %bb.9: ; %itofp-sw-bb +; SDAG-NEXT: v_lshlrev_b64 v[2:3], 1, v[2:3] +; SDAG-NEXT: v_lshrrev_b32_e32 v3, 31, v1 +; SDAG-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; SDAG-NEXT: v_or_b32_e32 v2, v2, v3 +; SDAG-NEXT: ; %bb.10: ; %itofp-sw-epilog ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: v_lshrrev_b32_e32 v0, 2, v5 -; SDAG-NEXT: v_and_or_b32 v0, v0, 1, v5 +; SDAG-NEXT: v_lshrrev_b32_e32 v3, 2, v0 +; SDAG-NEXT: v_and_or_b32 v0, v3, 1, v0 ; SDAG-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v6, vcc -; SDAG-NEXT: v_addc_co_u32_e32 v2, vcc, 0, v7, vcc +; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc +; SDAG-NEXT: v_addc_co_u32_e32 v2, vcc, 0, v2, vcc ; SDAG-NEXT: v_lshrrev_b64 v[4:5], 2, v[0:1] ; SDAG-NEXT: v_and_b32_e32 v3, 0x800000, v1 ; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 0, v3 -; SDAG-NEXT: v_alignbit_b32 v12, v2, v1, 2 +; SDAG-NEXT: v_alignbit_b32 v9, v2, v1, 2 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc -; SDAG-NEXT: ; %bb.13: ; %itofp-if-then20 +; SDAG-NEXT: ; %bb.11: ; %itofp-if-then20 ; SDAG-NEXT: v_lshrrev_b64 v[4:5], 3, v[0:1] -; SDAG-NEXT: v_alignbit_b32 v12, v2, v1, 3 -; SDAG-NEXT: v_mov_b32_e32 v9, v10 -; SDAG-NEXT: ; %bb.14: ; %Flow +; SDAG-NEXT: v_alignbit_b32 v9, v2, v1, 3 +; SDAG-NEXT: v_mov_b32_e32 v6, v7 +; SDAG-NEXT: ; %bb.12: ; %Flow ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: .LBB3_15: ; %Flow7 +; SDAG-NEXT: .LBB3_13: ; %Flow4 ; SDAG-NEXT: s_or_b64 exec, exec, s[8:9] -; SDAG-NEXT: v_and_b32_e32 v0, 0xfffff, v12 -; SDAG-NEXT: v_lshl_or_b32 v0, v9, 20, v0 +; SDAG-NEXT: v_and_b32_e32 v0, 0xfffff, v9 +; SDAG-NEXT: v_lshl_or_b32 v0, v6, 20, v0 ; SDAG-NEXT: v_add_u32_e32 v5, 0x3ff00000, v0 -; SDAG-NEXT: .LBB3_16: ; %Flow8 +; SDAG-NEXT: .LBB3_14: ; %Flow5 ; SDAG-NEXT: s_or_b64 exec, exec, s[6:7] ; SDAG-NEXT: v_mov_b32_e32 v0, v4 ; SDAG-NEXT: v_mov_b32_e32 v1, v5 @@ -1059,7 +953,7 @@ define double @uitofp_i128_to_f64(i128 %x) { ; GISEL-NEXT: v_mov_b32_e32 v4, s4 ; GISEL-NEXT: v_mov_b32_e32 v5, s5 ; GISEL-NEXT: s_and_saveexec_b64 s[6:7], vcc -; GISEL-NEXT: s_cbranch_execz .LBB3_16 +; GISEL-NEXT: s_cbranch_execz .LBB3_14 ; GISEL-NEXT: ; %bb.1: ; %itofp-if-end ; GISEL-NEXT: v_ffbh_u32_e32 v5, v0 ; GISEL-NEXT: v_ffbh_u32_e32 v4, v1 @@ -1071,139 +965,125 @@ define double @uitofp_i128_to_f64(i128 %x) { ; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[2:3] ; GISEL-NEXT: v_add_u32_e32 v4, 64, v4 ; GISEL-NEXT: v_min_u32_e32 v5, v5, v6 -; GISEL-NEXT: v_cndmask_b32_e32 v10, v5, v4, vcc -; GISEL-NEXT: v_sub_u32_e32 v8, 0x7f, v10 -; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 53, v8 -; GISEL-NEXT: ; implicit-def: $vgpr6 +; GISEL-NEXT: v_cndmask_b32_e32 v8, v5, v4, vcc +; GISEL-NEXT: v_sub_u32_e32 v7, 0x80, v8 +; GISEL-NEXT: v_sub_u32_e32 v6, 0x7f, v8 +; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 53, v7 +; GISEL-NEXT: ; implicit-def: $vgpr9 ; GISEL-NEXT: ; implicit-def: $vgpr4_vgpr5 ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc ; GISEL-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; GISEL-NEXT: ; %bb.2: ; %itofp-if-else -; GISEL-NEXT: v_add_u32_e32 v2, 0xffffffb5, v10 +; GISEL-NEXT: v_add_u32_e32 v2, 0xffffffb5, v8 ; GISEL-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 ; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v0, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v6, 0, v1, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v9, 0, v1, vcc +; GISEL-NEXT: ; implicit-def: $vgpr7 ; GISEL-NEXT: ; implicit-def: $vgpr0 -; GISEL-NEXT: ; implicit-def: $vgpr10 -; GISEL-NEXT: ; %bb.3: ; %Flow6 +; GISEL-NEXT: ; implicit-def: $vgpr8 +; GISEL-NEXT: ; %bb.3: ; %Flow3 ; GISEL-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; GISEL-NEXT: s_cbranch_execz .LBB3_15 +; GISEL-NEXT: s_cbranch_execz .LBB3_13 ; GISEL-NEXT: ; %bb.4: ; %NodeBlock -; GISEL-NEXT: v_sub_u32_e32 v9, 0x80, v10 -; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 55, v9 -; GISEL-NEXT: s_mov_b64 s[10:11], 0 -; GISEL-NEXT: s_mov_b64 s[4:5], 0 +; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 55, v7 +; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc +; GISEL-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; GISEL-NEXT: s_cbranch_execz .LBB3_8 +; GISEL-NEXT: ; %bb.5: ; %LeafBlock +; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 55, v7 ; GISEL-NEXT: s_and_saveexec_b64 s[12:13], vcc -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: ; %bb.5: ; %LeafBlock1 -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 55, v9 -; GISEL-NEXT: s_andn2_b64 s[4:5], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.6: ; %Flow3 -; GISEL-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; GISEL-NEXT: ; %bb.7: ; %LeafBlock -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 54, v9 -; GISEL-NEXT: s_andn2_b64 s[10:11], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, -1 -; GISEL-NEXT: s_or_b64 s[10:11], s[10:11], s[14:15] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.8: ; %Flow4 -; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: v_mov_b32_e32 v7, v3 -; GISEL-NEXT: v_mov_b32_e32 v6, v2 -; GISEL-NEXT: v_mov_b32_e32 v5, v1 -; GISEL-NEXT: v_mov_b32_e32 v4, v0 -; GISEL-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: s_cbranch_execz .LBB3_10 -; GISEL-NEXT: ; %bb.9: ; %itofp-sw-default -; GISEL-NEXT: v_sub_u32_e32 v13, 0x49, v10 -; GISEL-NEXT: v_sub_u32_e32 v6, 64, v13 +; GISEL-NEXT: s_cbranch_execz .LBB3_7 +; GISEL-NEXT: ; %bb.6: ; %itofp-sw-default +; GISEL-NEXT: v_sub_u32_e32 v13, 0x49, v8 +; GISEL-NEXT: v_sub_u32_e32 v9, 64, v13 ; GISEL-NEXT: v_lshrrev_b64 v[4:5], v13, v[0:1] -; GISEL-NEXT: v_lshlrev_b64 v[6:7], v6, v[2:3] +; GISEL-NEXT: v_lshlrev_b64 v[9:10], v9, v[2:3] ; GISEL-NEXT: v_subrev_u32_e32 v14, 64, v13 ; GISEL-NEXT: v_lshrrev_b64 v[11:12], v13, v[2:3] -; GISEL-NEXT: v_or_b32_e32 v6, v4, v6 -; GISEL-NEXT: v_or_b32_e32 v7, v5, v7 +; GISEL-NEXT: v_or_b32_e32 v9, v4, v9 +; GISEL-NEXT: v_or_b32_e32 v10, v5, v10 ; GISEL-NEXT: v_lshrrev_b64 v[4:5], v14, v[2:3] ; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v13 -; GISEL-NEXT: v_add_u32_e32 v14, 55, v10 -; GISEL-NEXT: v_cndmask_b32_e32 v5, v5, v7, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v7, 0, v12, vcc -; GISEL-NEXT: v_sub_u32_e32 v12, 64, v14 -; GISEL-NEXT: v_cndmask_b32_e32 v4, v4, v6, vcc +; GISEL-NEXT: v_add_u32_e32 v8, 55, v8 +; GISEL-NEXT: v_cndmask_b32_e32 v4, v4, v9, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v5, v5, v10, vcc ; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v13 -; GISEL-NEXT: v_cndmask_b32_e32 v6, 0, v11, vcc -; GISEL-NEXT: v_lshrrev_b64 v[10:11], v14, -1 +; GISEL-NEXT: v_cndmask_b32_e32 v10, 0, v11, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v11, 0, v12, vcc +; GISEL-NEXT: v_sub_u32_e32 v12, 64, v8 +; GISEL-NEXT: v_cndmask_b32_e64 v14, v4, v0, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e64 v9, v5, v1, s[4:5] +; GISEL-NEXT: v_lshrrev_b64 v[4:5], v8, -1 ; GISEL-NEXT: v_lshlrev_b64 v[12:13], v12, -1 -; GISEL-NEXT: v_subrev_u32_e32 v15, 64, v14 -; GISEL-NEXT: v_or_b32_e32 v16, v10, v12 -; GISEL-NEXT: v_or_b32_e32 v17, v11, v13 +; GISEL-NEXT: v_subrev_u32_e32 v15, 64, v8 +; GISEL-NEXT: v_or_b32_e32 v16, v4, v12 +; GISEL-NEXT: v_or_b32_e32 v17, v5, v13 ; GISEL-NEXT: v_lshrrev_b64 v[12:13], v15, -1 -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v14 -; GISEL-NEXT: v_cndmask_b32_e64 v4, v4, v0, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e64 v5, v5, v1, s[4:5] +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v8 ; GISEL-NEXT: v_cndmask_b32_e32 v12, v12, v16, vcc ; GISEL-NEXT: v_cndmask_b32_e32 v13, v13, v17, vcc -; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v14 -; GISEL-NEXT: v_cndmask_b32_e32 v10, 0, v10, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v11, 0, v11, vcc -; GISEL-NEXT: v_cndmask_b32_e64 v12, v12, -1, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e64 v13, v13, -1, s[4:5] -; GISEL-NEXT: v_and_b32_e32 v10, v10, v2 -; GISEL-NEXT: v_and_b32_e32 v11, v11, v3 -; GISEL-NEXT: v_and_or_b32 v10, v12, v0, v10 -; GISEL-NEXT: v_and_or_b32 v11, v13, v1, v11 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[10:11] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[10:11], exec -; GISEL-NEXT: v_cndmask_b32_e64 v10, 0, 1, vcc -; GISEL-NEXT: s_and_b64 s[10:11], exec, 0 -; GISEL-NEXT: v_or_b32_e32 v4, v4, v10 -; GISEL-NEXT: s_or_b64 s[10:11], s[4:5], s[10:11] -; GISEL-NEXT: .LBB3_10: ; %Flow5 +; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v8 +; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v5, 0, v5, vcc +; GISEL-NEXT: v_cndmask_b32_e64 v8, v12, -1, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e64 v12, v13, -1, s[4:5] +; GISEL-NEXT: v_and_b32_e32 v2, v4, v2 +; GISEL-NEXT: v_and_b32_e32 v3, v5, v3 +; GISEL-NEXT: v_and_or_b32 v0, v8, v0, v2 +; GISEL-NEXT: v_and_or_b32 v1, v12, v1, v3 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; GISEL-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; GISEL-NEXT: v_or_b32_e32 v8, v14, v0 +; GISEL-NEXT: v_mov_b32_e32 v0, v8 +; GISEL-NEXT: v_mov_b32_e32 v1, v9 +; GISEL-NEXT: v_mov_b32_e32 v2, v10 +; GISEL-NEXT: v_mov_b32_e32 v3, v11 +; GISEL-NEXT: .LBB3_7: ; %Flow1 ; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; GISEL-NEXT: ; %bb.11: ; %itofp-sw-bb -; GISEL-NEXT: v_lshlrev_b64 v[6:7], 1, v[2:3] -; GISEL-NEXT: v_lshlrev_b64 v[4:5], 1, v[0:1] +; GISEL-NEXT: .LBB3_8: ; %Flow2 +; GISEL-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; GISEL-NEXT: ; %bb.9: ; %itofp-sw-bb +; GISEL-NEXT: v_lshlrev_b64 v[8:9], 1, v[0:1] +; GISEL-NEXT: v_lshlrev_b64 v[10:11], 1, v[2:3] ; GISEL-NEXT: v_lshrrev_b32_e32 v0, 31, v1 -; GISEL-NEXT: v_or_b32_e32 v6, v6, v0 -; GISEL-NEXT: ; %bb.12: ; %itofp-sw-epilog +; GISEL-NEXT: v_or_b32_e32 v10, v10, v0 +; GISEL-NEXT: v_mov_b32_e32 v0, v8 +; GISEL-NEXT: v_mov_b32_e32 v1, v9 +; GISEL-NEXT: v_mov_b32_e32 v2, v10 +; GISEL-NEXT: v_mov_b32_e32 v3, v11 +; GISEL-NEXT: ; %bb.10: ; %itofp-sw-epilog ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: v_bfe_u32 v0, v4, 2, 1 -; GISEL-NEXT: v_or_b32_e32 v0, v4, v0 +; GISEL-NEXT: v_bfe_u32 v4, v0, 2, 1 +; GISEL-NEXT: v_or_b32_e32 v0, v0, v4 ; GISEL-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v5, vcc -; GISEL-NEXT: v_addc_co_u32_e32 v2, vcc, 0, v6, vcc +; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc +; GISEL-NEXT: v_addc_co_u32_e32 v2, vcc, 0, v2, vcc +; GISEL-NEXT: v_addc_co_u32_e32 v3, vcc, 0, v3, vcc +; GISEL-NEXT: v_mov_b32_e32 v8, 0 +; GISEL-NEXT: v_and_b32_e32 v9, 0x800000, v1 ; GISEL-NEXT: v_lshrrev_b64 v[4:5], 2, v[0:1] -; GISEL-NEXT: v_addc_co_u32_e32 v3, vcc, 0, v7, vcc -; GISEL-NEXT: v_mov_b32_e32 v5, 0 -; GISEL-NEXT: v_and_b32_e32 v6, 0x800000, v1 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[5:6] -; GISEL-NEXT: v_lshlrev_b64 v[5:6], 30, v[2:3] -; GISEL-NEXT: v_lshrrev_b32_e32 v6, 2, v1 -; GISEL-NEXT: v_or_b32_e32 v6, v6, v5 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[8:9] +; GISEL-NEXT: v_lshlrev_b64 v[8:9], 30, v[2:3] +; GISEL-NEXT: v_lshrrev_b32_e32 v5, 2, v1 +; GISEL-NEXT: v_or_b32_e32 v9, v5, v8 ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc -; GISEL-NEXT: ; %bb.13: ; %itofp-if-then20 +; GISEL-NEXT: ; %bb.11: ; %itofp-if-then20 ; GISEL-NEXT: v_lshlrev_b64 v[2:3], 29, v[2:3] ; GISEL-NEXT: v_lshrrev_b64 v[4:5], 3, v[0:1] ; GISEL-NEXT: v_lshrrev_b32_e32 v0, 3, v1 -; GISEL-NEXT: v_or_b32_e32 v6, v0, v2 -; GISEL-NEXT: v_mov_b32_e32 v8, v9 -; GISEL-NEXT: ; %bb.14: ; %Flow +; GISEL-NEXT: v_or_b32_e32 v9, v0, v2 +; GISEL-NEXT: v_mov_b32_e32 v6, v7 +; GISEL-NEXT: ; %bb.12: ; %Flow ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: .LBB3_15: ; %Flow7 +; GISEL-NEXT: .LBB3_13: ; %Flow4 ; GISEL-NEXT: s_or_b64 exec, exec, s[8:9] ; GISEL-NEXT: v_mov_b32_e32 v0, 0x3ff00000 -; GISEL-NEXT: v_lshl_add_u32 v0, v8, 20, v0 -; GISEL-NEXT: v_and_b32_e32 v1, 0xfffff, v6 +; GISEL-NEXT: v_lshl_add_u32 v0, v6, 20, v0 +; GISEL-NEXT: v_and_b32_e32 v1, 0xfffff, v9 ; GISEL-NEXT: v_and_or_b32 v4, v4, -1, 0 ; GISEL-NEXT: v_or3_b32 v5, v1, v0, 0 -; GISEL-NEXT: .LBB3_16: ; %Flow8 +; GISEL-NEXT: .LBB3_14: ; %Flow5 ; GISEL-NEXT: s_or_b64 exec, exec, s[6:7] ; GISEL-NEXT: v_mov_b32_e32 v0, v4 ; GISEL-NEXT: v_mov_b32_e32 v1, v5 @@ -1221,7 +1101,7 @@ define half @sitofp_i128_to_f16(i128 %x) { ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] ; SDAG-NEXT: v_mov_b32_e32 v4, 0 ; SDAG-NEXT: s_and_saveexec_b64 s[6:7], vcc -; SDAG-NEXT: s_cbranch_execz .LBB4_16 +; SDAG-NEXT: s_cbranch_execz .LBB4_14 ; SDAG-NEXT: ; %bb.1: ; %itofp-if-end ; SDAG-NEXT: v_ashrrev_i32_e32 v5, 31, v3 ; SDAG-NEXT: v_xor_b32_e32 v0, v5, v0 @@ -1242,113 +1122,101 @@ define half @sitofp_i128_to_f16(i128 %x) { ; SDAG-NEXT: v_min_u32_e32 v6, v6, v7 ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] ; SDAG-NEXT: v_add_u32_e32 v6, 64, v6 -; SDAG-NEXT: v_cndmask_b32_e32 v9, v6, v2, vcc -; SDAG-NEXT: v_sub_u32_e32 v2, 0x7f, v9 -; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 25, v2 -; SDAG-NEXT: ; implicit-def: $vgpr6 +; SDAG-NEXT: v_cndmask_b32_e32 v7, v6, v2, vcc +; SDAG-NEXT: v_sub_u32_e32 v6, 0x80, v7 +; SDAG-NEXT: v_sub_u32_e32 v2, 0x7f, v7 +; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 25, v6 +; SDAG-NEXT: ; implicit-def: $vgpr8 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc ; SDAG-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; SDAG-NEXT: ; %bb.2: ; %itofp-if-else -; SDAG-NEXT: v_add_u32_e32 v4, 0xffffff98, v9 +; SDAG-NEXT: v_add_u32_e32 v4, 0xffffff98, v7 ; SDAG-NEXT: v_lshlrev_b64 v[0:1], v4, v[0:1] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v4 -; SDAG-NEXT: v_cndmask_b32_e32 v6, 0, v0, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v8, 0, v0, vcc +; SDAG-NEXT: ; implicit-def: $vgpr6 ; SDAG-NEXT: ; implicit-def: $vgpr0_vgpr1 -; SDAG-NEXT: ; implicit-def: $vgpr9 +; SDAG-NEXT: ; implicit-def: $vgpr7 ; SDAG-NEXT: ; implicit-def: $vgpr4_vgpr5 -; SDAG-NEXT: ; %bb.3: ; %Flow6 +; SDAG-NEXT: ; %bb.3: ; %Flow3 ; SDAG-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; SDAG-NEXT: s_cbranch_execz .LBB4_15 +; SDAG-NEXT: s_cbranch_execz .LBB4_13 ; SDAG-NEXT: ; %bb.4: ; %NodeBlock -; SDAG-NEXT: v_sub_u32_e32 v8, 0x80, v9 -; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 25, v8 -; SDAG-NEXT: s_mov_b64 s[10:11], 0 -; SDAG-NEXT: s_mov_b64 s[4:5], 0 +; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 25, v6 +; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc +; SDAG-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; SDAG-NEXT: s_cbranch_execz .LBB4_8 +; SDAG-NEXT: ; %bb.5: ; %LeafBlock +; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 26, v6 ; SDAG-NEXT: s_and_saveexec_b64 s[12:13], vcc -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: ; %bb.5: ; %LeafBlock1 -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 26, v8 -; SDAG-NEXT: s_and_b64 s[4:5], vcc, exec -; SDAG-NEXT: ; %bb.6: ; %Flow3 -; SDAG-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; SDAG-NEXT: ; %bb.7: ; %LeafBlock -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 25, v8 -; SDAG-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; SDAG-NEXT: s_and_b64 s[14:15], vcc, exec -; SDAG-NEXT: s_mov_b64 s[10:11], exec -; SDAG-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; SDAG-NEXT: ; %bb.8: ; %Flow4 -; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: v_mov_b32_e32 v7, v1 -; SDAG-NEXT: v_mov_b32_e32 v6, v0 -; SDAG-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: s_cbranch_execz .LBB4_10 -; SDAG-NEXT: ; %bb.9: ; %itofp-sw-default -; SDAG-NEXT: v_sub_u32_e32 v12, 0x66, v9 +; SDAG-NEXT: s_cbranch_execz .LBB4_7 +; SDAG-NEXT: ; %bb.6: ; %itofp-sw-default +; SDAG-NEXT: v_sub_u32_e32 v12, 0x66, v7 ; SDAG-NEXT: v_sub_u32_e32 v10, 64, v12 -; SDAG-NEXT: v_lshrrev_b64 v[6:7], v12, v[0:1] +; SDAG-NEXT: v_lshrrev_b64 v[8:9], v12, v[0:1] ; SDAG-NEXT: v_lshlrev_b64 v[10:11], v10, v[4:5] -; SDAG-NEXT: v_sub_u32_e32 v13, 38, v9 -; SDAG-NEXT: v_or_b32_e32 v11, v7, v11 -; SDAG-NEXT: v_or_b32_e32 v10, v6, v10 -; SDAG-NEXT: v_lshrrev_b64 v[6:7], v13, v[4:5] +; SDAG-NEXT: v_sub_u32_e32 v13, 38, v7 +; SDAG-NEXT: v_or_b32_e32 v11, v9, v11 +; SDAG-NEXT: v_or_b32_e32 v10, v8, v10 +; SDAG-NEXT: v_lshrrev_b64 v[8:9], v13, v[4:5] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v12 -; SDAG-NEXT: v_add_u32_e32 v14, 26, v9 -; SDAG-NEXT: v_cndmask_b32_e32 v7, v7, v11, vcc +; SDAG-NEXT: v_add_u32_e32 v14, 26, v7 +; SDAG-NEXT: v_cndmask_b32_e32 v9, v9, v11, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v12 -; SDAG-NEXT: v_cndmask_b32_e32 v6, v6, v10, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v8, v8, v10, vcc ; SDAG-NEXT: v_lshrrev_b64 v[10:11], v13, v[0:1] ; SDAG-NEXT: v_lshlrev_b64 v[12:13], v14, v[4:5] -; SDAG-NEXT: v_subrev_u32_e32 v9, 38, v9 -; SDAG-NEXT: v_cndmask_b32_e64 v15, v6, v0, s[4:5] -; SDAG-NEXT: v_or_b32_e32 v6, v13, v11 -; SDAG-NEXT: v_or_b32_e32 v11, v12, v10 -; SDAG-NEXT: v_lshlrev_b64 v[9:10], v9, v[0:1] +; SDAG-NEXT: v_subrev_u32_e32 v7, 38, v7 +; SDAG-NEXT: v_cndmask_b32_e64 v15, v8, v0, s[4:5] +; SDAG-NEXT: v_lshlrev_b64 v[7:8], v7, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v9, v9, v1, s[4:5] +; SDAG-NEXT: v_or_b32_e32 v11, v13, v11 +; SDAG-NEXT: v_or_b32_e32 v10, v12, v10 ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v14 -; SDAG-NEXT: v_cndmask_b32_e64 v7, v7, v1, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v6, v10, v6, vcc +; SDAG-NEXT: v_lshlrev_b64 v[0:1], v14, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e32 v8, v8, v11, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v14 -; SDAG-NEXT: v_cndmask_b32_e64 v10, v6, v5, s[4:5] -; SDAG-NEXT: v_lshlrev_b64 v[5:6], v14, v[0:1] -; SDAG-NEXT: v_cndmask_b32_e32 v9, v9, v11, vcc -; SDAG-NEXT: v_cndmask_b32_e64 v4, v9, v4, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v6, 0, v6, vcc -; SDAG-NEXT: v_cndmask_b32_e32 v9, 0, v5, vcc -; SDAG-NEXT: v_or_b32_e32 v5, v6, v10 -; SDAG-NEXT: v_or_b32_e32 v4, v9, v4 -; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] -; SDAG-NEXT: s_andn2_b64 s[10:11], s[10:11], exec -; SDAG-NEXT: v_cndmask_b32_e64 v4, 0, 1, vcc -; SDAG-NEXT: v_or_b32_e32 v6, v15, v4 -; SDAG-NEXT: .LBB4_10: ; %Flow5 +; SDAG-NEXT: v_cndmask_b32_e32 v7, v7, v10, vcc +; SDAG-NEXT: v_cndmask_b32_e64 v5, v8, v5, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e64 v4, v7, v4, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v0, 0, v0, vcc +; SDAG-NEXT: v_or_b32_e32 v1, v1, v5 +; SDAG-NEXT: v_or_b32_e32 v0, v0, v4 +; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; SDAG-NEXT: v_or_b32_e32 v8, v15, v0 +; SDAG-NEXT: v_mov_b32_e32 v0, v8 +; SDAG-NEXT: v_mov_b32_e32 v1, v9 +; SDAG-NEXT: .LBB4_7: ; %Flow1 ; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; SDAG-NEXT: ; %bb.11: ; %itofp-sw-bb -; SDAG-NEXT: v_lshlrev_b64 v[6:7], 1, v[0:1] -; SDAG-NEXT: ; %bb.12: ; %itofp-sw-epilog +; SDAG-NEXT: .LBB4_8: ; %Flow2 +; SDAG-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; SDAG-NEXT: ; %bb.9: ; %itofp-sw-bb +; SDAG-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; SDAG-NEXT: ; %bb.10: ; %itofp-sw-epilog ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: v_lshrrev_b32_e32 v0, 2, v6 -; SDAG-NEXT: v_and_or_b32 v0, v0, 1, v6 +; SDAG-NEXT: v_lshrrev_b32_e32 v4, 2, v0 +; SDAG-NEXT: v_and_or_b32 v0, v4, 1, v0 ; SDAG-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v7, vcc +; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc ; SDAG-NEXT: v_and_b32_e32 v4, 0x4000000, v0 ; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 0, v4 -; SDAG-NEXT: v_alignbit_b32 v6, v1, v0, 2 +; SDAG-NEXT: v_alignbit_b32 v8, v1, v0, 2 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc -; SDAG-NEXT: ; %bb.13: ; %itofp-if-then20 -; SDAG-NEXT: v_alignbit_b32 v6, v1, v0, 3 -; SDAG-NEXT: v_mov_b32_e32 v2, v8 -; SDAG-NEXT: ; %bb.14: ; %Flow +; SDAG-NEXT: ; %bb.11: ; %itofp-if-then20 +; SDAG-NEXT: v_alignbit_b32 v8, v1, v0, 3 +; SDAG-NEXT: v_mov_b32_e32 v2, v6 +; SDAG-NEXT: ; %bb.12: ; %Flow ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: .LBB4_15: ; %Flow7 +; SDAG-NEXT: .LBB4_13: ; %Flow4 ; SDAG-NEXT: s_or_b64 exec, exec, s[8:9] ; SDAG-NEXT: v_and_b32_e32 v0, 0x80000000, v3 ; SDAG-NEXT: v_lshl_add_u32 v1, v2, 23, 1.0 -; SDAG-NEXT: v_and_b32_e32 v2, 0x7fffff, v6 +; SDAG-NEXT: v_and_b32_e32 v2, 0x7fffff, v8 ; SDAG-NEXT: v_or3_b32 v0, v2, v0, v1 ; SDAG-NEXT: v_cvt_f16_f32_e32 v4, v0 -; SDAG-NEXT: .LBB4_16: ; %Flow8 +; SDAG-NEXT: .LBB4_14: ; %Flow5 ; SDAG-NEXT: s_or_b64 exec, exec, s[6:7] ; SDAG-NEXT: v_mov_b32_e32 v0, v4 ; SDAG-NEXT: s_setpc_b64 s[30:31] @@ -1362,145 +1230,127 @@ define half @sitofp_i128_to_f16(i128 %x) { ; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] ; GISEL-NEXT: v_mov_b32_e32 v4, s4 ; GISEL-NEXT: s_and_saveexec_b64 s[6:7], vcc -; GISEL-NEXT: s_cbranch_execz .LBB4_16 +; GISEL-NEXT: s_cbranch_execz .LBB4_14 ; GISEL-NEXT: ; %bb.1: ; %itofp-if-end -; GISEL-NEXT: v_ashrrev_i32_e32 v8, 31, v3 -; GISEL-NEXT: v_xor_b32_e32 v0, v8, v0 -; GISEL-NEXT: v_xor_b32_e32 v1, v8, v1 -; GISEL-NEXT: v_sub_co_u32_e32 v0, vcc, v0, v8 -; GISEL-NEXT: v_xor_b32_e32 v2, v8, v2 -; GISEL-NEXT: v_subb_co_u32_e32 v1, vcc, v1, v8, vcc -; GISEL-NEXT: v_xor_b32_e32 v3, v8, v3 -; GISEL-NEXT: v_subb_co_u32_e32 v6, vcc, v2, v8, vcc -; GISEL-NEXT: v_subb_co_u32_e32 v7, vcc, v3, v8, vcc -; GISEL-NEXT: v_ffbh_u32_e32 v3, v0 -; GISEL-NEXT: v_ffbh_u32_e32 v2, v1 -; GISEL-NEXT: v_add_u32_e32 v3, 32, v3 -; GISEL-NEXT: v_ffbh_u32_e32 v4, v6 -; GISEL-NEXT: v_min_u32_e32 v2, v2, v3 -; GISEL-NEXT: v_ffbh_u32_e32 v3, v7 -; GISEL-NEXT: v_add_u32_e32 v4, 32, v4 -; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[6:7] -; GISEL-NEXT: v_add_u32_e32 v2, 64, v2 -; GISEL-NEXT: v_min_u32_e32 v3, v3, v4 -; GISEL-NEXT: v_cndmask_b32_e32 v11, v3, v2, vcc -; GISEL-NEXT: v_sub_u32_e32 v9, 0x7f, v11 -; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 24, v9 -; GISEL-NEXT: ; implicit-def: $vgpr2 +; GISEL-NEXT: v_ashrrev_i32_e32 v6, 31, v3 +; GISEL-NEXT: v_xor_b32_e32 v0, v6, v0 +; GISEL-NEXT: v_xor_b32_e32 v1, v6, v1 +; GISEL-NEXT: v_sub_co_u32_e32 v0, vcc, v0, v6 +; GISEL-NEXT: v_xor_b32_e32 v2, v6, v2 +; GISEL-NEXT: v_subb_co_u32_e32 v1, vcc, v1, v6, vcc +; GISEL-NEXT: v_xor_b32_e32 v3, v6, v3 +; GISEL-NEXT: v_subb_co_u32_e32 v2, vcc, v2, v6, vcc +; GISEL-NEXT: v_ffbh_u32_e32 v5, v0 +; GISEL-NEXT: v_subb_co_u32_e32 v3, vcc, v3, v6, vcc +; GISEL-NEXT: v_ffbh_u32_e32 v4, v1 +; GISEL-NEXT: v_add_u32_e32 v5, 32, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v7, v2 +; GISEL-NEXT: v_min_u32_e32 v4, v4, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v5, v3 +; GISEL-NEXT: v_add_u32_e32 v7, 32, v7 +; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[2:3] +; GISEL-NEXT: v_add_u32_e32 v4, 64, v4 +; GISEL-NEXT: v_min_u32_e32 v5, v5, v7 +; GISEL-NEXT: v_cndmask_b32_e32 v5, v5, v4, vcc +; GISEL-NEXT: v_sub_u32_e32 v8, 0x80, v5 +; GISEL-NEXT: v_sub_u32_e32 v7, 0x7f, v5 +; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 24, v8 +; GISEL-NEXT: ; implicit-def: $vgpr4 ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc ; GISEL-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; GISEL-NEXT: ; %bb.2: ; %itofp-if-else -; GISEL-NEXT: v_add_u32_e32 v2, 0xffffff98, v11 +; GISEL-NEXT: v_add_u32_e32 v2, 0xffffff98, v5 ; GISEL-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 -; GISEL-NEXT: v_cndmask_b32_e32 v2, 0, v0, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v0, vcc +; GISEL-NEXT: ; implicit-def: $vgpr8 ; GISEL-NEXT: ; implicit-def: $vgpr0 -; GISEL-NEXT: ; implicit-def: $vgpr11 -; GISEL-NEXT: ; implicit-def: $vgpr6 -; GISEL-NEXT: ; %bb.3: ; %Flow6 +; GISEL-NEXT: ; implicit-def: $vgpr5 +; GISEL-NEXT: ; implicit-def: $vgpr2 +; GISEL-NEXT: ; %bb.3: ; %Flow3 ; GISEL-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; GISEL-NEXT: s_cbranch_execz .LBB4_15 +; GISEL-NEXT: s_cbranch_execz .LBB4_13 ; GISEL-NEXT: ; %bb.4: ; %NodeBlock -; GISEL-NEXT: v_sub_u32_e32 v10, 0x80, v11 -; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 26, v10 -; GISEL-NEXT: s_mov_b64 s[10:11], 0 -; GISEL-NEXT: s_mov_b64 s[4:5], 0 +; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 26, v8 +; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc +; GISEL-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; GISEL-NEXT: s_cbranch_execz .LBB4_8 +; GISEL-NEXT: ; %bb.5: ; %LeafBlock +; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 26, v8 ; GISEL-NEXT: s_and_saveexec_b64 s[12:13], vcc -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: ; %bb.5: ; %LeafBlock1 -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 26, v10 -; GISEL-NEXT: s_andn2_b64 s[4:5], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.6: ; %Flow3 -; GISEL-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; GISEL-NEXT: ; %bb.7: ; %LeafBlock -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 25, v10 -; GISEL-NEXT: s_andn2_b64 s[10:11], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, -1 -; GISEL-NEXT: s_or_b64 s[10:11], s[10:11], s[14:15] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.8: ; %Flow4 -; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: v_mov_b32_e32 v5, v3 -; GISEL-NEXT: v_mov_b32_e32 v4, v2 -; GISEL-NEXT: v_mov_b32_e32 v3, v1 -; GISEL-NEXT: v_mov_b32_e32 v2, v0 -; GISEL-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: s_cbranch_execz .LBB4_10 -; GISEL-NEXT: ; %bb.9: ; %itofp-sw-default -; GISEL-NEXT: v_sub_u32_e32 v12, 0x66, v11 -; GISEL-NEXT: v_sub_u32_e32 v4, 64, v12 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v12, v[0:1] -; GISEL-NEXT: v_lshlrev_b64 v[4:5], v4, v[6:7] -; GISEL-NEXT: v_subrev_u32_e32 v13, 64, v12 -; GISEL-NEXT: v_or_b32_e32 v4, v2, v4 -; GISEL-NEXT: v_or_b32_e32 v5, v3, v5 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v13, v[6:7] -; GISEL-NEXT: v_add_u32_e32 v13, 26, v11 -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v12 -; GISEL-NEXT: v_sub_u32_e32 v11, 64, v13 -; GISEL-NEXT: v_cndmask_b32_e32 v2, v2, v4, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v3, v3, v5, vcc -; GISEL-NEXT: v_cmp_eq_u32_e32 vcc, 0, v12 -; GISEL-NEXT: v_lshrrev_b64 v[4:5], v13, -1 +; GISEL-NEXT: s_cbranch_execz .LBB4_7 +; GISEL-NEXT: ; %bb.6: ; %itofp-sw-default +; GISEL-NEXT: v_sub_u32_e32 v4, 0x66, v5 +; GISEL-NEXT: v_sub_u32_e32 v11, 64, v4 +; GISEL-NEXT: v_lshrrev_b64 v[9:10], v4, v[0:1] +; GISEL-NEXT: v_lshlrev_b64 v[11:12], v11, v[2:3] +; GISEL-NEXT: v_subrev_u32_e32 v13, 64, v4 +; GISEL-NEXT: v_or_b32_e32 v11, v9, v11 +; GISEL-NEXT: v_or_b32_e32 v12, v10, v12 +; GISEL-NEXT: v_lshrrev_b64 v[9:10], v13, v[2:3] +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v4 +; GISEL-NEXT: v_add_u32_e32 v5, 26, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v9, v9, v11, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v10, v10, v12, vcc +; GISEL-NEXT: v_cmp_eq_u32_e32 vcc, 0, v4 +; GISEL-NEXT: v_sub_u32_e32 v11, 64, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v13, v9, v0, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v4, v10, v1, vcc +; GISEL-NEXT: v_lshrrev_b64 v[9:10], v5, -1 ; GISEL-NEXT: v_lshlrev_b64 v[11:12], v11, -1 -; GISEL-NEXT: v_subrev_u32_e32 v14, 64, v13 -; GISEL-NEXT: v_or_b32_e32 v15, v4, v11 -; GISEL-NEXT: v_or_b32_e32 v16, v5, v12 +; GISEL-NEXT: v_subrev_u32_e32 v14, 64, v5 +; GISEL-NEXT: v_or_b32_e32 v15, v9, v11 +; GISEL-NEXT: v_or_b32_e32 v16, v10, v12 ; GISEL-NEXT: v_lshrrev_b64 v[11:12], v14, -1 -; GISEL-NEXT: v_cndmask_b32_e32 v2, v2, v0, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v3, v3, v1, vcc -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v13 +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v5 ; GISEL-NEXT: v_cndmask_b32_e32 v11, v11, v15, vcc ; GISEL-NEXT: v_cndmask_b32_e32 v12, v12, v16, vcc -; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v13 -; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v5, 0, v5, vcc -; GISEL-NEXT: v_cndmask_b32_e64 v11, v11, -1, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e64 v12, v12, -1, s[4:5] -; GISEL-NEXT: v_and_b32_e32 v4, v4, v6 -; GISEL-NEXT: v_and_b32_e32 v5, v5, v7 -; GISEL-NEXT: v_and_or_b32 v4, v11, v0, v4 -; GISEL-NEXT: v_and_or_b32 v5, v12, v1, v5 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[10:11], exec -; GISEL-NEXT: v_cndmask_b32_e64 v4, 0, 1, vcc -; GISEL-NEXT: s_and_b64 s[10:11], exec, 0 -; GISEL-NEXT: v_or_b32_e32 v2, v2, v4 -; GISEL-NEXT: s_or_b64 s[10:11], s[4:5], s[10:11] -; GISEL-NEXT: .LBB4_10: ; %Flow5 +; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v9, 0, v9, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v10, 0, v10, vcc +; GISEL-NEXT: v_cndmask_b32_e64 v5, v11, -1, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e64 v11, v12, -1, s[4:5] +; GISEL-NEXT: v_and_b32_e32 v2, v9, v2 +; GISEL-NEXT: v_and_b32_e32 v3, v10, v3 +; GISEL-NEXT: v_and_or_b32 v0, v5, v0, v2 +; GISEL-NEXT: v_and_or_b32 v1, v11, v1, v3 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; GISEL-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; GISEL-NEXT: v_or_b32_e32 v3, v13, v0 +; GISEL-NEXT: v_mov_b32_e32 v0, v3 +; GISEL-NEXT: v_mov_b32_e32 v1, v4 +; GISEL-NEXT: v_mov_b32_e32 v2, v5 +; GISEL-NEXT: v_mov_b32_e32 v3, v6 +; GISEL-NEXT: .LBB4_7: ; %Flow1 ; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; GISEL-NEXT: ; %bb.11: ; %itofp-sw-bb -; GISEL-NEXT: v_lshlrev_b64 v[2:3], 1, v[0:1] -; GISEL-NEXT: ; %bb.12: ; %itofp-sw-epilog +; GISEL-NEXT: .LBB4_8: ; %Flow2 +; GISEL-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; GISEL-NEXT: ; %bb.9: ; %itofp-sw-bb +; GISEL-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; GISEL-NEXT: ; %bb.10: ; %itofp-sw-epilog ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: v_bfe_u32 v0, v2, 2, 1 -; GISEL-NEXT: v_or_b32_e32 v0, v2, v0 +; GISEL-NEXT: v_bfe_u32 v2, v0, 2, 1 +; GISEL-NEXT: v_or_b32_e32 v0, v0, v2 ; GISEL-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v3, vcc -; GISEL-NEXT: v_lshrrev_b64 v[2:3], 2, v[0:1] -; GISEL-NEXT: v_and_b32_e32 v3, 0x4000000, v0 -; GISEL-NEXT: v_mov_b32_e32 v4, 0 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[3:4] +; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc +; GISEL-NEXT: v_and_b32_e32 v2, 0x4000000, v0 +; GISEL-NEXT: v_mov_b32_e32 v3, 0 +; GISEL-NEXT: v_lshrrev_b64 v[4:5], 2, v[0:1] +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc -; GISEL-NEXT: ; %bb.13: ; %itofp-if-then20 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], 3, v[0:1] -; GISEL-NEXT: v_mov_b32_e32 v9, v10 -; GISEL-NEXT: ; %bb.14: ; %Flow +; GISEL-NEXT: ; %bb.11: ; %itofp-if-then20 +; GISEL-NEXT: v_lshrrev_b64 v[4:5], 3, v[0:1] +; GISEL-NEXT: v_mov_b32_e32 v7, v8 +; GISEL-NEXT: ; %bb.12: ; %Flow ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: .LBB4_15: ; %Flow7 +; GISEL-NEXT: .LBB4_13: ; %Flow4 ; GISEL-NEXT: s_or_b64 exec, exec, s[8:9] -; GISEL-NEXT: v_and_b32_e32 v0, 0x80000000, v8 -; GISEL-NEXT: v_lshl_add_u32 v1, v9, 23, 1.0 -; GISEL-NEXT: v_and_b32_e32 v2, 0x7fffff, v2 +; GISEL-NEXT: v_and_b32_e32 v0, 0x80000000, v6 +; GISEL-NEXT: v_lshl_add_u32 v1, v7, 23, 1.0 +; GISEL-NEXT: v_and_b32_e32 v2, 0x7fffff, v4 ; GISEL-NEXT: v_or3_b32 v0, v2, v0, v1 ; GISEL-NEXT: v_cvt_f16_f32_e32 v4, v0 -; GISEL-NEXT: .LBB4_16: ; %Flow8 +; GISEL-NEXT: .LBB4_14: ; %Flow5 ; GISEL-NEXT: s_or_b64 exec, exec, s[6:7] ; GISEL-NEXT: v_mov_b32_e32 v0, v4 ; GISEL-NEXT: s_setpc_b64 s[30:31] @@ -1517,7 +1367,7 @@ define half @uitofp_i128_to_f16(i128 %x) { ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] ; SDAG-NEXT: v_mov_b32_e32 v4, 0 ; SDAG-NEXT: s_and_saveexec_b64 s[6:7], vcc -; SDAG-NEXT: s_cbranch_execz .LBB5_16 +; SDAG-NEXT: s_cbranch_execz .LBB5_14 ; SDAG-NEXT: ; %bb.1: ; %itofp-if-end ; SDAG-NEXT: v_ffbh_u32_e32 v4, v2 ; SDAG-NEXT: v_add_u32_e32 v4, 32, v4 @@ -1529,114 +1379,100 @@ define half @uitofp_i128_to_f16(i128 %x) { ; SDAG-NEXT: v_min_u32_e32 v5, v5, v6 ; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] ; SDAG-NEXT: v_add_u32_e32 v5, 64, v5 -; SDAG-NEXT: v_cndmask_b32_e32 v8, v5, v4, vcc -; SDAG-NEXT: v_sub_u32_e32 v6, 0x7f, v8 -; SDAG-NEXT: v_mov_b32_e32 v5, v1 -; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 25, v6 -; SDAG-NEXT: v_mov_b32_e32 v4, v0 -; SDAG-NEXT: ; implicit-def: $vgpr9 +; SDAG-NEXT: v_cndmask_b32_e32 v6, v5, v4, vcc +; SDAG-NEXT: v_sub_u32_e32 v5, 0x80, v6 +; SDAG-NEXT: v_sub_u32_e32 v4, 0x7f, v6 +; SDAG-NEXT: v_cmp_gt_i32_e32 vcc, 25, v5 +; SDAG-NEXT: ; implicit-def: $vgpr7 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc ; SDAG-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; SDAG-NEXT: ; %bb.2: ; %itofp-if-else -; SDAG-NEXT: v_add_u32_e32 v2, 0xffffff98, v8 +; SDAG-NEXT: v_add_u32_e32 v2, 0xffffff98, v6 ; SDAG-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 -; SDAG-NEXT: v_cndmask_b32_e32 v9, 0, v0, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v7, 0, v0, vcc +; SDAG-NEXT: ; implicit-def: $vgpr5 ; SDAG-NEXT: ; implicit-def: $vgpr0_vgpr1 -; SDAG-NEXT: ; implicit-def: $vgpr8 +; SDAG-NEXT: ; implicit-def: $vgpr6 ; SDAG-NEXT: ; implicit-def: $vgpr2_vgpr3 -; SDAG-NEXT: ; implicit-def: $vgpr4_vgpr5 -; SDAG-NEXT: ; %bb.3: ; %Flow6 +; SDAG-NEXT: ; %bb.3: ; %Flow3 ; SDAG-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; SDAG-NEXT: s_cbranch_execz .LBB5_15 +; SDAG-NEXT: s_cbranch_execz .LBB5_13 ; SDAG-NEXT: ; %bb.4: ; %NodeBlock -; SDAG-NEXT: v_sub_u32_e32 v7, 0x80, v8 -; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 25, v7 -; SDAG-NEXT: s_mov_b64 s[10:11], 0 -; SDAG-NEXT: s_mov_b64 s[4:5], 0 +; SDAG-NEXT: v_cmp_lt_i32_e32 vcc, 25, v5 +; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc +; SDAG-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; SDAG-NEXT: s_cbranch_execz .LBB5_8 +; SDAG-NEXT: ; %bb.5: ; %LeafBlock +; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 26, v5 ; SDAG-NEXT: s_and_saveexec_b64 s[12:13], vcc -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: ; %bb.5: ; %LeafBlock1 -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 26, v7 -; SDAG-NEXT: s_and_b64 s[4:5], vcc, exec -; SDAG-NEXT: ; %bb.6: ; %Flow3 -; SDAG-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; SDAG-NEXT: ; %bb.7: ; %LeafBlock -; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 25, v7 -; SDAG-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; SDAG-NEXT: s_and_b64 s[14:15], vcc, exec -; SDAG-NEXT: s_mov_b64 s[10:11], exec -; SDAG-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; SDAG-NEXT: ; implicit-def: $vgpr4_vgpr5 -; SDAG-NEXT: ; %bb.8: ; %Flow4 -; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; SDAG-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; SDAG-NEXT: s_cbranch_execz .LBB5_10 -; SDAG-NEXT: ; %bb.9: ; %itofp-sw-default -; SDAG-NEXT: v_sub_u32_e32 v11, 0x66, v8 +; SDAG-NEXT: s_cbranch_execz .LBB5_7 +; SDAG-NEXT: ; %bb.6: ; %itofp-sw-default +; SDAG-NEXT: v_sub_u32_e32 v11, 0x66, v6 ; SDAG-NEXT: v_sub_u32_e32 v9, 64, v11 -; SDAG-NEXT: v_lshrrev_b64 v[4:5], v11, v[0:1] +; SDAG-NEXT: v_lshrrev_b64 v[7:8], v11, v[0:1] ; SDAG-NEXT: v_lshlrev_b64 v[9:10], v9, v[2:3] -; SDAG-NEXT: v_sub_u32_e32 v12, 38, v8 -; SDAG-NEXT: v_or_b32_e32 v10, v5, v10 -; SDAG-NEXT: v_or_b32_e32 v9, v4, v9 -; SDAG-NEXT: v_lshrrev_b64 v[4:5], v12, v[2:3] +; SDAG-NEXT: v_sub_u32_e32 v12, 38, v6 +; SDAG-NEXT: v_or_b32_e32 v10, v8, v10 +; SDAG-NEXT: v_or_b32_e32 v9, v7, v9 +; SDAG-NEXT: v_lshrrev_b64 v[7:8], v12, v[2:3] ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v11 -; SDAG-NEXT: v_add_u32_e32 v13, 26, v8 -; SDAG-NEXT: v_cndmask_b32_e32 v5, v5, v10, vcc +; SDAG-NEXT: v_add_u32_e32 v13, 26, v6 +; SDAG-NEXT: v_cndmask_b32_e32 v8, v8, v10, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v11 -; SDAG-NEXT: v_cndmask_b32_e32 v4, v4, v9, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v7, v7, v9, vcc ; SDAG-NEXT: v_lshrrev_b64 v[9:10], v12, v[0:1] ; SDAG-NEXT: v_lshlrev_b64 v[11:12], v13, v[2:3] -; SDAG-NEXT: v_subrev_u32_e32 v8, 38, v8 -; SDAG-NEXT: v_cndmask_b32_e64 v14, v4, v0, s[4:5] -; SDAG-NEXT: v_or_b32_e32 v4, v12, v10 -; SDAG-NEXT: v_or_b32_e32 v10, v11, v9 -; SDAG-NEXT: v_lshlrev_b64 v[8:9], v8, v[0:1] +; SDAG-NEXT: v_subrev_u32_e32 v6, 38, v6 +; SDAG-NEXT: v_cndmask_b32_e64 v14, v7, v0, s[4:5] +; SDAG-NEXT: v_lshlrev_b64 v[6:7], v6, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v8, v8, v1, s[4:5] +; SDAG-NEXT: v_or_b32_e32 v10, v12, v10 +; SDAG-NEXT: v_or_b32_e32 v9, v11, v9 ; SDAG-NEXT: v_cmp_gt_u32_e32 vcc, 64, v13 -; SDAG-NEXT: v_cndmask_b32_e64 v5, v5, v1, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v4, v9, v4, vcc +; SDAG-NEXT: v_lshlrev_b64 v[0:1], v13, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e32 v7, v7, v10, vcc ; SDAG-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v13 -; SDAG-NEXT: v_cndmask_b32_e64 v9, v4, v3, s[4:5] -; SDAG-NEXT: v_lshlrev_b64 v[3:4], v13, v[0:1] -; SDAG-NEXT: v_cndmask_b32_e32 v8, v8, v10, vcc -; SDAG-NEXT: v_cndmask_b32_e64 v2, v8, v2, s[4:5] -; SDAG-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc -; SDAG-NEXT: v_cndmask_b32_e32 v8, 0, v3, vcc -; SDAG-NEXT: v_or_b32_e32 v3, v4, v9 -; SDAG-NEXT: v_or_b32_e32 v2, v8, v2 -; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] -; SDAG-NEXT: s_andn2_b64 s[10:11], s[10:11], exec -; SDAG-NEXT: v_cndmask_b32_e64 v2, 0, 1, vcc -; SDAG-NEXT: v_or_b32_e32 v4, v14, v2 -; SDAG-NEXT: .LBB5_10: ; %Flow5 +; SDAG-NEXT: v_cndmask_b32_e32 v6, v6, v9, vcc +; SDAG-NEXT: v_cndmask_b32_e64 v3, v7, v3, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e64 v2, v6, v2, s[4:5] +; SDAG-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc +; SDAG-NEXT: v_cndmask_b32_e32 v0, 0, v0, vcc +; SDAG-NEXT: v_or_b32_e32 v1, v1, v3 +; SDAG-NEXT: v_or_b32_e32 v0, v0, v2 +; SDAG-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; SDAG-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; SDAG-NEXT: v_or_b32_e32 v7, v14, v0 +; SDAG-NEXT: v_mov_b32_e32 v0, v7 +; SDAG-NEXT: v_mov_b32_e32 v1, v8 +; SDAG-NEXT: .LBB5_7: ; %Flow1 ; SDAG-NEXT: s_or_b64 exec, exec, s[12:13] -; SDAG-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; SDAG-NEXT: ; %bb.11: ; %itofp-sw-bb -; SDAG-NEXT: v_lshlrev_b64 v[4:5], 1, v[0:1] -; SDAG-NEXT: ; %bb.12: ; %itofp-sw-epilog +; SDAG-NEXT: .LBB5_8: ; %Flow2 +; SDAG-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; SDAG-NEXT: ; %bb.9: ; %itofp-sw-bb +; SDAG-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; SDAG-NEXT: ; %bb.10: ; %itofp-sw-epilog ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: v_lshrrev_b32_e32 v0, 2, v4 -; SDAG-NEXT: v_and_or_b32 v0, v0, 1, v4 +; SDAG-NEXT: v_lshrrev_b32_e32 v2, 2, v0 +; SDAG-NEXT: v_and_or_b32 v0, v2, 1, v0 ; SDAG-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v5, vcc +; SDAG-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc ; SDAG-NEXT: v_and_b32_e32 v2, 0x4000000, v0 ; SDAG-NEXT: v_cmp_ne_u32_e32 vcc, 0, v2 -; SDAG-NEXT: v_alignbit_b32 v9, v1, v0, 2 +; SDAG-NEXT: v_alignbit_b32 v7, v1, v0, 2 ; SDAG-NEXT: s_and_saveexec_b64 s[4:5], vcc -; SDAG-NEXT: ; %bb.13: ; %itofp-if-then20 -; SDAG-NEXT: v_alignbit_b32 v9, v1, v0, 3 -; SDAG-NEXT: v_mov_b32_e32 v6, v7 -; SDAG-NEXT: ; %bb.14: ; %Flow +; SDAG-NEXT: ; %bb.11: ; %itofp-if-then20 +; SDAG-NEXT: v_alignbit_b32 v7, v1, v0, 3 +; SDAG-NEXT: v_mov_b32_e32 v4, v5 +; SDAG-NEXT: ; %bb.12: ; %Flow ; SDAG-NEXT: s_or_b64 exec, exec, s[4:5] -; SDAG-NEXT: .LBB5_15: ; %Flow7 +; SDAG-NEXT: .LBB5_13: ; %Flow4 ; SDAG-NEXT: s_or_b64 exec, exec, s[8:9] -; SDAG-NEXT: v_and_b32_e32 v0, 0x7fffff, v9 -; SDAG-NEXT: v_lshl_or_b32 v0, v6, 23, v0 +; SDAG-NEXT: v_and_b32_e32 v0, 0x7fffff, v7 +; SDAG-NEXT: v_lshl_or_b32 v0, v4, 23, v0 ; SDAG-NEXT: v_add_u32_e32 v0, 1.0, v0 ; SDAG-NEXT: v_cvt_f16_f32_e32 v4, v0 -; SDAG-NEXT: .LBB5_16: ; %Flow8 +; SDAG-NEXT: .LBB5_14: ; %Flow5 ; SDAG-NEXT: s_or_b64 exec, exec, s[6:7] ; SDAG-NEXT: v_mov_b32_e32 v0, v4 ; SDAG-NEXT: s_setpc_b64 s[30:31] @@ -1644,145 +1480,125 @@ define half @uitofp_i128_to_f16(i128 %x) { ; GISEL-LABEL: uitofp_i128_to_f16: ; GISEL: ; %bb.0: ; %itofp-entry ; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) -; GISEL-NEXT: v_mov_b32_e32 v4, v2 -; GISEL-NEXT: v_mov_b32_e32 v5, v3 -; GISEL-NEXT: v_or_b32_e32 v2, v0, v4 -; GISEL-NEXT: v_or_b32_e32 v3, v1, v5 +; GISEL-NEXT: v_or_b32_e32 v4, v0, v2 +; GISEL-NEXT: v_or_b32_e32 v5, v1, v3 ; GISEL-NEXT: s_mov_b32 s4, 0 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] -; GISEL-NEXT: v_mov_b32_e32 v2, s4 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[4:5] +; GISEL-NEXT: v_mov_b32_e32 v4, s4 ; GISEL-NEXT: s_and_saveexec_b64 s[6:7], vcc -; GISEL-NEXT: s_cbranch_execz .LBB5_16 +; GISEL-NEXT: s_cbranch_execz .LBB5_14 ; GISEL-NEXT: ; %bb.1: ; %itofp-if-end -; GISEL-NEXT: v_ffbh_u32_e32 v3, v0 -; GISEL-NEXT: v_ffbh_u32_e32 v2, v1 -; GISEL-NEXT: v_add_u32_e32 v3, 32, v3 -; GISEL-NEXT: v_ffbh_u32_e32 v6, v4 -; GISEL-NEXT: v_min_u32_e32 v2, v2, v3 -; GISEL-NEXT: v_ffbh_u32_e32 v3, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v5, v0 +; GISEL-NEXT: v_ffbh_u32_e32 v4, v1 +; GISEL-NEXT: v_add_u32_e32 v5, 32, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v6, v2 +; GISEL-NEXT: v_min_u32_e32 v4, v4, v5 +; GISEL-NEXT: v_ffbh_u32_e32 v5, v3 ; GISEL-NEXT: v_add_u32_e32 v6, 32, v6 -; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[4:5] -; GISEL-NEXT: v_add_u32_e32 v2, 64, v2 -; GISEL-NEXT: v_min_u32_e32 v3, v3, v6 -; GISEL-NEXT: v_cndmask_b32_e32 v12, v3, v2, vcc -; GISEL-NEXT: v_sub_u32_e32 v10, 0x7f, v12 -; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 24, v10 -; GISEL-NEXT: ; implicit-def: $vgpr2 +; GISEL-NEXT: v_cmp_eq_u64_e32 vcc, 0, v[2:3] +; GISEL-NEXT: v_add_u32_e32 v4, 64, v4 +; GISEL-NEXT: v_min_u32_e32 v5, v5, v6 +; GISEL-NEXT: v_cndmask_b32_e32 v5, v5, v4, vcc +; GISEL-NEXT: v_sub_u32_e32 v7, 0x80, v5 +; GISEL-NEXT: v_sub_u32_e32 v6, 0x7f, v5 +; GISEL-NEXT: v_cmp_ge_i32_e32 vcc, 24, v7 +; GISEL-NEXT: ; implicit-def: $vgpr4 ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc ; GISEL-NEXT: s_xor_b64 s[4:5], exec, s[4:5] ; GISEL-NEXT: ; %bb.2: ; %itofp-if-else -; GISEL-NEXT: v_add_u32_e32 v2, 0xffffff98, v12 +; GISEL-NEXT: v_add_u32_e32 v2, 0xffffff98, v5 ; GISEL-NEXT: v_lshlrev_b64 v[0:1], v2, v[0:1] ; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v2 -; GISEL-NEXT: v_cndmask_b32_e32 v2, 0, v0, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v4, 0, v0, vcc +; GISEL-NEXT: ; implicit-def: $vgpr7 ; GISEL-NEXT: ; implicit-def: $vgpr0 -; GISEL-NEXT: ; implicit-def: $vgpr12 -; GISEL-NEXT: ; implicit-def: $vgpr4 -; GISEL-NEXT: ; %bb.3: ; %Flow6 +; GISEL-NEXT: ; implicit-def: $vgpr5 +; GISEL-NEXT: ; implicit-def: $vgpr2 +; GISEL-NEXT: ; %bb.3: ; %Flow3 ; GISEL-NEXT: s_andn2_saveexec_b64 s[8:9], s[4:5] -; GISEL-NEXT: s_cbranch_execz .LBB5_15 +; GISEL-NEXT: s_cbranch_execz .LBB5_13 ; GISEL-NEXT: ; %bb.4: ; %NodeBlock -; GISEL-NEXT: v_sub_u32_e32 v11, 0x80, v12 -; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 26, v11 -; GISEL-NEXT: s_mov_b64 s[10:11], 0 -; GISEL-NEXT: s_mov_b64 s[4:5], 0 +; GISEL-NEXT: v_cmp_le_i32_e32 vcc, 26, v7 +; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc +; GISEL-NEXT: s_xor_b64 s[10:11], exec, s[4:5] +; GISEL-NEXT: s_cbranch_execz .LBB5_8 +; GISEL-NEXT: ; %bb.5: ; %LeafBlock +; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 26, v7 ; GISEL-NEXT: s_and_saveexec_b64 s[12:13], vcc -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: ; %bb.5: ; %LeafBlock1 -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 26, v11 -; GISEL-NEXT: s_andn2_b64 s[4:5], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.6: ; %Flow3 -; GISEL-NEXT: s_andn2_saveexec_b64 s[12:13], s[12:13] -; GISEL-NEXT: ; %bb.7: ; %LeafBlock -; GISEL-NEXT: v_cmp_ne_u32_e32 vcc, 25, v11 -; GISEL-NEXT: s_andn2_b64 s[10:11], 0, exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, -1 -; GISEL-NEXT: s_or_b64 s[10:11], s[10:11], s[14:15] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[4:5], exec -; GISEL-NEXT: s_and_b64 s[14:15], exec, vcc -; GISEL-NEXT: s_or_b64 s[4:5], s[4:5], s[14:15] -; GISEL-NEXT: ; %bb.8: ; %Flow4 -; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: v_mov_b32_e32 v9, v3 -; GISEL-NEXT: v_mov_b32_e32 v7, v1 -; GISEL-NEXT: v_mov_b32_e32 v6, v0 -; GISEL-NEXT: v_mov_b32_e32 v8, v2 -; GISEL-NEXT: s_and_saveexec_b64 s[12:13], s[4:5] -; GISEL-NEXT: s_xor_b64 s[12:13], exec, s[12:13] -; GISEL-NEXT: s_cbranch_execz .LBB5_10 -; GISEL-NEXT: ; %bb.9: ; %itofp-sw-default -; GISEL-NEXT: v_sub_u32_e32 v8, 0x66, v12 -; GISEL-NEXT: v_sub_u32_e32 v6, 64, v8 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v8, v[0:1] -; GISEL-NEXT: v_lshlrev_b64 v[6:7], v6, v[4:5] -; GISEL-NEXT: v_subrev_u32_e32 v9, 64, v8 -; GISEL-NEXT: v_or_b32_e32 v6, v2, v6 -; GISEL-NEXT: v_or_b32_e32 v7, v3, v7 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v9, v[4:5] -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v8 -; GISEL-NEXT: v_add_u32_e32 v12, 26, v12 -; GISEL-NEXT: v_cndmask_b32_e32 v2, v2, v6, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v3, v3, v7, vcc -; GISEL-NEXT: v_cmp_eq_u32_e32 vcc, 0, v8 -; GISEL-NEXT: v_sub_u32_e32 v8, 64, v12 -; GISEL-NEXT: v_cndmask_b32_e32 v6, v2, v0, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v7, v3, v1, vcc -; GISEL-NEXT: v_lshrrev_b64 v[2:3], v12, -1 -; GISEL-NEXT: v_lshlrev_b64 v[8:9], v8, -1 -; GISEL-NEXT: v_subrev_u32_e32 v13, 64, v12 -; GISEL-NEXT: v_or_b32_e32 v14, v2, v8 -; GISEL-NEXT: v_or_b32_e32 v15, v3, v9 -; GISEL-NEXT: v_lshrrev_b64 v[8:9], v13, -1 -; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v12 -; GISEL-NEXT: v_cndmask_b32_e32 v8, v8, v14, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v9, v9, v15, vcc -; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v12 -; GISEL-NEXT: v_cndmask_b32_e32 v2, 0, v2, vcc -; GISEL-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc -; GISEL-NEXT: v_cndmask_b32_e64 v8, v8, -1, s[4:5] -; GISEL-NEXT: v_cndmask_b32_e64 v9, v9, -1, s[4:5] -; GISEL-NEXT: v_and_b32_e32 v2, v2, v4 -; GISEL-NEXT: v_and_b32_e32 v3, v3, v5 -; GISEL-NEXT: v_and_or_b32 v2, v8, v0, v2 -; GISEL-NEXT: v_and_or_b32 v3, v9, v1, v3 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] -; GISEL-NEXT: s_andn2_b64 s[4:5], s[10:11], exec -; GISEL-NEXT: v_cndmask_b32_e64 v2, 0, 1, vcc -; GISEL-NEXT: s_and_b64 s[10:11], exec, 0 -; GISEL-NEXT: v_or_b32_e32 v6, v6, v2 -; GISEL-NEXT: s_or_b64 s[10:11], s[4:5], s[10:11] -; GISEL-NEXT: .LBB5_10: ; %Flow5 +; GISEL-NEXT: s_cbranch_execz .LBB5_7 +; GISEL-NEXT: ; %bb.6: ; %itofp-sw-default +; GISEL-NEXT: v_sub_u32_e32 v4, 0x66, v5 +; GISEL-NEXT: v_sub_u32_e32 v10, 64, v4 +; GISEL-NEXT: v_lshrrev_b64 v[8:9], v4, v[0:1] +; GISEL-NEXT: v_lshlrev_b64 v[10:11], v10, v[2:3] +; GISEL-NEXT: v_subrev_u32_e32 v12, 64, v4 +; GISEL-NEXT: v_or_b32_e32 v10, v8, v10 +; GISEL-NEXT: v_or_b32_e32 v11, v9, v11 +; GISEL-NEXT: v_lshrrev_b64 v[8:9], v12, v[2:3] +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v4 +; GISEL-NEXT: v_add_u32_e32 v5, 26, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v8, v8, v10, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v9, v9, v11, vcc +; GISEL-NEXT: v_cmp_eq_u32_e32 vcc, 0, v4 +; GISEL-NEXT: v_sub_u32_e32 v10, 64, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v12, v8, v0, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v4, v9, v1, vcc +; GISEL-NEXT: v_lshrrev_b64 v[8:9], v5, -1 +; GISEL-NEXT: v_lshlrev_b64 v[10:11], v10, -1 +; GISEL-NEXT: v_subrev_u32_e32 v13, 64, v5 +; GISEL-NEXT: v_or_b32_e32 v14, v8, v10 +; GISEL-NEXT: v_or_b32_e32 v15, v9, v11 +; GISEL-NEXT: v_lshrrev_b64 v[10:11], v13, -1 +; GISEL-NEXT: v_cmp_gt_u32_e32 vcc, 64, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v10, v10, v14, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v11, v11, v15, vcc +; GISEL-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5 +; GISEL-NEXT: v_cndmask_b32_e32 v8, 0, v8, vcc +; GISEL-NEXT: v_cndmask_b32_e32 v9, 0, v9, vcc +; GISEL-NEXT: v_cndmask_b32_e64 v5, v10, -1, s[4:5] +; GISEL-NEXT: v_cndmask_b32_e64 v10, v11, -1, s[4:5] +; GISEL-NEXT: v_and_b32_e32 v2, v8, v2 +; GISEL-NEXT: v_and_b32_e32 v3, v9, v3 +; GISEL-NEXT: v_and_or_b32 v0, v5, v0, v2 +; GISEL-NEXT: v_and_or_b32 v1, v10, v1, v3 +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1] +; GISEL-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc +; GISEL-NEXT: v_or_b32_e32 v3, v12, v0 +; GISEL-NEXT: v_mov_b32_e32 v0, v3 +; GISEL-NEXT: v_mov_b32_e32 v1, v4 +; GISEL-NEXT: v_mov_b32_e32 v2, v5 +; GISEL-NEXT: v_mov_b32_e32 v3, v6 +; GISEL-NEXT: .LBB5_7: ; %Flow1 ; GISEL-NEXT: s_or_b64 exec, exec, s[12:13] -; GISEL-NEXT: s_and_saveexec_b64 s[4:5], s[10:11] -; GISEL-NEXT: ; %bb.11: ; %itofp-sw-bb -; GISEL-NEXT: v_lshlrev_b64 v[6:7], 1, v[0:1] -; GISEL-NEXT: ; %bb.12: ; %itofp-sw-epilog +; GISEL-NEXT: .LBB5_8: ; %Flow2 +; GISEL-NEXT: s_andn2_saveexec_b64 s[4:5], s[10:11] +; GISEL-NEXT: ; %bb.9: ; %itofp-sw-bb +; GISEL-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1] +; GISEL-NEXT: ; %bb.10: ; %itofp-sw-epilog ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: v_bfe_u32 v0, v6, 2, 1 -; GISEL-NEXT: v_or_b32_e32 v0, v6, v0 +; GISEL-NEXT: v_bfe_u32 v2, v0, 2, 1 +; GISEL-NEXT: v_or_b32_e32 v0, v0, v2 ; GISEL-NEXT: v_add_co_u32_e32 v0, vcc, 1, v0 -; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v7, vcc -; GISEL-NEXT: v_lshrrev_b64 v[2:3], 2, v[0:1] -; GISEL-NEXT: v_and_b32_e32 v3, 0x4000000, v0 -; GISEL-NEXT: v_mov_b32_e32 v4, 0 -; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[3:4] +; GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v1, vcc +; GISEL-NEXT: v_and_b32_e32 v2, 0x4000000, v0 +; GISEL-NEXT: v_mov_b32_e32 v3, 0 +; GISEL-NEXT: v_lshrrev_b64 v[4:5], 2, v[0:1] +; GISEL-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[2:3] ; GISEL-NEXT: s_and_saveexec_b64 s[4:5], vcc -; GISEL-NEXT: ; %bb.13: ; %itofp-if-then20 -; GISEL-NEXT: v_lshrrev_b64 v[2:3], 3, v[0:1] -; GISEL-NEXT: v_mov_b32_e32 v10, v11 -; GISEL-NEXT: ; %bb.14: ; %Flow +; GISEL-NEXT: ; %bb.11: ; %itofp-if-then20 +; GISEL-NEXT: v_lshrrev_b64 v[4:5], 3, v[0:1] +; GISEL-NEXT: v_mov_b32_e32 v6, v7 +; GISEL-NEXT: ; %bb.12: ; %Flow ; GISEL-NEXT: s_or_b64 exec, exec, s[4:5] -; GISEL-NEXT: .LBB5_15: ; %Flow7 +; GISEL-NEXT: .LBB5_13: ; %Flow4 ; GISEL-NEXT: s_or_b64 exec, exec, s[8:9] -; GISEL-NEXT: v_lshl_add_u32 v0, v10, 23, 1.0 +; GISEL-NEXT: v_lshl_add_u32 v0, v6, 23, 1.0 ; GISEL-NEXT: v_mov_b32_e32 v1, 0x7fffff -; GISEL-NEXT: v_and_or_b32 v0, v2, v1, v0 -; GISEL-NEXT: v_cvt_f16_f32_e32 v2, v0 -; GISEL-NEXT: .LBB5_16: ; %Flow8 +; GISEL-NEXT: v_and_or_b32 v0, v4, v1, v0 +; GISEL-NEXT: v_cvt_f16_f32_e32 v4, v0 +; GISEL-NEXT: .LBB5_14: ; %Flow5 ; GISEL-NEXT: s_or_b64 exec, exec, s[6:7] -; GISEL-NEXT: v_mov_b32_e32 v0, v2 +; GISEL-NEXT: v_mov_b32_e32 v0, v4 ; GISEL-NEXT: s_setpc_b64 s[30:31] %cvt = uitofp i128 %x to half ret half %cvt -- GitLab From c957715d721b70591ef32dd9609d6d96de9f0554 Mon Sep 17 00:00:00 2001 From: Simon Pilgrim Date: Fri, 15 Mar 2024 14:12:59 +0000 Subject: [PATCH 006/782] [X86] isGuaranteedNotToBeUndefOrPoisonForTargetNode - generalize shuffle decoding to support more target shuffles in the future. --- llvm/lib/Target/X86/X86ISelLowering.cpp | 32 +++++++++++++++++-------- 1 file changed, 22 insertions(+), 10 deletions(-) diff --git a/llvm/lib/Target/X86/X86ISelLowering.cpp b/llvm/lib/Target/X86/X86ISelLowering.cpp index dbfcb3752ae8..b9a87f9024c7 100644 --- a/llvm/lib/Target/X86/X86ISelLowering.cpp +++ b/llvm/lib/Target/X86/X86ISelLowering.cpp @@ -42665,7 +42665,6 @@ SDValue X86TargetLowering::SimplifyMultipleUseDemandedBitsForTargetNode( bool X86TargetLowering::isGuaranteedNotToBeUndefOrPoisonForTargetNode( SDValue Op, const APInt &DemandedElts, const SelectionDAG &DAG, bool PoisonOnly, unsigned Depth) const { - unsigned EltsBits = Op.getScalarValueSizeInBits(); unsigned NumElts = DemandedElts.getBitWidth(); // TODO: Add more target shuffles. @@ -42673,15 +42672,28 @@ bool X86TargetLowering::isGuaranteedNotToBeUndefOrPoisonForTargetNode( case X86ISD::PSHUFD: case X86ISD::VPERMILPI: { SmallVector Mask; - DecodePSHUFMask(NumElts, EltsBits, Op.getConstantOperandVal(1), Mask); - - APInt DemandedSrcElts = APInt::getZero(NumElts); - for (unsigned I = 0; I != NumElts; ++I) - if (DemandedElts[I]) - DemandedSrcElts.setBit(Mask[I]); - - return DAG.isGuaranteedNotToBeUndefOrPoison( - Op.getOperand(0), DemandedSrcElts, PoisonOnly, Depth + 1); + SmallVector Ops; + if (getTargetShuffleMask(Op.getNode(), Op.getSimpleValueType(), true, Ops, + Mask)) { + SmallVector DemandedSrcElts(Ops.size(), + APInt::getZero(NumElts)); + for (auto M : enumerate(Mask)) { + if (!DemandedElts[M.index()] || M.value() == SM_SentinelZero) + continue; + if (M.value() == SM_SentinelUndef) + return false; + assert(0 <= M.value() && M.value() < (int)(Ops.size() * NumElts) && + "Shuffle mask index out of range"); + DemandedSrcElts[M.value() / NumElts].setBit(M.value() % NumElts); + } + for (auto Op : enumerate(Ops)) + if (!DemandedSrcElts[Op.index()].isZero() && + !DAG.isGuaranteedNotToBeUndefOrPoison( + Op.value(), DemandedSrcElts[Op.index()], PoisonOnly, Depth + 1)) + return false; + return true; + } + break; } } return TargetLowering::isGuaranteedNotToBeUndefOrPoisonForTargetNode( -- GitLab From bf3f86623c4712bff146b5d168cad5f7e658ac56 Mon Sep 17 00:00:00 2001 From: Jay Foad Date: Fri, 15 Mar 2024 13:39:00 +0000 Subject: [PATCH 007/782] [AMDGPU] Simplify GFX11 and GFX12 FLAT saddr field definition It is simpler to define this field correctly in the base class for the Reals for each architecture, than to override it in subclasses for different addressing modes. --- llvm/lib/Target/AMDGPU/FLATInstructions.td | 21 ++++++--------------- 1 file changed, 6 insertions(+), 15 deletions(-) diff --git a/llvm/lib/Target/AMDGPU/FLATInstructions.td b/llvm/lib/Target/AMDGPU/FLATInstructions.td index 594754492564..c6a0d6e89f44 100644 --- a/llvm/lib/Target/AMDGPU/FLATInstructions.td +++ b/llvm/lib/Target/AMDGPU/FLATInstructions.td @@ -174,7 +174,7 @@ class VFLAT_Real op, FLAT_Pseudo ps, string opName = ps.Mnemonic> : bits<8> vaddr; bits<24> offset; - let Inst{6-0} = !if(ps.enabled_saddr, saddr, 0x7f); + let Inst{6-0} = !if(ps.enabled_saddr, saddr, SGPR_NULL_gfx11plus.Index); let Inst{21-14} = op; let Inst{31-26} = 0x3b; let Inst{39-32} = !if(ps.has_vdst, vdst, ?); @@ -2353,6 +2353,7 @@ class FLAT_Real_gfx11 op, FLAT_Pseudo ps, string opName = ps.Mnemonic> let Inst{14} = !if(ps.has_glc, cpol{CPolBit.GLC}, ps.glcValue); let Inst{15} = cpol{CPolBit.SLC}; let Inst{17-16} = seg; + let Inst{54-48} = !if(ps.enabled_saddr, saddr, SGPR_NULL_gfx11plus.Index); let Inst{55} = ps.sve; } @@ -2363,15 +2364,11 @@ multiclass FLAT_Aliases_gfx11 { multiclass FLAT_Real_Base_gfx11 op, string ps, string opName, int renamed = false> : FLAT_Aliases_gfx11 { - def _gfx11 : FLAT_Real_gfx11(ps), opName> { - let Inst{54-48} = SGPR_NULL_gfx11plus.Index; - } + def _gfx11 : FLAT_Real_gfx11(ps), opName>; } multiclass FLAT_Real_RTN_gfx11 op, string ps, string opName> { - def _RTN_gfx11 : FLAT_Real_gfx11(ps#"_RTN"), opName> { - let Inst{54-48} = SGPR_NULL_gfx11plus.Index; - } + def _RTN_gfx11 : FLAT_Real_gfx11(ps#"_RTN"), opName>; } multiclass FLAT_Real_SADDR_gfx11 op, string ps, string opName> { @@ -2384,7 +2381,6 @@ multiclass FLAT_Real_SADDR_RTN_gfx11 op, string ps, string opName> { multiclass FLAT_Real_ST_gfx11 op, string ps, string opName> { def _ST_gfx11 : FLAT_Real_gfx11(ps#"_ST"), opName> { - let Inst{54-48} = SGPR_NULL_gfx11plus.Index; let OtherPredicates = [HasFlatScratchSTMode]; } } @@ -2579,15 +2575,11 @@ multiclass VFLAT_Aliases_gfx12 op, string ps = NAME, string opName = !tolower(NAME), int renamed = false, string alias = ""> : VFLAT_Aliases_gfx12 { - def _gfx12 : VFLAT_Real_gfx12(ps), opName> { - let Inst{6-0} = !cast(SGPR_NULL_gfx11plus.HWEncoding); - } + def _gfx12 : VFLAT_Real_gfx12(ps), opName>; } multiclass VFLAT_Real_RTN_gfx12 op, string ps, string opName> { - def _RTN_gfx12 : VFLAT_Real_gfx12(ps#"_RTN"), opName> { - let Inst{6-0} = !cast(SGPR_NULL_gfx11plus.HWEncoding); - } + def _RTN_gfx12 : VFLAT_Real_gfx12(ps#"_RTN"), opName>; } multiclass VFLAT_Real_SADDR_gfx12 op, string ps, string opName> { @@ -2600,7 +2592,6 @@ multiclass VFLAT_Real_SADDR_RTN_gfx12 op, string ps, string opName> { multiclass VFLAT_Real_ST_gfx12 op, string ps, string opName> { def _ST_gfx12 : VFLAT_Real_gfx12(ps#"_ST"), opName> { - let Inst{6-0} = !cast(SGPR_NULL_gfx11plus.HWEncoding); let OtherPredicates = [HasFlatScratchSTMode]; } } -- GitLab From 65284be2992fc7c6feafc44dda7c0f00df7aacfb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Valentin=20Clement=20=28=E3=83=90=E3=83=AC=E3=83=B3?= =?UTF-8?q?=E3=82=BF=E3=82=A4=E3=83=B3=20=E3=82=AF=E3=83=AC=E3=83=A1?= =?UTF-8?q?=E3=83=B3=29?= Date: Fri, 15 Mar 2024 07:22:22 -0700 Subject: [PATCH 008/782] [flang][cuda] Lower dim3 grid z correctly on calls (#85346) --- flang/lib/Lower/ConvertCall.cpp | 6 ++++-- flang/test/Lower/CUDA/cuda-kernel-calls.cuf | 6 ++++-- 2 files changed, 8 insertions(+), 4 deletions(-) diff --git a/flang/lib/Lower/ConvertCall.cpp b/flang/lib/Lower/ConvertCall.cpp index 990912195d14..95569337a06e 100644 --- a/flang/lib/Lower/ConvertCall.cpp +++ b/flang/lib/Lower/ConvertCall.cpp @@ -416,7 +416,7 @@ std::pair Fortran::lower::genCallOpAndResult( mlir::Type i32Ty = builder.getI32Type(); mlir::Value one = builder.createIntegerConstant(loc, i32Ty, 1); - mlir::Value grid_x, grid_y; + mlir::Value grid_x, grid_y, grid_z; if (caller.getCallDescription().chevrons()[0].GetType()->category() == Fortran::common::TypeCategory::Integer) { // If grid is an integer, it is converted to dim3(grid,1,1). Since z is @@ -426,11 +426,13 @@ std::pair Fortran::lower::genCallOpAndResult( fir::getBase(converter.genExprValue( caller.getCallDescription().chevrons()[0], stmtCtx))); grid_y = one; + grid_z = one; } else { auto dim3Addr = converter.genExprAddr( caller.getCallDescription().chevrons()[0], stmtCtx); grid_x = readDim3Value(builder, loc, fir::getBase(dim3Addr), "x"); grid_y = readDim3Value(builder, loc, fir::getBase(dim3Addr), "y"); + grid_z = readDim3Value(builder, loc, fir::getBase(dim3Addr), "z"); } mlir::Value block_x, block_y, block_z; @@ -466,7 +468,7 @@ std::pair Fortran::lower::genCallOpAndResult( caller.getCallDescription().chevrons()[3], stmtCtx))); builder.create( - loc, funcType.getResults(), funcSymbolAttr, grid_x, grid_y, one, + loc, funcType.getResults(), funcSymbolAttr, grid_x, grid_y, grid_z, block_x, block_y, block_z, bytes, stream, operands); callNumResults = 0; } else if (caller.requireDispatchCall()) { diff --git a/flang/test/Lower/CUDA/cuda-kernel-calls.cuf b/flang/test/Lower/CUDA/cuda-kernel-calls.cuf index d5dabaa1df96..55b5246e59a0 100644 --- a/flang/test/Lower/CUDA/cuda-kernel-calls.cuf +++ b/flang/test/Lower/CUDA/cuda-kernel-calls.cuf @@ -20,13 +20,15 @@ contains call dev_kernel0<<<10, 20>>>() ! CHECK: fir.cuda_kernel_launch @_QMtest_callPdev_kernel0<<<%c10{{.*}}, %c1{{.*}}, %c1{{.*}}, %c20{{.*}}, %c1{{.*}}, %c1{{.*}}>>>() - call dev_kernel0<<< __builtin_dim3(1,1), __builtin_dim3(32,1,1) >>> + call dev_kernel0<<< __builtin_dim3(1,1,4), __builtin_dim3(32,1,1) >>> ! CHECK: %[[ADDR_DIM3_GRID:.*]] = fir.address_of(@_QQro._QM__fortran_builtinsT__builtin_dim3.{{.*}}) : !fir.ref> ! CHECK: %[[DIM3_GRID:.*]]:2 = hlfir.declare %[[ADDR_DIM3_GRID]] {fortran_attrs = #fir.var_attrs, uniq_name = "_QQro._QM__fortran_builtinsT__builtin_dim3.0"} : (!fir.ref>) -> (!fir.ref>, !fir.ref>) ! CHECK: %[[GRID_X:.*]] = hlfir.designate %[[DIM3_GRID]]#1{"x"} : (!fir.ref>) -> !fir.ref ! CHECK: %[[GRID_X_LOAD:.*]] = fir.load %[[GRID_X]] : !fir.ref ! CHECK: %[[GRID_Y:.*]] = hlfir.designate %[[DIM3_GRID]]#1{"y"} : (!fir.ref>) -> !fir.ref ! CHECK: %[[GRID_Y_LOAD:.*]] = fir.load %[[GRID_Y]] : !fir.ref +! CHECK: %[[GRID_Z:.*]] = hlfir.designate %[[DIM3_GRID]]#1{"z"} : (!fir.ref>) -> !fir.ref +! CHECK: %[[GRID_Z_LOAD:.*]] = fir.load %[[GRID_Z]] : !fir.ref ! CHECK: %[[ADDR_DIM3_BLOCK:.*]] = fir.address_of(@_QQro._QM__fortran_builtinsT__builtin_dim3.{{.*}}) : !fir.ref> ! CHECK: %[[DIM3_BLOCK:.*]]:2 = hlfir.declare %[[ADDR_DIM3_BLOCK]] {fortran_attrs = #fir.var_attrs, uniq_name = "_QQro._QM__fortran_builtinsT__builtin_dim3.1"} : (!fir.ref>) -> (!fir.ref>, !fir.ref>) ! CHECK: %[[BLOCK_X:.*]] = hlfir.designate %[[DIM3_BLOCK]]#1{"x"} : (!fir.ref>) -> !fir.ref @@ -35,7 +37,7 @@ contains ! CHECK: %[[BLOCK_Y_LOAD:.*]] = fir.load %[[BLOCK_Y]] : !fir.ref ! CHECK: %[[BLOCK_Z:.*]] = hlfir.designate %[[DIM3_BLOCK]]#1{"z"} : (!fir.ref>) -> !fir.ref ! CHECK: %[[BLOCK_Z_LOAD:.*]] = fir.load %[[BLOCK_Z]] : !fir.ref -! CHECK: fir.cuda_kernel_launch @_QMtest_callPdev_kernel0<<<%[[GRID_X_LOAD]], %[[GRID_Y_LOAD]], %c1{{.*}}, %[[BLOCK_X_LOAD]], %[[BLOCK_Y_LOAD]], %[[BLOCK_Z_LOAD]]>>>() +! CHECK: fir.cuda_kernel_launch @_QMtest_callPdev_kernel0<<<%[[GRID_X_LOAD]], %[[GRID_Y_LOAD]], %[[GRID_Z_LOAD]], %[[BLOCK_X_LOAD]], %[[BLOCK_Y_LOAD]], %[[BLOCK_Z_LOAD]]>>>() call dev_kernel0<<<10, 20, 2>>>() ! CHECK: fir.cuda_kernel_launch @_QMtest_callPdev_kernel0<<<%c10{{.*}}, %c1{{.*}}, %c1{{.*}}, %c20{{.*}}, %c1{{.*}}, %c1{{.*}}, %c2{{.*}}>>>() -- GitLab From 0e0bfacff71859d1f9212205f8f873d47029d3fb Mon Sep 17 00:00:00 2001 From: yonghong-song Date: Fri, 15 Mar 2024 07:24:28 -0700 Subject: [PATCH 009/782] [BPF] Add support for may_goto insn (#85358) Alexei added may_goto insn in [1]. The asm syntax for may_goto looks like may_goto