[VPlan] Support to vectorize inner loops with VPlan native path enabled

[Kazhuu]

Summary

When no explicit loop is marked to be vectorized and VPlan native path is
enabled. The innermost loop will be vectorized using the inner loop vectorizer.
Previously this setup caused errors like this to happen:
https://bugs.llvm.org/show_bug.cgi?id=42592.

This patch fixes that and add a test case to test inner loop vectorization
happens when VPlan native path is enabled.

[fhahn]

Thanks for the patch!

llvm/test/Transforms/LoopVectorize/vplan-vectorize-inner-loop.ll

5 ↗

(On Diff #326046)

Can you add a reference to the bug report (PR42592)?

62 ↗

(On Diff #326046)

I think the outer loop could be simplified and only contain the bare minimum (outer induction variable & checks)

70 ↗

(On Diff #326046)

Do we need a reduction here? Might be simpler to just have a simple store instead?

[Kazhuu]

Add bare minimum test case that I was able to use to reproduce the issues and it passes when the patch is applied. Also add one more assert fix that I noticed came up when testing this.

[fhahn]

llvm/lib/Transforms/Vectorize/LoopVectorize.cpp

4447–4449

I think we should probably have a UseVPlanNativePath variable in ILV, which is true if EnableVPlanNativePath && !OrigLoop->isInnerMost() and replace all checks of EnableVPlanNativePath with checking UseVPlanNativePath.

5188

Why is this needed? If we handle inner loops with the regular ILV even in the VPlanNativePath, the assert should not change I think. Is it possible that we need to update other parts to set the right decision for ILV in VPlanNativePath?

[Kazhuu]

llvm/lib/Transforms/Vectorize/LoopVectorize.cpp

5188

I needed to add this when inner loop contained load instruction that is uniform on all iterations . This load is also in the test case I added in the loop body:

%d = load float, float* %arrayidx5.i, align 4

I guess technically LICM should lift this load up from the loop but I experienced this error when using PoCL with CPU backend and polybench OpenCL testbech using normal -O3 pass pipeline. I guess I could post it as a separate issue as well if needed. The assert line reported to fail was this:

opt: /home/kazooie/extra/programming/llvm-project/llvm/lib/Transforms/Vectorize/LoopVectorize.cpp:5184: llvm::LoopVectorizationCostModel::collectLoopUniforms(llvm::ElementCount)::<lambda(llvm::Instruction*, llvm::ElementCount)>: Assertion `WideningDecision == CM_Scalarize' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace.

And to my knowledge where scatter/gather widening decision is coming from is form the fact that VPlan doesn't support recognizing uniform memory strides, so it is using scatter/gather instead of loads.

[Kazhuu]

Add UseVPlanNative path boolean to inform when to use VPlan native path, removed unnecessary assert modification from earlier and update the test case. Now vectorizing inner loop produces the same vectorized code when VPlan native path flag is enabled and not enabled.

[Kazhuu]

llvm/lib/Transforms/Vectorize/LoopVectorize.cpp

4447–4449

Added the boolean to LoopVectorizationCostModel class instead because the flags were used there as well and to avoid having the same flag in both ILV and LoopVectorizationCostModel classes.

5188

Yeah you were right this was unnecessary indeed when the VPlan flag usage is fixed.

[a.elovikov]

llvm/lib/Transforms/Vectorize/LoopVectorize.cpp

4447–4449

For my education, why shouldn't we use VPlan native path for inner loop vectorization? Isn't it a simpler way to make the path stable enough (it's already covered by an options that is disabled by default and is served for development purposes only)? I'd expect inner loop vectorization should be a subset of what it should be able to do.

In other words, do we expect it to be a temporary workaround to enable more work on VPlan native path right now, or do I miss some design decision that would explain why this should be a long term fix?

To clarify - I'm not going to block the review/request changes to the patch because of this. It is solely for the purpose of me getting a better understanding of the ongoing VPlan development and future plans.

[Kazhuu]

llvm/lib/Transforms/Vectorize/LoopVectorize.cpp

4447–4449

I wanted to fix this because of the crashes it causes which in my opinion are not the good thing have. To be honest I'm not aware of the long term plan that well either. Maybe @Florian or someone who knows the plan better can shed some light on this?

[Kazhuu]

llvm/lib/Transforms/Vectorize/LoopVectorize.cpp

4447–4449

Oh no sorry. I meant @fhahn :D

[fhahn]

llvm/lib/Transforms/Vectorize/LoopVectorize.cpp

1837

I'd not specify the conditions when this is set true. I'd say something like Controls whether the VPlan native path is used or not.

nit: use docygen comment ///

4447–4449

In other words, do we expect it to be a temporary workaround to enable more work on VPlan native path right now, or do I miss some design decision that would explain why this should be a long term fix?

Unfortunately the VPlan native path is extremely constrained at the moment, exclusively relying on user annotation for legality checks. Most recent work has been in the other direction, VPlanization of the inner loop vectorizer, so in a way the term 'VPlan native path' is becoming more inaccurate and in the medium term the 'plan native path' will be more like 'VPlan outer loop vectorization'.

In terms of user experience, I think the behavior of the patch is much more user-friendly: currently the VPlan native path also runs on inner loops without requiring any pragmas, causing crashes as outlined in the bug-report. With the current patch, inner loops should go through the existing inner loop vectorizer and only outer loops explicitly marked for vectorization will go through the VPlan native path. This should make it more use-able for people who want to experiment with it on larger code-bases, where they only want to use it for certain loops.

I don't think we are losing much, as we are *very* far off from the VPlan native path being ready and as I said work is underway VPlanizing the inner loop vectorizer towards unifying the paths. This also means the focus of the VPlan native path can remain brining up pieces specifically needed for outer loops.

[Kazhuu]

Fix coment.

[Kazhuu]

ping

[Kazhuu]

ping

llvm/lib/Transforms/Vectorize/LoopVectorize.cpp

LoopVectorizationCostModel(ScalarEpilogueLowering SEL, Loop *L,
                              const TargetTransformInfo &TTI,
                              const TargetLibraryInfo *TLI, DemandedBits *DB,
                              AssumptionCache *AC,
                              OptimizationRemarkEmitter *ORE, const Function *F,
                              const LoopVectorizeHints *Hints,
                              InterleavedAccessInfo &IAI)
       : ScalarEpilogueStatus(SEL), TheLoop(L), PSE(PSE), LI(LI), Legal(Legal),
         TTI(TTI), TLI(TLI), DB(DB), AC(AC), ORE(ORE), TheFunction(F),
-        Hints(Hints), InterleaveInfo(IAI) {}
+        Hints(Hints), InterleaveInfo(IAI),
+        UseVPlanNativePath(EnableVPlanNativePath && !L->isInnermost()) {}
 
   /// \return An upper bound for the vectorization factor, or None if
   /// vectorization and interleaving should be avoided up front.
   Optional<ElementCount> computeMaxVF(ElementCount UserVF, unsigned UserIC);
 
   /// \return True if runtime checks are required for vectorization, and false
   /// otherwise.
   bool runtimeChecksRequired();
public:
   /// \returns True if it is more profitable to scalarize instruction \p I for
   /// vectorization factor \p VF.
   bool isProfitableToScalarize(Instruction *I, ElementCount VF) const {
     assert(VF.isVector() &&
            "Profitable to scalarize relevant only for VF > 1.");
 
     // Cost model is not run in the VPlan-native path - return conservative
     // result until this changes.
-    if (EnableVPlanNativePath)
+    if (UseVPlanNativePath)
       return false;
 
     auto Scalars = InstsToScalarize.find(VF);
     assert(Scalars != InstsToScalarize.end() &&
            "VF not yet analyzed for scalarization profitability");
     return Scalars->second.find(I) != Scalars->second.end();
   }
 
   /// Returns true if \p I is known to be uniform after vectorization.
   bool isUniformAfterVectorization(Instruction *I, ElementCount VF) const {
     if (VF.isScalar())
       return true;
 
     // Cost model is not run in the VPlan-native path - return conservative
     // result until this changes.
-    if (EnableVPlanNativePath)
+    if (UseVPlanNativePath)
       return false;
 
     auto UniformsPerVF = Uniforms.find(VF);
     assert(UniformsPerVF != Uniforms.end() &&
            "VF not yet analyzed for uniformity");
     return UniformsPerVF->second.count(I);
   }
 
   /// Returns true if \p I is known to be scalar after vectorization.
   bool isScalarAfterVectorization(Instruction *I, ElementCount VF) const {
     if (VF.isScalar())
       return true;
 
     // Cost model is not run in the VPlan-native path - return conservative
     // result until this changes.
-    if (EnableVPlanNativePath)
+    if (UseVPlanNativePath)
       return false;
 
     auto ScalarsPerVF = Scalars.find(VF);
     assert(ScalarsPerVF != Scalars.end() &&
            "Scalar values are not calculated for VF");
     return ScalarsPerVF->second.count(I);
   }
 
public:
 
   /// Return the cost model decision for the given instruction \p I and vector
   /// width \p VF. Return CM_Unknown if this instruction did not pass
   /// through the cost modeling.
   InstWidening getWideningDecision(Instruction *I, ElementCount VF) {
     assert(VF.isVector() && "Expected VF to be a vector VF");
     // Cost model is not run in the VPlan-native path - return conservative
     // result until this changes.
-    if (EnableVPlanNativePath)
+    if (UseVPlanNativePath)
       return CM_GatherScatter;
 
     std::pair<Instruction *, ElementCount> InstOnVF = std::make_pair(I, VF);
     auto Itr = WideningDecisions.find(InstOnVF);
     if (Itr == WideningDecisions.end())
       return CM_Unknown;
     return Itr->second.first;
   }
public:
   /// Values to ignore in the cost model.
   SmallPtrSet<const Value *, 16> ValuesToIgnore;
 
   /// Values to ignore in the cost model when VF > 1.
   SmallPtrSet<const Value *, 16> VecValuesToIgnore;
 
   /// Profitable vector factors.
   SmallVector<VectorizationFactor, 8> ProfitableVFs;
+
+  /// Controls whether the VPlan native path is used or not.
+  bool UseVPlanNativePath;
 };
 
 } // end namespace llvm
 
 // Return true if \p OuterLp is an outer loop annotated with hints for explicit
 // vectorization. The loop needs to be annotated with #pragma omp simd
 // simdlen(#) or #pragma clang vectorize(enable) vectorize_width(#). If the
 // vector length information is not provided, vectorization is not considered
ReplaceInstWithInst(
       SCEVCheckBlock->getTerminator(),
       BranchInst::Create(Bypass, LoopVectorPreHeader, SCEVCheck));
   LoopBypassBlocks.push_back(SCEVCheckBlock);
   AddedSafetyChecks = true;
 }
 
 void InnerLoopVectorizer::emitMemRuntimeChecks(Loop *L, BasicBlock *Bypass) {
   // VPlan-native path does not do any analysis for runtime checks currently.
-  if (EnableVPlanNativePath)
+  if (Cost->UseVPlanNativePath)
     return;
 
   // Reuse existing vector loop preheader for runtime memory checks.
   // Note that new preheader block is generated for vector loop.
   BasicBlock *const MemCheckBlock = L->getLoopPreheader();
 
   // Generate the code that checks in runtime if arrays overlap. We put the
   // checks into a separate block to make the more common case of few elements

 void InnerLoopVectorizer::fixVectorizedLoop(VPTransformState &State) {
   // Insert truncates and extends for any truncated instructions as hints to
   // InstCombine.
   if (VF.isVector())
     truncateToMinimalBitwidths(State);
 
   // Fix widened non-induction PHIs by setting up the PHI operands.
   if (OrigPHIsToFix.size()) {
-    assert(EnableVPlanNativePath &&
+    assert(Cost->UseVPlanNativePath &&
            "Unexpected non-induction PHIs for fixup in non VPlan-native path");
     fixNonInductionPHIs(State);
   }
 
   // At this point every instruction in the original loop is widened to a
   // vector form. Now we need to fix the recurrences in the loop. These PHI
   // nodes are currently empty because we did not want to introduce cycles.
   // This is the second stage of vectorizing recurrences.
void InnerLoopVectorizer::widenGEP(GetElementPtrInst *GEP, VPValue *VPDef,
   }
 }
 
 void InnerLoopVectorizer::widenPHIInstruction(Instruction *PN,
                                               RecurrenceDescriptor *RdxDesc,
                                               VPValue *StartVPV, VPValue *Def,
                                               VPTransformState &State) {
   PHINode *P = cast<PHINode>(PN);
-  if (EnableVPlanNativePath) {
-    // Currently we enter here in the VPlan-native path for non-induction
-    // PHIs where all control flow is uniform. We simply widen these PHIs.
+  if (Cost->UseVPlanNativePath) {
+    // We enter here in the VPlan-native path and when the loop is not the
+    // innermost loop. We handle non-induction PHIs here and simply widen them.
     // Create a vector phi with no operands - the vector phi operands will be
     // set at the end of vector code generation.
     Type *VecTy = (State.VF.isScalar())
                       ? PN->getType()
                       : VectorType::get(PN->getType(), State.VF);
     Value *VecPhi = Builder.CreatePHI(VecTy, PN->getNumOperands(), "vec.phi");
     State.set(Def, VecPhi, 0);
     OrigPHIsToFix.push_back(P);
void LoopVectorizationCostModel::collectLoopUniforms(ElementCount VF) {
   auto isUniformDecision = [&](Instruction *I, ElementCount VF) {
     InstWidening WideningDecision = getWideningDecision(I, VF);
     assert(WideningDecision != CM_Unknown &&
            "Widening decision should be ready at this moment");
 
     // A uniform memory op is itself uniform.  We exclude uniform stores
     // here as they demand the last lane, not the first one.
     if (isa<LoadInst>(I) && Legal->isUniformMemOp(*I)) {
       assert(WideningDecision == CM_Scalarize);
       return true;
     }
 
     return (WideningDecision == CM_Widen ||
             WideningDecision == CM_Widen_Reverse ||
             WideningDecision == CM_Interleave);
   };
 

llvm/lib/Transforms/Vectorize/VPlan.cpp

void VPlan::execute(VPTransformState *State) {
   }
 
   // 3. Merge the temporary latch created with the last basic-block filled.
   BasicBlock *LastBB = State->CFG.PrevBB;
   // Connect LastBB to VectorLatchBB to facilitate their merge.
   assert((EnableVPlanNativePath ||
           isa<UnreachableInst>(LastBB->getTerminator())) &&
          "Expected InnerLoop VPlan CFG to terminate with unreachable");
-  assert((!EnableVPlanNativePath || isa<BranchInst>(LastBB->getTerminator())) &&
+  assert((!EnableVPlanNativePath ||
+          (L->isInnermost() || isa<BranchInst>(LastBB->getTerminator()))) &&
          "Expected VPlan CFG to terminate with branch in NativePath");
   LastBB->getTerminator()->eraseFromParent();
   BranchInst::Create(VectorLatchBB, LastBB);
 
   // Merge LastBB with Latch.
   bool Merged = MergeBlockIntoPredecessor(VectorLatchBB, nullptr, State->LI);
   (void)Merged;
   assert(Merged && "Could not merge last basic block with latch.");

llvm/test/Transforms/LoopVectorize/vplan-vectorize-single-loop.ll

+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt -loop-vectorize -force-vector-width=4 -enable-vplan-native-path -S %s | FileCheck %s
+
+; Test that when VPlan native path is enabled and no explicit loop is marked to
+; be vectorized that innermost loop will be vectorized wihtout any issues. The
+; result of the vectorization should be the same as using inner loop
+; vectorization wihtout enabling VPlan native path flag.
+; See PR42592 (https://bugs.llvm.org/show_bug.cgi?id=42592).
+
+target triple = "x86_64-unknown-linux-gnu"
+define void @kernel(float* nocapture readonly %0, float* nocapture readonly %1, float* nocapture %2, i64 %3, i64 %4) {
+; CHECK-LABEL: @kernel(
+; CHECK-NEXT:  entry:
+; CHECK-NEXT:    [[TMP5:%.*]] = bitcast float* [[TMP2:%.*]] to i8*
+; CHECK-NEXT:    [[TMP6:%.*]] = bitcast float* [[TMP0:%.*]] to i8*
+; CHECK-NEXT:    [[MIN_ITERS_CHECK:%.*]] = icmp ult i64 [[TMP4:%.*]], 4
+; CHECK-NEXT:    br i1 [[MIN_ITERS_CHECK]], label [[SCALAR_PH:%.*]], label [[VECTOR_MEMCHECK:%.*]]
+; CHECK:       vector.memcheck:
+; CHECK-NEXT:    [[SCEVGEP:%.*]] = getelementptr float, float* [[TMP2]], i64 [[TMP4]]
+; CHECK-NEXT:    [[SCEVGEP1:%.*]] = bitcast float* [[SCEVGEP]] to i8*
+; CHECK-NEXT:    [[SCEVGEP2:%.*]] = getelementptr float, float* [[TMP0]], i64 [[TMP4]]
+; CHECK-NEXT:    [[SCEVGEP23:%.*]] = bitcast float* [[SCEVGEP2]] to i8*
+; CHECK-NEXT:    [[SCEVGEP4:%.*]] = getelementptr float, float* [[TMP1:%.*]], i64 [[TMP3:%.*]]
+; CHECK-NEXT:    [[SCEVGEP45:%.*]] = bitcast float* [[SCEVGEP4]] to i8*
+; CHECK-NEXT:    [[SCEVGEP6:%.*]] = getelementptr float, float* [[TMP1]], i64 [[TMP3]]
+; CHECK-NEXT:    [[SCEVGEP67:%.*]] = bitcast float* [[SCEVGEP6]] to i8*
+; CHECK-NEXT:    [[UGLYGEP:%.*]] = getelementptr i8, i8* [[SCEVGEP67]], i64 1
+; CHECK-NEXT:    [[BOUND0:%.*]] = icmp ult i8* [[TMP5]], [[SCEVGEP23]]
+; CHECK-NEXT:    [[BOUND1:%.*]] = icmp ult i8* [[TMP6]], [[SCEVGEP1]]
+; CHECK-NEXT:    [[FOUND_CONFLICT:%.*]] = and i1 [[BOUND0]], [[BOUND1]]
+; CHECK-NEXT:    [[BOUND08:%.*]] = icmp ult i8* [[TMP5]], [[UGLYGEP]]
+; CHECK-NEXT:    [[BOUND19:%.*]] = icmp ult i8* [[SCEVGEP45]], [[SCEVGEP1]]
+; CHECK-NEXT:    [[FOUND_CONFLICT10:%.*]] = and i1 [[BOUND08]], [[BOUND19]]
+; CHECK-NEXT:    [[CONFLICT_RDX:%.*]] = or i1 [[FOUND_CONFLICT]], [[FOUND_CONFLICT10]]
+; CHECK-NEXT:    [[MEMCHECK_CONFLICT:%.*]] = and i1 [[CONFLICT_RDX]], true
+; CHECK-NEXT:    br i1 [[MEMCHECK_CONFLICT]], label [[SCALAR_PH]], label [[VECTOR_PH:%.*]]
+; CHECK:       vector.ph:
+; CHECK-NEXT:    [[N_MOD_VF:%.*]] = urem i64 [[TMP4]], 4
+; CHECK-NEXT:    [[N_VEC:%.*]] = sub i64 [[TMP4]], [[N_MOD_VF]]
+; CHECK-NEXT:    br label [[VECTOR_BODY:%.*]]
+; CHECK:       vector.body:
+; CHECK-NEXT:    [[INDEX:%.*]] = phi i64 [ 0, [[VECTOR_PH]] ], [ [[INDEX_NEXT:%.*]], [[VECTOR_BODY]] ]
+; CHECK-NEXT:    [[VEC_IND:%.*]] = phi <4 x i64> [ <i64 0, i64 1, i64 2, i64 3>, [[VECTOR_PH]] ], [ [[VEC_IND_NEXT:%.*]], [[VECTOR_BODY]] ]
+; CHECK-NEXT:    [[TMP7:%.*]] = add i64 [[INDEX]], 0
+; CHECK-NEXT:    [[TMP8:%.*]] = add i64 [[INDEX]], 1
+; CHECK-NEXT:    [[TMP9:%.*]] = add i64 [[INDEX]], 2
+; CHECK-NEXT:    [[TMP10:%.*]] = add i64 [[INDEX]], 3
+; CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds float, float* [[TMP0]], i64 [[TMP7]]
+; CHECK-NEXT:    [[TMP12:%.*]] = getelementptr inbounds float, float* [[TMP1]], i64 [[TMP3]]
+; CHECK-NEXT:    [[TMP13:%.*]] = add nuw <4 x i64> [[VEC_IND]], <i64 2, i64 2, i64 2, i64 2>
+; CHECK-NEXT:    [[TMP14:%.*]] = getelementptr inbounds float, float* [[TMP11]], i32 0
+; CHECK-NEXT:    [[TMP15:%.*]] = bitcast float* [[TMP14]] to <4 x float>*
+; CHECK-NEXT:    [[WIDE_LOAD:%.*]] = load <4 x float>, <4 x float>* [[TMP15]], align 4, !alias.scope !0
+; CHECK-NEXT:    [[TMP16:%.*]] = load float, float* [[TMP12]], align 4, !alias.scope !3
+; CHECK-NEXT:    [[BROADCAST_SPLATINSERT:%.*]] = insertelement <4 x float> poison, float [[TMP16]], i32 0
+; CHECK-NEXT:    [[BROADCAST_SPLAT:%.*]] = shufflevector <4 x float> [[BROADCAST_SPLATINSERT]], <4 x float> poison, <4 x i32> zeroinitializer
+; CHECK-NEXT:    [[TMP17:%.*]] = fdiv <4 x float> [[WIDE_LOAD]], [[BROADCAST_SPLAT]]
+; CHECK-NEXT:    [[TMP18:%.*]] = getelementptr inbounds float, float* [[TMP2]], i64 [[TMP7]]
+; CHECK-NEXT:    [[TMP19:%.*]] = getelementptr inbounds float, float* [[TMP18]], i32 0
+; CHECK-NEXT:    [[TMP20:%.*]] = bitcast float* [[TMP19]] to <4 x float>*
+; CHECK-NEXT:    store <4 x float> [[TMP17]], <4 x float>* [[TMP20]], align 4, !alias.scope !5, !noalias !7
+; CHECK-NEXT:    [[INDEX_NEXT]] = add i64 [[INDEX]], 4
+; CHECK-NEXT:    [[VEC_IND_NEXT]] = add <4 x i64> [[VEC_IND]], <i64 4, i64 4, i64 4, i64 4>
+; CHECK-NEXT:    [[TMP21:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
+; CHECK-NEXT:    br i1 [[TMP21]], label [[MIDDLE_BLOCK:%.*]], label [[VECTOR_BODY]], [[LOOP8:!llvm.loop !.*]]
+; CHECK:       middle.block:
+; CHECK-NEXT:    [[CMP_N:%.*]] = icmp eq i64 [[TMP4]], [[N_VEC]]
+; CHECK-NEXT:    br i1 [[CMP_N]], label [[KERNEL_EXIT:%.*]], label [[SCALAR_PH]]
+; CHECK:       scalar.ph:
+; CHECK-NEXT:    [[BC_RESUME_VAL:%.*]] = phi i64 [ [[N_VEC]], [[MIDDLE_BLOCK]] ], [ 0, [[ENTRY:%.*]] ], [ 0, [[VECTOR_MEMCHECK]] ]
+; CHECK-NEXT:    br label [[FOR_BODY:%.*]]
+; CHECK:       for.body:
+; CHECK-NEXT:    [[I_0:%.*]] = phi i64 [ [[BC_RESUME_VAL]], [[SCALAR_PH]] ], [ [[R:%.*]], [[FOR_BODY]] ]
+; CHECK-NEXT:    [[ARRAYIDX_I:%.*]] = getelementptr inbounds float, float* [[TMP0]], i64 [[I_0]]
+; CHECK-NEXT:    [[ARRAYIDX5_I:%.*]] = getelementptr inbounds float, float* [[TMP1]], i64 [[TMP3]]
+; CHECK-NEXT:    [[AA:%.*]] = add nuw i64 [[I_0]], 2
+; CHECK-NEXT:    [[C:%.*]] = load float, float* [[ARRAYIDX_I]], align 4
+; CHECK-NEXT:    [[D:%.*]] = load float, float* [[ARRAYIDX5_I]], align 4
+; CHECK-NEXT:    [[DIV_I:%.*]] = fdiv float [[C]], [[D]]
+; CHECK-NEXT:    [[ARRAYIDX9_I:%.*]] = getelementptr inbounds float, float* [[TMP2]], i64 [[I_0]]
+; CHECK-NEXT:    store float [[DIV_I]], float* [[ARRAYIDX9_I]], align 4
+; CHECK-NEXT:    [[R]] = add nuw i64 [[I_0]], 1
+; CHECK-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i64 [[R]], [[TMP4]]
+; CHECK-NEXT:    br i1 [[EXITCOND_NOT]], label [[KERNEL_EXIT]], label [[FOR_BODY]], [[LOOP10:!llvm.loop !.*]]
+; CHECK:       kernel.exit:
+; CHECK-NEXT:    ret void
+;
+entry:
+  br label %for.body
+
+for.body:
+  %i.0 = phi i64 [ 0, %entry ], [ %r, %for.body ]
+  %arrayidx.i = getelementptr inbounds float, float* %0, i64 %i.0
+  %arrayidx5.i = getelementptr inbounds float, float* %1, i64 %3
+  %aa = add nuw i64 %i.0, 2
+  %c = load float, float* %arrayidx.i, align 4
+  %d = load float, float* %arrayidx5.i, align 4
+  %div.i = fdiv float %c, %d
+  %arrayidx9.i = getelementptr inbounds float, float* %2, i64 %i.0
+  store float %div.i, float* %arrayidx9.i, align 4
+  %r = add nuw i64 %i.0, 1
+  %exitcond.not = icmp eq i64 %r, %4
+  br i1 %exitcond.not, label %kernel.exit, label %for.body
+
+kernel.exit:
+  ret void
+}