From f4efef02fdad93da7b0d62be0ecdaf5aee8a24b0 Mon Sep 17 00:00:00 2001 From: "William S. Moses" Date: Sat, 10 Oct 2026 19:48:34 -0500 Subject: [PATCH] Update jax-ml/jax to commit 6aaf234b1ee98800fcb80b9014a052a66dadbd8a Rebase the patches onto the XLA (ed6af492) and LLVM (b08e0059) it pins: - xla.patch: the rocm BUILD.tpl hunk takes its new context; the topology_util.cc hunk (#3376) is upstream and dropped. - llvm_amdgpu_bb_prolog.patch: new context only; llvm/llvm-project#225433 is still open. - llvm_mlir_import_inrange_width.patch (#3367) and llvm_orc_unw_revert.patch are upstream and no longer applied. - The Windows guard appended to stablehlo's temporary.patch is anchored where XLA's own InterpreterOps.cpp change now leaves evalRunParallelOp. And build against that LLVM: - vendored Polly: IRInserter::InsertHelper no longer overrides, as the IRBuilder inserter is no longer virtual (llvm/llvm-project#229700). - populateGpuToROCDLConversionPatterns takes a ROCDL::TargetInfo in place of an amdgpu::Chipset (llvm/llvm-project#223562); the default one is the same unknown target. - raiseimmerse3: the new arith fold of or(x, or(y, x)) changes the order affine-cfg reaches the compare in, which is now left as an affine.if of true/false rather than turned into a cmpi eq; the same happens on main when the input is given that fold already. - CompareOpCanon: XLA's stablehlo patch gives the reference compareOp the op's optional compare type; pass it through. Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_016zErYp7upmqr4NHfhod9UD --- patches/llvm_amdgpu_bb_prolog.patch | 146 +++++++++--------- patches/xla.patch | 33 +--- .../jax/Passes/ConvertPolygeistToLLVM.cpp | 2 +- src/enzyme_ad/jax/Passes/EnzymeHLOOpt.cpp | 2 +- .../polly/include/polly/CodeGen/IRBuilder.h | 2 +- test/lit_tests/raising/raiseimmerse3.mlir | 14 +- workspace.bzl | 10 +- 7 files changed, 98 insertions(+), 111 deletions(-) diff --git a/patches/llvm_amdgpu_bb_prolog.patch b/patches/llvm_amdgpu_bb_prolog.patch index e21a98310b..c15449eefa 100644 --- a/patches/llvm_amdgpu_bb_prolog.patch +++ b/patches/llvm_amdgpu_bb_prolog.patch @@ -1,5 +1,5 @@ diff --git a/llvm/include/llvm/CodeGen/MachineBasicBlock.h b/llvm/include/llvm/CodeGen/MachineBasicBlock.h -index 2155629..4d88ec8 100644 +index 215562968158..4d88ec8cb212 100644 --- a/llvm/include/llvm/CodeGen/MachineBasicBlock.h +++ b/llvm/include/llvm/CodeGen/MachineBasicBlock.h @@ -903,19 +903,42 @@ public: @@ -54,7 +54,7 @@ index 2155629..4d88ec8 100644 /// block. If a terminator does not exist, it returns end(). LLVM_ABI iterator getFirstTerminator(); diff --git a/llvm/include/llvm/CodeGen/MachineInstr.h b/llvm/include/llvm/CodeGen/MachineInstr.h -index f66dff6..544e288 100644 +index f66dff6e5bb3..544e288cf0bb 100644 --- a/llvm/include/llvm/CodeGen/MachineInstr.h +++ b/llvm/include/llvm/CodeGen/MachineInstr.h @@ -129,8 +129,10 @@ public: @@ -70,10 +70,10 @@ index f66dff6..544e288 100644 static constexpr uint32_t getPoisonGeneratingFlags() { diff --git a/llvm/include/llvm/CodeGen/TargetInstrInfo.h b/llvm/include/llvm/CodeGen/TargetInstrInfo.h -index b013511..94db7f1 100644 +index 3191b9e50a86..ce2c432d92f7 100644 --- a/llvm/include/llvm/CodeGen/TargetInstrInfo.h +++ b/llvm/include/llvm/CodeGen/TargetInstrInfo.h -@@ -2190,15 +2190,6 @@ public: +@@ -2184,15 +2184,6 @@ public: return Inst.isReturn() && Inst.isCall(); } @@ -90,7 +90,7 @@ index b013511..94db7f1 100644 /// range of a register in register allocation. virtual unsigned getLiveRangeSplitOpcode(Register Reg, diff --git a/llvm/lib/CodeGen/FixupStatepointCallerSaved.cpp b/llvm/lib/CodeGen/FixupStatepointCallerSaved.cpp -index 8d2f9bf..2be3def 100644 +index 8d2f9bfb6c5c..2be3def1b09d 100644 --- a/llvm/lib/CodeGen/FixupStatepointCallerSaved.cpp +++ b/llvm/lib/CodeGen/FixupStatepointCallerSaved.cpp @@ -455,8 +455,7 @@ public: @@ -104,7 +104,7 @@ index 8d2f9bf..2be3def 100644 LLVM_DEBUG(dbgs() << "...also reload at EHPad " << printMBBReference(*EHPad) << "\n"); diff --git a/llvm/lib/CodeGen/InlineSpiller.cpp b/llvm/lib/CodeGen/InlineSpiller.cpp -index f3682a2..6e78bf5 100644 +index f3682a2e2480..6e78bf50c372 100644 --- a/llvm/lib/CodeGen/InlineSpiller.cpp +++ b/llvm/lib/CodeGen/InlineSpiller.cpp @@ -457,7 +457,7 @@ bool InlineSpiller::hoistSpillInsideBB(LiveInterval &SpillLI, @@ -169,7 +169,7 @@ index f3682a2..6e78bf5 100644 for (const MachineInstr &MI : make_range(MIS.begin(), MII)) getVDefInterval(MI, LIS); diff --git a/llvm/lib/CodeGen/LiveRangeEdit.cpp b/llvm/lib/CodeGen/LiveRangeEdit.cpp -index 1d6b835..1a9c9df 100644 +index 1d6b835f1646..1a9c9df02701 100644 --- a/llvm/lib/CodeGen/LiveRangeEdit.cpp +++ b/llvm/lib/CodeGen/LiveRangeEdit.cpp @@ -91,6 +91,8 @@ SlotIndex LiveRangeEdit::rematerializeAt( @@ -182,10 +182,10 @@ index 1d6b835..1a9c9df 100644 ++NumReMaterialization; diff --git a/llvm/lib/CodeGen/MIRParser/MILexer.cpp b/llvm/lib/CodeGen/MIRParser/MILexer.cpp -index c220047..f81e055 100644 +index a07047a25bb6..c2c1cc71499f 100644 --- a/llvm/lib/CodeGen/MIRParser/MILexer.cpp +++ b/llvm/lib/CodeGen/MIRParser/MILexer.cpp -@@ -304,6 +304,7 @@ static MIToken::TokenKind getIdentifierKind(StringRef Identifier) { +@@ -306,6 +306,7 @@ static MIToken::TokenKind getIdentifierKind(StringRef Identifier) { .Case("noconvergent", MIToken::kw_noconvergent) .Case("mmra", MIToken::kw_mmra) .Case("lr-split", MIToken::kw_lr_split) @@ -194,10 +194,10 @@ index c220047..f81e055 100644 } diff --git a/llvm/lib/CodeGen/MIRParser/MILexer.h b/llvm/lib/CodeGen/MIRParser/MILexer.h -index 452a7bb..8272a76 100644 +index 6c27eff9cd4d..e56849773809 100644 --- a/llvm/lib/CodeGen/MIRParser/MILexer.h +++ b/llvm/lib/CodeGen/MIRParser/MILexer.h -@@ -156,6 +156,7 @@ struct MIToken { +@@ -158,6 +158,7 @@ struct MIToken { kw_noconvergent, kw_mmra, kw_lr_split, @@ -206,10 +206,10 @@ index 452a7bb..8272a76 100644 // Metadata types. kw_distinct, diff --git a/llvm/lib/CodeGen/MIRParser/MIParser.cpp b/llvm/lib/CodeGen/MIRParser/MIParser.cpp -index bd67f06..028fc36 100644 +index 1a4bcbd0345b..482ccf20bc26 100644 --- a/llvm/lib/CodeGen/MIRParser/MIParser.cpp +++ b/llvm/lib/CodeGen/MIRParser/MIParser.cpp -@@ -1414,7 +1414,8 @@ bool MIParser::parseInstruction(unsigned &OpCode, unsigned &Flags) { +@@ -1426,7 +1426,8 @@ bool MIParser::parseInstruction(unsigned &OpCode, unsigned &Flags) { Token.is(MIToken::kw_samesign) || Token.is(MIToken::kw_inbounds) || Token.is(MIToken::kw_nonnull) || @@ -219,7 +219,7 @@ index bd67f06..028fc36 100644 // clang-format on // Mine frame and fast math flags if (Token.is(MIToken::kw_frame_setup)) -@@ -1461,6 +1462,8 @@ bool MIParser::parseInstruction(unsigned &OpCode, unsigned &Flags) { +@@ -1473,6 +1474,8 @@ bool MIParser::parseInstruction(unsigned &OpCode, unsigned &Flags) { Flags |= MachineInstr::NonNull; if (Token.is(MIToken::kw_lr_split)) Flags |= MachineInstr::LRSplit; @@ -229,7 +229,7 @@ index bd67f06..028fc36 100644 lex(); } diff --git a/llvm/lib/CodeGen/MIRPrinter.cpp b/llvm/lib/CodeGen/MIRPrinter.cpp -index 7ac0c40..ef6c7e6 100644 +index def751e6df6b..e4a5ad20e2dc 100644 --- a/llvm/lib/CodeGen/MIRPrinter.cpp +++ b/llvm/lib/CodeGen/MIRPrinter.cpp @@ -908,6 +908,8 @@ static void printMI(raw_ostream &OS, MFPrintState &State, @@ -242,7 +242,7 @@ index 7ac0c40..ef6c7e6 100644 // NOTE: Please add new MIFlags also to the MI_FLAGS_STR in // llvm/utils/UpdateTestChecks/mir.py. diff --git a/llvm/lib/CodeGen/MachineBasicBlock.cpp b/llvm/lib/CodeGen/MachineBasicBlock.cpp -index 006a9fc..9106794 100644 +index a5ba8823830e..505f2551a716 100644 --- a/llvm/lib/CodeGen/MachineBasicBlock.cpp +++ b/llvm/lib/CodeGen/MachineBasicBlock.cpp @@ -210,13 +210,24 @@ MachineBasicBlock::iterator MachineBasicBlock::getFirstNonPHI() { @@ -291,7 +291,7 @@ index 006a9fc..9106794 100644 // FIXME: This needs to change if we wish to bundle labels / dbg_values // inside the bundle. diff --git a/llvm/lib/CodeGen/MachineInstr.cpp b/llvm/lib/CodeGen/MachineInstr.cpp -index 4978967..ce7dac9 100644 +index 4978967c7cb1..ce7dac9fcbe9 100644 --- a/llvm/lib/CodeGen/MachineInstr.cpp +++ b/llvm/lib/CodeGen/MachineInstr.cpp @@ -1920,6 +1920,8 @@ void MachineInstr::print(raw_ostream &OS, ModuleSlotTracker &MST, @@ -304,7 +304,7 @@ index 4978967..ce7dac9 100644 // Print the opcode name. if (TII) diff --git a/llvm/lib/CodeGen/MachineSink.cpp b/llvm/lib/CodeGen/MachineSink.cpp -index cbd2598..834effe 100644 +index c3caac2d9715..58e249c49d6b 100644 --- a/llvm/lib/CodeGen/MachineSink.cpp +++ b/llvm/lib/CodeGen/MachineSink.cpp @@ -343,7 +343,7 @@ static bool blockPrologueInterferes(const MachineBasicBlock *BB, @@ -326,7 +326,7 @@ index cbd2598..834effe 100644 "can't sink.\n"); continue; diff --git a/llvm/lib/CodeGen/MachineVerifier.cpp b/llvm/lib/CodeGen/MachineVerifier.cpp -index cae5012..9cecf07 100644 +index cae5012b41c3..9cecf0709a84 100644 --- a/llvm/lib/CodeGen/MachineVerifier.cpp +++ b/llvm/lib/CodeGen/MachineVerifier.cpp @@ -159,6 +159,7 @@ struct MachineVerifier { @@ -367,10 +367,10 @@ index cae5012..9cecf07 100644 if (MI->isInlineAsm()) verifyInlineAsm(MI); diff --git a/llvm/lib/CodeGen/PHIElimination.cpp b/llvm/lib/CodeGen/PHIElimination.cpp -index de69ca8..cd95778 100644 +index 2a3d7dfc6e34..a85e151abead 100644 --- a/llvm/lib/CodeGen/PHIElimination.cpp +++ b/llvm/lib/CodeGen/PHIElimination.cpp -@@ -423,6 +423,7 @@ void PHIEliminationImpl::LowerPHINode(MachineBasicBlock &MBB, +@@ -440,6 +440,7 @@ void PHIEliminationImpl::LowerPHINode(MachineBasicBlock &MBB, PHICopy = TII->createPHIDestinationCopy( MBB, AfterPHIsIt, MPhi->getDebugLoc(), IncomingReg, DestReg); } @@ -379,10 +379,10 @@ index de69ca8..cd95778 100644 if (MPhi->peekDebugInstrNum() && IncomingReg) { // If referred to by debug-info, store where this PHI was. diff --git a/llvm/lib/CodeGen/RegAllocFast.cpp b/llvm/lib/CodeGen/RegAllocFast.cpp -index 078b3f3..86db6f2 100644 +index 76237fa260d3..16896af81444 100644 --- a/llvm/lib/CodeGen/RegAllocFast.cpp +++ b/llvm/lib/CodeGen/RegAllocFast.cpp -@@ -616,7 +616,8 @@ void RegAllocFastImpl::spill(MachineBasicBlock::iterator Before, +@@ -641,7 +641,8 @@ void RegAllocFastImpl::spill(MachineBasicBlock::iterator Before, LLVM_DEBUG(dbgs() << " to stack slot #" << FI << '\n'); const TargetRegisterClass &RC = *MRI->getRegClass(VirtReg); @@ -392,7 +392,7 @@ index 078b3f3..86db6f2 100644 ++NumStores; MachineBasicBlock::iterator FirstTerm = MBB->getFirstTerminator(); -@@ -673,7 +674,8 @@ void RegAllocFastImpl::reload(MachineBasicBlock::iterator Before, +@@ -698,7 +699,8 @@ void RegAllocFastImpl::reload(MachineBasicBlock::iterator Before, << printReg(PhysReg, TRI) << '\n'); int FI = getStackSpaceFor(VirtReg); const TargetRegisterClass &RC = *MRI->getRegClass(VirtReg); @@ -402,7 +402,7 @@ index 078b3f3..86db6f2 100644 ++NumLoads; } -@@ -691,7 +693,7 @@ MachineBasicBlock::iterator RegAllocFastImpl::getMBBBeginInsertionPoint( +@@ -716,7 +718,7 @@ MachineBasicBlock::iterator RegAllocFastImpl::getMBBBeginInsertionPoint( } // Skip prologues and inlineasm_br spills to place reloads afterwards. @@ -411,7 +411,7 @@ index 078b3f3..86db6f2 100644 break; // However if a prolog instruction reads a register that needs to be -@@ -1106,7 +1108,8 @@ bool RegAllocFastImpl::defineLiveThroughVirtReg(MachineInstr &MI, +@@ -1159,7 +1161,8 @@ bool RegAllocFastImpl::defineLiveThroughVirtReg(MachineInstr &MI, << printReg(PrevReg, TRI) << '\n'); BuildMI(*MBB, InsertBefore, MI.getDebugLoc(), TII->get(TargetOpcode::COPY), PrevReg) @@ -421,7 +421,7 @@ index 078b3f3..86db6f2 100644 } MachineOperand &MO = MI.getOperand(OpNum); if (MO.getSubReg() && !MO.isUndef()) { -@@ -1172,7 +1175,8 @@ bool RegAllocFastImpl::defineVirtReg(MachineInstr &MI, unsigned OpNum, +@@ -1225,7 +1228,8 @@ bool RegAllocFastImpl::defineVirtReg(MachineInstr &MI, unsigned OpNum, if (MO.isMBB()) { MachineBasicBlock *Succ = MO.getMBB(); TII->storeRegToStackSlot(*Succ, Succ->begin(), PhysReg, Kill, FI, @@ -432,7 +432,7 @@ index 078b3f3..86db6f2 100644 Succ->addLiveIn(PhysReg); } diff --git a/llvm/lib/CodeGen/RegAllocGreedy.cpp b/llvm/lib/CodeGen/RegAllocGreedy.cpp -index af4fc60..b7b41fb 100644 +index 51faa9dd9103..248ab92b6eeb 100644 --- a/llvm/lib/CodeGen/RegAllocGreedy.cpp +++ b/llvm/lib/CodeGen/RegAllocGreedy.cpp @@ -844,8 +844,7 @@ bool RAGreedy::addThroughConstraints(InterferenceCache::Cursor Intf, @@ -446,7 +446,7 @@ index af4fc60..b7b41fb 100644 ? Indexes->getMBBEndIdx(MBB) : LIS->getInstructionIndex(*InsertPt); diff --git a/llvm/lib/CodeGen/SplitKit.cpp b/llvm/lib/CodeGen/SplitKit.cpp -index 404cfe4..ac43e67 100644 +index 404cfe43289d..ac43e67b690a 100644 --- a/llvm/lib/CodeGen/SplitKit.cpp +++ b/llvm/lib/CodeGen/SplitKit.cpp @@ -672,6 +672,8 @@ VNInfo *SplitEditor::defFromParent(unsigned RegIdx, const VNInfo *ParentVNI, @@ -479,10 +479,10 @@ index 404cfe4..ac43e67 100644 LLVM_DEBUG(dump()); return VNI->def; diff --git a/llvm/lib/CodeGen/TargetInstrInfo.cpp b/llvm/lib/CodeGen/TargetInstrInfo.cpp -index 92fc628..abb36cb 100644 +index 0b87900194b8..947190612d9e 100644 --- a/llvm/lib/CodeGen/TargetInstrInfo.cpp +++ b/llvm/lib/CodeGen/TargetInstrInfo.cpp -@@ -788,6 +788,10 @@ MachineInstr *TargetInstrInfo::foldMemoryOperand(MachineInstr &MI, +@@ -789,6 +789,10 @@ MachineInstr *TargetInstrInfo::foldMemoryOperand(MachineInstr &MI, // call instructions. We need copy it form old instruction. NewMI->cloneInstrSymbols(MF, MI); @@ -493,7 +493,7 @@ index 92fc628..abb36cb 100644 return NewMI; } -@@ -801,17 +805,21 @@ MachineInstr *TargetInstrInfo::foldMemoryOperand(MachineInstr &MI, +@@ -802,17 +806,21 @@ MachineInstr *TargetInstrInfo::foldMemoryOperand(MachineInstr &MI, const MachineOperand &MO = MI.getOperand(1 - Ops[0]); MachineBasicBlock::iterator Pos = MI; @@ -519,7 +519,7 @@ index 92fc628..abb36cb 100644 return &*--Pos; } diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp -index 831f3e3..7f0dcf3 100644 +index 948b0bc0fbaa..45c50039b3ed 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp @@ -1818,9 +1818,12 @@ bool AMDGPUInstructionSelector::selectEndCfIntrinsic(MachineInstr &MI) const { @@ -539,10 +539,10 @@ index 831f3e3..7f0dcf3 100644 Register Reg = MI.getOperand(1).getReg(); MI.eraseFromParent(); diff --git a/llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp b/llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp -index 605e84b..686a02c 100644 +index ca7256d945d3..0fcfbd3fb37d 100644 --- a/llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp +++ b/llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp -@@ -443,10 +443,9 @@ static bool isReachable(const MachineInstr *From, +@@ -439,10 +439,9 @@ static bool isReachable(const MachineInstr *From, } // Return the first non-prologue instruction in the block. @@ -555,7 +555,7 @@ index 605e84b..686a02c 100644 ++I; return I; -@@ -558,7 +557,7 @@ static bool hoistAndMergeSGPRInits(unsigned Reg, +@@ -561,7 +560,7 @@ static bool hoistAndMergeSGPRInits(unsigned Reg, continue; } @@ -564,7 +564,7 @@ index 605e84b..686a02c 100644 if (!interferes(MI1, I) && !interferes(MI2, I)) { LLVM_DEBUG(dbgs() << "Erasing from " -@@ -597,11 +596,11 @@ static bool hoistAndMergeSGPRInits(unsigned Reg, +@@ -600,11 +599,11 @@ static bool hoistAndMergeSGPRInits(unsigned Reg, auto &Defs = Init.second; for (auto *MI : Defs) { auto *MBB = MI->getParent(); @@ -579,10 +579,10 @@ index 605e84b..686a02c 100644 auto R = std::next(MI->getReverseIterator()); diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp -index be0424d..82472a4 100644 +index 2cd2e3035c38..a1459b00fc31 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp -@@ -19862,6 +19862,12 @@ void SITargetLowering::AdjustInstrPostInstrSelection(MachineInstr &MI, +@@ -19871,6 +19871,12 @@ void SITargetLowering::AdjustInstrPostInstrSelection(MachineInstr &MI, MachineFunction *MF = MI.getMF(); MachineRegisterInfo &MRI = MF->getRegInfo(); @@ -596,10 +596,10 @@ index be0424d..82472a4 100644 // Make sure constant bus requirements are respected. TII->legalizeOperandsVOP3(MRI, MI); diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp -index ad7073f..e0b4225 100644 +index 72cf6f7ec4d3..372c67c873c5 100644 --- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp -@@ -1800,10 +1800,11 @@ void SIInstrInfo::storeRegToStackSlotImpl( +@@ -1798,10 +1798,11 @@ void SIInstrInfo::storeRegToStackSlotImpl( } BuildMI(MBB, MI, DL, OpDesc) @@ -615,7 +615,7 @@ index ad7073f..e0b4225 100644 return; } -@@ -1813,11 +1814,12 @@ void SIInstrInfo::storeRegToStackSlotImpl( +@@ -1811,11 +1812,12 @@ void SIInstrInfo::storeRegToStackSlotImpl( MFI->setHasSpilledVGPRs(); BuildMI(MBB, MI, DL, get(Opcode)) @@ -633,7 +633,7 @@ index ad7073f..e0b4225 100644 } void SIInstrInfo::storeRegToStackSlot( -@@ -2008,9 +2010,10 @@ void SIInstrInfo::loadRegFromStackSlot(MachineBasicBlock &MBB, +@@ -2006,9 +2008,10 @@ void SIInstrInfo::loadRegFromStackSlot(MachineBasicBlock &MBB, } BuildMI(MBB, MI, DL, OpDesc, DestReg) @@ -647,7 +647,7 @@ index ad7073f..e0b4225 100644 return; } -@@ -2021,7 +2024,8 @@ void SIInstrInfo::loadRegFromStackSlot(MachineBasicBlock &MBB, +@@ -2019,7 +2022,8 @@ void SIInstrInfo::loadRegFromStackSlot(MachineBasicBlock &MBB, .addFrameIndex(FrameIndex) // vaddr .addReg(MFI->getStackPtrOffsetReg()) // scratch_offset .addImm(0) // offset @@ -657,7 +657,7 @@ index ad7073f..e0b4225 100644 } void SIInstrInfo::insertNoop(MachineBasicBlock &MBB, -@@ -10561,51 +10565,6 @@ unsigned SIInstrInfo::getLiveRangeSplitOpcode(Register SrcReg, +@@ -10518,51 +10522,6 @@ unsigned SIInstrInfo::getLiveRangeSplitOpcode(Register SrcReg, return AMDGPU::COPY; } @@ -710,10 +710,10 @@ index ad7073f..e0b4225 100644 SIInstrInfo::getAddNoCarry(MachineBasicBlock &MBB, MachineBasicBlock::iterator I, diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h -index dd4133d..914b85d 100644 +index f589d0299442..a843c362244d 100644 --- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h +++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h -@@ -924,6 +924,14 @@ public: +@@ -922,6 +922,14 @@ public: static bool isSpill(const MachineInstr &MI) { return isSpill(MI.getDesc()); } @@ -728,7 +728,7 @@ index dd4133d..914b85d 100644 static bool isWWMRegSpillOpcode(uint32_t Opcode) { return Opcode == AMDGPU::SI_SPILL_WWM_V32_SAVE || Opcode == AMDGPU::SI_SPILL_WWM_AV32_SAVE || -@@ -1737,11 +1745,6 @@ public: +@@ -1722,11 +1730,6 @@ public: unsigned getLiveRangeSplitOpcode(Register Reg, const MachineFunction &MF) const override; @@ -741,7 +741,7 @@ index dd4133d..914b85d 100644 MachineBasicBlock::iterator InsPt, const DebugLoc &DL, Register Src, diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td -index c797698..6e3cdd9 100644 +index ba02d0192a7a..250613537385 100644 --- a/llvm/lib/Target/AMDGPU/SIInstructions.td +++ b/llvm/lib/Target/AMDGPU/SIInstructions.td @@ -630,6 +630,7 @@ def SI_LOOP : CFPseudoInstSI < @@ -753,18 +753,18 @@ index c797698..6e3cdd9 100644 let isReMaterializable = 1; let hasSideEffects = 1; diff --git a/llvm/lib/Target/AMDGPU/SILowerControlFlow.cpp b/llvm/lib/Target/AMDGPU/SILowerControlFlow.cpp -index 8d7bd32..4e841a7 100644 +index c54ac39b3444..3c59b3dcbad9 100644 --- a/llvm/lib/Target/AMDGPU/SILowerControlFlow.cpp +++ b/llvm/lib/Target/AMDGPU/SILowerControlFlow.cpp -@@ -324,6 +324,7 @@ void SILowerControlFlow::emitElse(MachineInstr &MI) { +@@ -315,6 +315,7 @@ void SILowerControlFlow::emitElse(MachineInstr &MI) { BuildMI(MBB, Start, DL, TII->get(LMC.OrSaveExecOpc), SaveReg) .add(MI.getOperand(1)); // Saved EXEC setImpSCCDefDead(*OrSaveExec, /*IsDead=*/true); + OrSaveExec->setFlag(MachineInstr::BBProlog); - if (LV) - LV->replaceKillInstruction(SrcReg, MI, *OrSaveExec); -@@ -530,6 +531,10 @@ MachineBasicBlock *SILowerControlFlow::emitEndCf(MachineInstr &MI) { + MachineBasicBlock *DestBB = MI.getOperand(2).getMBB(); + +@@ -509,6 +510,10 @@ MachineBasicBlock *SILowerControlFlow::emitEndCf(MachineInstr &MI) { .addReg(LMC.ExecReg) .add(MI.getOperand(0)); copySCCDefDead(*NewMI, MI.getOperand(2)); @@ -773,10 +773,10 @@ index 8d7bd32..4e841a7 100644 + else + NewMI->setFlag(MachineInstr::BBProlog); - if (LV) { - LV->replaceKillInstruction(DataReg, MI, *NewMI); + LoweredEndCf.insert(NewMI); + diff --git a/llvm/lib/Target/AMDGPU/SILowerSGPRSpills.cpp b/llvm/lib/Target/AMDGPU/SILowerSGPRSpills.cpp -index d0a57e1..7f61cb4 100644 +index d0a57e163d29..7f61cb4708d5 100644 --- a/llvm/lib/Target/AMDGPU/SILowerSGPRSpills.cpp +++ b/llvm/lib/Target/AMDGPU/SILowerSGPRSpills.cpp @@ -132,6 +132,8 @@ static void insertCSRSaves(const GCNSubtarget &ST, MachineBasicBlock &SaveBlock, @@ -799,10 +799,10 @@ index d0a57e1..7f61cb4 100644 // Add WWM flag to the virtual register. FuncInfo->setFlag(Reg, AMDGPU::VirtRegFlag::WWM_REG); diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp -index ced2628..9f13c03 100644 +index e755e953528c..6c37ea1dcb88 100644 --- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp -@@ -2172,6 +2172,7 @@ bool SIRegisterInfo::spillSGPR(MachineBasicBlock::iterator MI, int Index, +@@ -2168,6 +2168,7 @@ bool SIRegisterInfo::spillSGPR(MachineBasicBlock::iterator MI, int Index, "undef spill should have been deleted earlier"); SGPRSpillBuilder SB(*this, *ST.getInstrInfo(), isWave32, MI, Index, RS); @@ -810,7 +810,7 @@ index ced2628..9f13c03 100644 ArrayRef VGPRSpills = SpillToPhysVGPRLane ? SB.MFI.getSGPRSpillToPhysicalVGPRLanes(Index) -@@ -2316,6 +2317,9 @@ bool SIRegisterInfo::spillSGPR(MachineBasicBlock::iterator MI, int Index, +@@ -2312,6 +2313,9 @@ bool SIRegisterInfo::spillSGPR(MachineBasicBlock::iterator MI, int Index, SB.restore(); } @@ -820,7 +820,7 @@ index ced2628..9f13c03 100644 MI->eraseFromParent(); SB.MFI.addToSpilledSGPRs(SB.NumSubRegs); -@@ -2330,6 +2334,7 @@ bool SIRegisterInfo::restoreSGPR(MachineBasicBlock::iterator MI, int Index, +@@ -2326,6 +2330,7 @@ bool SIRegisterInfo::restoreSGPR(MachineBasicBlock::iterator MI, int Index, LiveIntervals *LIS, bool OnlyToVGPR, bool SpillToPhysVGPRLane) const { SGPRSpillBuilder SB(*this, *ST.getInstrInfo(), isWave32, MI, Index, RS); @@ -828,7 +828,7 @@ index ced2628..9f13c03 100644 ArrayRef VGPRSpills = SpillToPhysVGPRLane ? SB.MFI.getSGPRSpillToPhysicalVGPRLanes(Index) -@@ -2397,6 +2402,9 @@ bool SIRegisterInfo::restoreSGPR(MachineBasicBlock::iterator MI, int Index, +@@ -2393,6 +2398,9 @@ bool SIRegisterInfo::restoreSGPR(MachineBasicBlock::iterator MI, int Index, SB.restore(); } @@ -839,10 +839,10 @@ index ced2628..9f13c03 100644 if (LIS) diff --git a/llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp b/llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp -index 3a495eb..976f5e5 100644 +index 5c41a2a563e0..d05adf077f4a 100644 --- a/llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp +++ b/llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp -@@ -1238,6 +1238,7 @@ void SIWholeQuadMode::toExact(MachineBasicBlock &MBB, +@@ -1234,6 +1234,7 @@ void SIWholeQuadMode::toExact(MachineBasicBlock &MBB, .setOperandDead(3); } @@ -850,7 +850,7 @@ index 3a495eb..976f5e5 100644 LIS->InsertMachineInstrInMaps(*MI); LIS->removeAllRegUnitsForPhysReg(AMDGPU::EXEC); StateTransition[MI] = StateExact; -@@ -1258,6 +1259,7 @@ void SIWholeQuadMode::toWQM(MachineBasicBlock &MBB, +@@ -1254,6 +1255,7 @@ void SIWholeQuadMode::toWQM(MachineBasicBlock &MBB, .setOperandDead(2); } @@ -858,7 +858,7 @@ index 3a495eb..976f5e5 100644 LIS->InsertMachineInstrInMaps(*MI); StateTransition[MI] = StateWQM; } -@@ -1281,6 +1283,7 @@ void SIWholeQuadMode::toStrictMode(MachineBasicBlock &MBB, +@@ -1277,6 +1279,7 @@ void SIWholeQuadMode::toStrictMode(MachineBasicBlock &MBB, .addImm(-1) .setOperandDead(3); } @@ -866,7 +866,7 @@ index 3a495eb..976f5e5 100644 LIS->InsertMachineInstrInMaps(*MI); StateTransition[MI] = StrictStateNeeded; } -@@ -1306,6 +1309,7 @@ void SIWholeQuadMode::fromStrictMode(MachineBasicBlock &MBB, +@@ -1302,6 +1305,7 @@ void SIWholeQuadMode::fromStrictMode(MachineBasicBlock &MBB, BuildMI(MBB, Before, DL, TII->get(AMDGPU::EXIT_STRICT_WQM), LMC.ExecReg) .addReg(SavedOrig); } @@ -874,7 +874,7 @@ index 3a495eb..976f5e5 100644 LIS->InsertMachineInstrInMaps(*MI); StateTransition[MI] = NonStrictState; } -@@ -1626,7 +1630,8 @@ void SIWholeQuadMode::lowerInitExec(MachineInstr &MI) { +@@ -1622,7 +1626,8 @@ void SIWholeQuadMode::lowerInitExec(MachineInstr &MI) { MachineInstr *SaveExec = BuildMI(*MBB, MBB->begin(), MI.getDebugLoc(), TII->get(LMC.OrSaveExecOpc), EntryExec) .addImm(-1) @@ -884,7 +884,7 @@ index 3a495eb..976f5e5 100644 // Replace all uses of MI's destination reg with EntryExec. MRI->replaceRegWith(MI.getOperand(0).getReg(), EntryExec); -@@ -1648,7 +1653,8 @@ void SIWholeQuadMode::lowerInitExec(MachineInstr &MI) { +@@ -1644,7 +1649,8 @@ void SIWholeQuadMode::lowerInitExec(MachineInstr &MI) { // This should be before all vector instructions. MachineInstr *InitMI = BuildMI(*MBB, MBB->begin(), MI.getDebugLoc(), TII->get(LMC.MovOpc), LMC.ExecReg) @@ -894,7 +894,7 @@ index 3a495eb..976f5e5 100644 if (LIS) { LIS->RemoveMachineInstrFromMaps(MI); LIS->InsertMachineInstrInMaps(*InitMI); -@@ -1681,6 +1687,8 @@ void SIWholeQuadMode::lowerInitExec(MachineInstr &MI) { +@@ -1677,6 +1683,8 @@ void SIWholeQuadMode::lowerInitExec(MachineInstr &MI) { // If first instruction is definition then move pointer after it. FirstMI = &*std::next(FirstMI->getIterator()); } @@ -903,7 +903,7 @@ index 3a495eb..976f5e5 100644 } } -@@ -1701,6 +1709,8 @@ void SIWholeQuadMode::lowerInitExec(MachineInstr &MI) { +@@ -1697,6 +1705,8 @@ void SIWholeQuadMode::lowerInitExec(MachineInstr &MI) { .addImm(WavefrontSize); auto CmovMI = BuildMI(*MBB, FirstMI, DL, TII->get(LMC.CMovOpc), LMC.ExecReg).addImm(-1); @@ -912,7 +912,7 @@ index 3a495eb..976f5e5 100644 if (!LIS) { MI.eraseFromParent(); -@@ -1773,6 +1783,7 @@ bool SIWholeQuadMode::run(MachineFunction &MF) { +@@ -1769,6 +1779,7 @@ bool SIWholeQuadMode::run(MachineFunction &MF) { MachineInstr *MI = BuildMI(Entry, EntryMI, DebugLoc(), TII->get(AMDGPU::COPY), LiveMaskReg) .addReg(LMC.ExecReg); @@ -920,7 +920,7 @@ index 3a495eb..976f5e5 100644 LIS->InsertMachineInstrInMaps(*MI); Changed = true; } -@@ -1807,6 +1818,7 @@ bool SIWholeQuadMode::run(MachineFunction &MF) { +@@ -1803,6 +1814,7 @@ bool SIWholeQuadMode::run(MachineFunction &MF) { BuildMI(Entry, EntryMI, DebugLoc(), TII->get(LMC.WQMOpc), LMC.ExecReg) .addReg(LMC.ExecReg) .setOperandDead(2); diff --git a/patches/xla.patch b/patches/xla.patch index 2f5e11210b..95627dde1f 100644 --- a/patches/xla.patch +++ b/patches/xla.patch @@ -1,9 +1,9 @@ diff --git a/third_party/gpus/rocm/BUILD.tpl b/third_party/gpus/rocm/BUILD.tpl -index 6e05471a99..c29d37acd3 100644 +index 09432ba..c2bf315 100644 --- a/third_party/gpus/rocm/BUILD.tpl +++ b/third_party/gpus/rocm/BUILD.tpl -@@ -468,6 +468,36 @@ cc_library( - data = [":system_libs_data"], +@@ -570,6 +570,36 @@ rocm_lib_import( + deps = [":system_libs"], ) +sh_binary( @@ -40,10 +40,10 @@ index 6e05471a99..c29d37acd3 100644 name = "toolchain_data", srcs = glob( diff --git a/third_party/rocm_device_libs/build_defs.bzl b/third_party/rocm_device_libs/build_defs.bzl -index 5b9aee4f08..f55aabe5ad 100644 +index 49c0119..bd4ea0c 100644 --- a/third_party/rocm_device_libs/build_defs.bzl +++ b/third_party/rocm_device_libs/build_defs.bzl -@@ -111,6 +111,7 @@ def _bitcode_library_impl(ctx): +@@ -124,6 +124,7 @@ def _bitcode_library_impl(ctx): inputs = [strip_out], outputs = [final_bc], arguments = [strip_out.path, "-o", final_bc.path], @@ -51,7 +51,7 @@ index 5b9aee4f08..f55aabe5ad 100644 progress_message = "Preparing final bitcode for {}".format(ctx.label.name), mnemonic = "BitcodeFinalize", ) -@@ -127,17 +128,17 @@ bitcode_library = rule( +@@ -139,17 +140,17 @@ bitcode_library = rule( "hdrs": attr.label_list(allow_files = [".h"]), "file_specific_flags": attr.string_list_dict(), "_clang": attr.label( @@ -73,10 +73,10 @@ index 5b9aee4f08..f55aabe5ad 100644 cfg = "exec", ), diff --git a/third_party/rocm_device_libs/rocm_device_libs.BUILD b/third_party/rocm_device_libs/rocm_device_libs.BUILD -index ee162b63ab..98a55618a4 100644 +index 1709a6b..8994c5f 100644 --- a/third_party/rocm_device_libs/rocm_device_libs.BUILD +++ b/third_party/rocm_device_libs/rocm_device_libs.BUILD -@@ -28,36 +28,12 @@ cc_binary( +@@ -43,36 +43,12 @@ cc_binary( ], ) @@ -117,20 +117,3 @@ index ee162b63ab..98a55618a4 100644 - }, + actual = "@local_config_rocm//rocm:ocml" ) - -diff --git a/xla/pjrt/distributed/topology_util.cc b/xla/pjrt/distributed/topology_util.cc ---- a/xla/pjrt/distributed/topology_util.cc -+++ b/xla/pjrt/distributed/topology_util.cc -@@ -307,8 +307,10 @@ absl::Status ExchangeTopologies(absl::string_view platform, int node_id, - if (num_nodes == 1) { - LocalTopologyProto* topology = global_topology->add_processes(); - *topology = local_topology; -- for (DeviceProto& device : *topology->mutable_devices()) { -- device.set_global_device_id(device.local_device_ordinal()); -+ if (assign_global_device_ids) { -+ for (int i = 0; i < topology->devices_size(); ++i) { -+ topology->mutable_devices(i)->set_global_device_id(i); -+ } - } - return absl::OkStatus(); - } diff --git a/src/enzyme_ad/jax/Passes/ConvertPolygeistToLLVM.cpp b/src/enzyme_ad/jax/Passes/ConvertPolygeistToLLVM.cpp index aa7745ead9..3e38bfe0cb 100644 --- a/src/enzyme_ad/jax/Passes/ConvertPolygeistToLLVM.cpp +++ b/src/enzyme_ad/jax/Passes/ConvertPolygeistToLLVM.cpp @@ -4525,7 +4525,7 @@ populateCStyleGPUFuncLoweringPatterns(RewritePatternSet &patterns, } else if (gpuTarget == "rocm") { mlir::populateGpuToROCDLConversionPatterns(typeConverter, patterns, mlir::gpu::amd::Runtime::HIP, - amdgpu::Chipset()); + ROCDL::TargetInfo()); } } } diff --git a/src/enzyme_ad/jax/Passes/EnzymeHLOOpt.cpp b/src/enzyme_ad/jax/Passes/EnzymeHLOOpt.cpp index 4817a0c0f6..6d93c5d72b 100644 --- a/src/enzyme_ad/jax/Passes/EnzymeHLOOpt.cpp +++ b/src/enzyme_ad/jax/Passes/EnzymeHLOOpt.cpp @@ -14631,7 +14631,7 @@ struct CompareOpCanon final ? stablehlo::makeTensor(rhs.resizeSplat(RankedTensorType::get( {}, rhs.getType().getElementType()))) : stablehlo::constantOp(rhs), - op.getComparisonDirection(), ty)); + op.getComparisonDirection(), op.getCompareType(), ty)); if (isSplat) out = out.resizeSplat(op.getType()); diff --git a/src/external/polly/include/polly/CodeGen/IRBuilder.h b/src/external/polly/include/polly/CodeGen/IRBuilder.h index 83b8e885ea..56ad16726f 100644 --- a/src/external/polly/include/polly/CodeGen/IRBuilder.h +++ b/src/external/polly/include/polly/CodeGen/IRBuilder.h @@ -123,7 +123,7 @@ class IRInserter final : public llvm::IRBuilderDefaultInserter { IRInserter(ScopAnnotator &A) : Annotator(&A) {} void InsertHelper(llvm::Instruction *I, const llvm::Twine &Name, - llvm::BasicBlock::iterator InsertPt) const override { + llvm::BasicBlock::iterator InsertPt) const { llvm::IRBuilderDefaultInserter::InsertHelper(I, Name, InsertPt); if (Annotator) Annotator->annotate(I); diff --git a/test/lit_tests/raising/raiseimmerse3.mlir b/test/lit_tests/raising/raiseimmerse3.mlir index 7edddf676e..a750a7e4b7 100644 --- a/test/lit_tests/raising/raiseimmerse3.mlir +++ b/test/lit_tests/raising/raiseimmerse3.mlir @@ -27,15 +27,19 @@ module { } // CHECK: func.func private @"##call__Z25gpu__mask_immersed_field_16CompilerMetadataI10StaticSizeI13_180__85__20_E12DynamicCheckvv7NDRangeILi3ES0_I11_12__6__20_ES0_I11_16__16__1_EvvEE11OffsetArrayI7Float64Li3E13CuTracedArrayIS9_Li3ELi1E13_194__99__34_EE5TupleI6Center4FaceSE_E20ImmersedBoundaryGridIS9_8Periodic14RightConnected7Bounded28OrthogonalSphericalShellGridIS9_SI_SJ_SK_28StaticVerticalDiscretizationIS8_IS9_Li1ESA_IS9_Li1ELi1E5_35__EES8_IS9_Li1ESA_IS9_Li1ELi1E5_34__EESO_SQ_ES8_IS9_Li2ESA_IS9_Li2ELi1E9_194__99_EE8TripolarI5Int64SV_SV_EvE16GridFittedBottomI5FieldISE_SE_vvvvS8_IS9_Li3ESA_IS9_Li3ELi1E12_194__99__1_EES9_vvvE23CenterImmersedConditionEvvvES9_#1335$par75"(%arg0: memref<34x99x194xf64, 1>, %arg1: memref<34xf64, 1>, %arg2: memref<1x99x194xf64, 1>, %arg3: i1) { -// CHECK-NEXT: %c0 = arith.constant 0 : index +// CHECK-NEXT: %false = arith.constant false +// CHECK-NEXT: %true = arith.constant true // CHECK-NEXT: %cst = arith.constant 0.000000e+00 : f64 // CHECK-NEXT: %cst_0 = arith.constant 2.000000e+00 : f64 // CHECK-NEXT: affine.parallel (%arg4) = (0) to (85) { -// CHECK-NEXT: %0 = arith.cmpi eq, %arg4, %c0 : index +// CHECK-NEXT: %0 = affine.if #set(%arg4) -> i1 { +// CHECK-NEXT: affine.yield %true : i1 +// CHECK-NEXT: } else { +// CHECK-NEXT: affine.yield %false : i1 +// CHECK-NEXT: } // CHECK-NEXT: %1 = arith.ori %0, %arg3 : i1 -// CHECK-NEXT: %2 = arith.ori %arg3, %1 : i1 -// CHECK-NEXT: %3 = arith.select %2, %cst, %cst_0 : f64 -// CHECK-NEXT: affine.store %3, %arg0[7, %arg4 + 7, 7] : memref<34x99x194xf64, 1> +// CHECK-NEXT: %2 = arith.select %1, %cst, %cst_0 : f64 +// CHECK-NEXT: affine.store %2, %arg0[7, %arg4 + 7, 7] : memref<34x99x194xf64, 1> // CHECK-NEXT: } // CHECK-NEXT: return // CHECK-NEXT: } diff --git a/workspace.bzl b/workspace.bzl index 5286c37ecf..ceab05bbb1 100644 --- a/workspace.bzl +++ b/workspace.bzl @@ -1,4 +1,4 @@ -JAX_COMMIT = "fbf6588d55c5662ccf98558015d5b2c4c4a7bf6f" +JAX_COMMIT = "6aaf234b1ee98800fcb80b9014a052a66dadbd8a" JAX_SHA256 = "" ENZYME_COMMIT = "56b02e7e94265e7af38fe95e9fbf9818ef1d09c0" @@ -70,10 +70,10 @@ echo " InterpreterFallback* fallback) {" >> third_party/stablehlo/temporary. echo "+#if (defined(_WIN32) || defined(__CYGWIN__))" >> third_party/stablehlo/temporary.patch echo "+ llvm::report_fatal_error(\\"Op not supported on windows due to std::future\\");" >> third_party/stablehlo/temporary.patch echo "+#else" >> third_party/stablehlo/temporary.patch -echo " llvm::DefaultThreadPool threadPool;" >> third_party/stablehlo/temporary.patch -echo " llvm::ThreadPoolTaskGroup taskGroup(threadPool);" >> third_party/stablehlo/temporary.patch +echo " uint32_t numReplicas = programs.size();" >> third_party/stablehlo/temporary.patch +echo " uint32_t numPartitions = programs[0].size();" >> third_party/stablehlo/temporary.patch echo " " >> third_party/stablehlo/temporary.patch -echo "@@ -211,6 +214,7 @@ SmallVector evalRunParallelOp(" >> third_party/stablehlo/temporary.patch +echo "@@ -223,6 +226,7 @@ SmallVector evalRunParallelOp(" >> third_party/stablehlo/temporary.patch echo " for (auto& output : taskOutputs) results.append(output);" >> third_party/stablehlo/temporary.patch echo " // TODO(#1725): Figure out how to test the outfeed queue." >> third_party/stablehlo/temporary.patch echo " return results;" >> third_party/stablehlo/temporary.patch @@ -105,7 +105,7 @@ echo " llvm::Error evalPrintOp(PrintOp& op, InterpreterValue operand) {" >> thir # place copies and spills above the exec restore of an if/else join, # miscompiling kernels (llvm/llvm-project#222368). Drop once XLA's LLVM # includes it. - sed -i.bak0 "s/llvm:generated.patch\\\",/llvm:generated.patch\\\", \\\"\\/\\/:patches\\/llvm_amdgpu_bb_prolog.patch\\\", \\\"\\/\\/:patches\\/llvm_mlir_import_inrange_width.patch\\\", \\\"\\/\\/:patches\\/llvm_orc_unw_revert.patch\\\", \\\"\\/\\/:patches\\/llvm_affine_parallel_signless_minmax.patch\\\",/g" third_party/llvm/workspace.bzl + sed -i.bak0 "s/llvm:generated.patch\\\",/llvm:generated.patch\\\", \\\"\\/\\/:patches\\/llvm_amdgpu_bb_prolog.patch\\\", \\\"\\/\\/:patches\\/llvm_affine_parallel_signless_minmax.patch\\\",/g" third_party/llvm/workspace.bzl """, """ sed -i.bak0 "s/tf_http_archive/http_archive/g" third_party/llvm/workspace.bzl