[AMDGPU] Break-up large PHIs for DAGISel

[Pierre-vh]

Summary

DAGISel uses CopyToReg/CopyFromReg to lower PHI nodes. With large PHIs, this can result in poor codegen.
This is because it introduces a need to have a build_vector before copying the PHI value, and that build_vector may have many undef elements. This can cause very high register pressure and abnormal stack usage in some cases.

This scalarization/phi "break-up" can be easily tuned/disabled through CL options in case it's not beneficial for some users.
It's also only enabled for DAGIsel and GlobalISel handles PHIs much better (as it works on the whole function).

This can both scalarize (break a vector into its elements) and simplify (break a vector into smaller, more manageable subvectors) PHIs.

Fixes SWDEV-321581

[arsenm]

Also could use a switch test where the same successor phi needs to list the same predecessor twice

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

1361

A few random formatting changes

1409

With the way fallback works, we can end up using the DAG and wouldn’t have used this. That’s probably ok though

1418

Power of 2 shouldn’t matter

llvm/test/CodeGen/AMDGPU/amdgpu-codegenprepare-scalarize-large-phis.ll

612 ↗

(On Diff #496438)

Should also include some FP, 8-bit and 16-bit element cases. We’d be bette riff breaking those into 32 or 64 bit pieces than full scalarization

[Pierre-vh]

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

1361

Is it fine if I run clang-format on this file and commit the result directly to get rid of them?

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

1361

If you use git clang-format it only touches diffed lines

[Pierre-vh]

llvm/test/CodeGen/AMDGPU/amdgpu-codegenprepare-scalarize-large-phis.ll

612 ↗

(On Diff #496438)

Should I also lower the threshold for this to kick in? Else 8 bit vectors won't be interesting unless they're >32 elts

[arsenm]

llvm/test/CodeGen/AMDGPU/amdgpu-codegenprepare-scalarize-large-phis.ll

612 ↗

(On Diff #496438)

We only really have 32-bit registers. Everything wider is a bit unnatural. GlobalISel breaks everything into 32 or 64 pieces. There may be second order benefits from a compiler machinery perspective for preserving larger pieces. There’s room for experimentation

[Pierre-vh]

I kept the threshold at 256 for now because lower values make the transform kick more often and it's not always clear to me whether it's a positive or negative thing. Some tests have more instructions, some have a bit less, etc.
I think 256 is good enough to get started, we can always tweak it in a separate commit if benchmarking proves that it's useful. What do you think?

Also the following tests are currently broken:

• test_mfma_loop_32xfloat in acc-ldst.ll: We get accvgpr read/writes instead of a load to agpr directly.
  • One incoming value from a load
  • One incoming value from the call to mfma later in the same block (loop)
• test_vccnz_ifcvt_triangle256 in early-if-convert.ll: s_branch_vccnz is not matched anymore.
  • One incoming value from a load
  • One incoming value from a vector add
• test_mfma_loop_zeroinit in mfma-loop.ll : accvgpr copies inside the loop
  • One incoming value is a zeroinit
  • One incoming value is from call to mfma later in the same block (loop)

I think we need an additional heuristic but I'm not sure what. Maybe we should not transform when one of the value comes from a load (we know there won't be an undef so it should be good) or comes from the same block (loop)?
I'm also not sure if extract/insert subvector with a constant operand (in the case of zeroinit) gets folded out; do I need to special-case it?

[Pierre-vh]

Add an optimization for Constant incoming values.
This allows us to avoid emitting a extract.vector of a zeroinit/constant vector.

[arsenm]

In D143731#4122195, @Pierre-vh wrote:

I kept the threshold at 256 for now because lower values make the transform kick more often and it's not always clear to me whether it's a positive or negative thing. Some tests have more instructions, some have a bit less, etc.
I think 256 is good enough to get started, we can always tweak it in a separate commit if benchmarking proves that it's useful. What do you think?

Also the following tests are currently broken:

• test_mfma_loop_32xfloat in acc-ldst.ll: We get accvgpr read/writes instead of a load to agpr directly.

Ideally this is the kind of case RegBankSelect is supposed to solve

I think we need an additional heuristic but I'm not sure what. Maybe we should not transform when one of the value comes from a load (we know there won't be an undef so it should be good) or comes from the same block (loop)?
I'm also not sure if extract/insert subvector with a constant operand (in the case of zeroinit) gets folded out; do I need to special-case it?

Really this is a register bank select problem. Fundamentally we always want 32-bit phis. Any regression from not using 32-bit phis is theoretically a failure with some other optimization

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

1416

This must be a new intrinsic, I didn't know it existed. Usually you would do a shufflevector to get the reduced elements

1439

Don't need = {}

[Pierre-vh]

In D143731#4124715, @arsenm wrote:

In D143731#4122195, @Pierre-vh wrote:

I kept the threshold at 256 for now because lower values make the transform kick more often and it's not always clear to me whether it's a positive or negative thing. Some tests have more instructions, some have a bit less, etc.
I think 256 is good enough to get started, we can always tweak it in a separate commit if benchmarking proves that it's useful. What do you think?

Also the following tests are currently broken:

• test_mfma_loop_32xfloat in acc-ldst.ll: We get accvgpr read/writes instead of a load to agpr directly.

Ideally this is the kind of case RegBankSelect is supposed to solve

I think we need an additional heuristic but I'm not sure what. Maybe we should not transform when one of the value comes from a load (we know there won't be an undef so it should be good) or comes from the same block (loop)?
I'm also not sure if extract/insert subvector with a constant operand (in the case of zeroinit) gets folded out; do I need to special-case it?

Really this is a register bank select problem. Fundamentally we always want 32-bit phis. Any regression from not using 32-bit phis is theoretically a failure with some other optimization

In this case I think it's a SIFoldOperand issue (or something strange happening higher that blocks SIFoldOperands). It's not folding the AGPR copies away.
If PHIs aren't broken, we get the areg copy in the entry block:

%181:vreg_1024_align2 = REG_SEQUENCE (..vreg128s)
%1:vreg_1024_align2 = COPY %181:vreg_1024_align2
%243:areg_1024_align2 = COPY %1:vreg_1024_align2, implicit $exec

But when PHIs are broken, we get it in the loop closer to V_MFMA and after the PHIs. I suspect SIFoldOperand doesn't handle the PHI node to "propagate" the areg class up until it reaches the store.

foldOperand: %122.sub0:areg_1024_align2
  %34:vgpr_32 = PHI %1:vgpr_32, %bb.0, %67:vgpr_32, %bb.1

I'm not sure how easy it is to make SIFoldOperand handle PHIs like that.
Or maybe the fix should be earlier in the pipeline (so aregs are assigned better)? There might be an opportunity to fix this in FixSGPRCopies:

%119:sgpr_1024 = REG_SEQUENCE (vgpr phi nodes)...
%123:areg_1024_align2 = COPY %119:sgpr_1024

We don't seem to fold sgpr -> agpr copies like we fold sgpr -> vgpr

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

1416

If shufflevector has better handling and gets folded out for sure then it might be simpler to use it.
Do I just do a shufflevector with the same operand twice + the mask of the indexes I want?

1439

warning: missing field 'IncomingValues' initializer [-Wmissing-field-initializers]

It's needed because I use brace initializers

[Pierre-vh]

Use unary shufflevector instead of extract_vector

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

1416

You can use poison for the second operand

[Pierre-vh]

• Add names to values created by the transform (makes spotting issues with it easier)
• Add comment to early-if-convert

[Pierre-vh]

Rebase on D144099

[Pierre-vh]

Rebase

[foad]

This is because it introduces a need to have a build_vector before copying the PHI value, and that build_vector may have many undef elements. This can cause very high register pressure and abnormal stack usage in some cases.

Do you have any insight into *why* undef elements in build_vector would cause high register pressure? Is there any chance of fixing this elsewhere, by teaching other parts of the compiler to handle undef elements better?

[Pierre-vh]

In D143731#4157819, @foad wrote:

This is because it introduces a need to have a build_vector before copying the PHI value, and that build_vector may have many undef elements. This can cause very high register pressure and abnormal stack usage in some cases.

Do you have any insight into *why* undef elements in build_vector would cause high register pressure? Is there any chance of fixing this elsewhere, by teaching other parts of the compiler to handle undef elements better?

In the particular case I was looking at, we had a 32xi32 vector (from 11xF64 legalized) with 10 undef elements. As the final operation of the BB was a BUILD_VECTOR + CopyToReg, it forced 22 values to be alive until that point and did a wasteful 32x32 copy for this 11xf64 vector. With this change, the undef values can be optimized out and the intermediate values (elements) no longer need to be all alive at the same point. This means they can be reordered and scratch usage goes back down to zero.

Before this patch, I had looked at optimizing this at the DAG level but it's too much effort to do there. As Matt put it earlier, PHIs/Copies larger than 32 bits are "unnatural" on AMDGPU anyway so there's no real benefit in keep around large PHIs.

[foad]

In D143731#4157821, @Pierre-vh wrote:

In D143731#4157819, @foad wrote:

This is because it introduces a need to have a build_vector before copying the PHI value, and that build_vector may have many undef elements. This can cause very high register pressure and abnormal stack usage in some cases.

Do you have any insight into *why* undef elements in build_vector would cause high register pressure? Is there any chance of fixing this elsewhere, by teaching other parts of the compiler to handle undef elements better?

In the particular case I was looking at, we had a 32xi32 vector (from 11xF64 legalized) with 10 undef elements. As the final operation of the BB was a BUILD_VECTOR + CopyToReg, it forced 22 values to be alive until that point and did a wasteful 32x32 copy for this 11xf64 vector. With this change, the undef values can be optimized out and the intermediate values (elements) no longer need to be all alive at the same point. This means they can be reordered and scratch usage goes back down to zero.

Before this patch, I had looked at optimizing this at the DAG level but it's too much effort to do there. As Matt put it earlier, PHIs/Copies larger than 32 bits are "unnatural" on AMDGPU anyway so there's no real benefit in keep around large PHIs.

Then why can't you scalarize them all? Why do we need yet another weird threshold in the backend?

[Pierre-vh]

Scalarizing everything may have some unintended side effects so that's why I left it at 256 for now. I'll check what happens when I lower it to 32 and report back.

[Pierre-vh]

• Reduce threshold to 32.
  • Lots more test changes. I disabled the transform in some of them to (hopefully) preserve the test's intent better.
  • loop-live-out-copy-undef-subrange.ll takes the biggest regression. I think we may not need to do this transform if an incoming value is an arg? Args will be copied anyway so we don't really gain anything in those cases. Thoughts?
• Preserve debug info better

[Pierre-vh]

Remove mention of "power of two" sizes

[Pierre-vh]

Ping

[Pierre-vh]

Ping - can this land?

[Pierre-vh]

Ping

[kzhuravl]

lgtm, thanks!

[Pierre-vh]

Rebase

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

static cl::opt<bool> WidenLoads(
   cl::init(false));
 
 static cl::opt<bool> Widen16BitOps(
   "amdgpu-codegenprepare-widen-16-bit-ops",
   cl::desc("Widen uniform 16-bit instructions to 32-bit in AMDGPUCodeGenPrepare"),
   cl::ReallyHidden,
   cl::init(true));
 
+static cl::opt<bool>
+    ScalarizeLargePHIs("amdgpu-codegenprepare-break-large-phis",
+                       cl::desc("Break large PHI nodes for DAGISel"),
+                       cl::ReallyHidden, cl::init(true));
+
+// TODO: 256 is a default chosen based on experimentation - can be tweaked
+//  if further testing proves it's beneficial to increase/decrease it.
+static cl::opt<unsigned> ScalarizeLargePHIsThreshold(
+    "amdgpu-codegenprepare-break-large-phis-threshold",
+    cl::desc("Minimum type size in bits for breaking large PHI nodes"),
+    cl::ReallyHidden, cl::init(256));
+
 static cl::opt<bool> UseMul24Intrin(
   "amdgpu-codegenprepare-mul24",
   cl::desc("Introduce mul24 intrinsics in AMDGPUCodeGenPrepare"),
   cl::ReallyHidden,
   cl::init(true));
 
 // Legalize 64-bit division by using the generic IR expansion.
 static cl::opt<bool> ExpandDiv64InIR(
public:
   bool visitFDiv(BinaryOperator &I);
   bool visitXor(BinaryOperator &I);
 
   bool visitInstruction(Instruction &I) { return false; }
   bool visitBinaryOperator(BinaryOperator &I);
   bool visitLoadInst(LoadInst &I);
   bool visitICmpInst(ICmpInst &I);
   bool visitSelectInst(SelectInst &I);
+  bool visitPHINode(PHINode &I);
 
   bool visitIntrinsicInst(IntrinsicInst &I);
   bool visitBitreverseIntrinsicInst(IntrinsicInst &I);
 
   bool doInitialization(Module &M) override;
   bool runOnFunction(Function &F) override;
 
   StringRef getPassName() const override { return "AMDGPU IR optimizations"; }
if (auto *Range = WidenLoad->getMetadata(LLVMContext::MD_range)) {
       if (Lower->isNullValue()) {
         WidenLoad->setMetadata(LLVMContext::MD_range, nullptr);
       } else {
         Metadata *LowAndHigh[] = {
           ConstantAsMetadata::get(ConstantInt::get(I32Ty, Lower->getValue().zext(32))),
           // Don't make assumptions about the high bits.
           ConstantAsMetadata::get(ConstantInt::get(I32Ty, 0))
         };
 
         WidenLoad->setMetadata(LLVMContext::MD_range,
                                MDNode::get(Mod->getContext(), LowAndHigh));
       }
     }
 
     int TySize = Mod->getDataLayout().getTypeSizeInBits(I.getType());
     Type *IntNTy = Builder.getIntNTy(TySize);
     Value *ValTrunc = Builder.CreateTrunc(WidenLoad, IntNTy);
bool AMDGPUCodeGenPrepare::visitSelectInst(SelectInst &I) {
 
   if (ST->has16BitInsts() && needsPromotionToI32(I.getType()) &&
       DA->isUniform(&I))
     Changed |= promoteUniformOpToI32(I);
 
   return Changed;
 }
 
+bool AMDGPUCodeGenPrepare::visitPHINode(PHINode &I) {
+  // Break-up large, non power-of-two sized PHIs.
+  //
+  // This is only helpful for DAGISel because it doesn't handle large PHIs as
+  // well as GlobalISel. DAGISel lowers PHIs by using CopyToReg/CopyFromReg.
+  // With large, odd-sized PHIs we may end up needing many `build_vector`
+  // operations with most elements being "undef". This inhibits a lot of
+  // optimization opportunities and can result in unreasonably high register
+  // pressure and the inevitable stack spilling.
+  if (!ScalarizeLargePHIs || getCGPassBuilderOption().EnableGlobalISelOption)
+    return false;
+
+  FixedVectorType *FVT = dyn_cast<FixedVectorType>(I.getType());
+  if (!FVT || DL->getTypeSizeInBits(FVT) < ScalarizeLargePHIsThreshold)
+    return false;
+
+  struct VectorSlice {
+    Type *Ty = nullptr;
+    unsigned Idx = 0;
+    unsigned NumElts = 0;
+    std::vector<Value *> IncomingValues = {};
+    PHINode *NewPHI = nullptr;
+  };
+
+  std::vector<VectorSlice> Slices;
+
+  Type *EltTy = FVT->getElementType();
+  {
+    unsigned Idx = 0;
+    // For 8/16 bits type, don't scalarize fully but break it up into as many
+    // 32-bit slices as we can, and scalarize the tail.
+    const unsigned EltSize = DL->getTypeSizeInBits(EltTy);
+    const unsigned NumElts = FVT->getNumElements();
+    if (EltSize == 8 || EltSize == 16) {
+      const unsigned SubVecSize = (32 / EltSize);
+      Type *SubVecTy = FixedVectorType::get(EltTy, SubVecSize);
+      for (unsigned End = alignDown(NumElts, SubVecSize); Idx < End;
+           Idx += SubVecSize)
+        Slices.push_back(VectorSlice{SubVecTy, Idx, SubVecSize});
+    }
+
+    // Scalarize all remaining elements.
+    for (; Idx < NumElts; ++Idx)
+      Slices.push_back(VectorSlice{EltTy, Idx, 1});
+  }
+
+  // Break up this PHI's incoming values.
+  for (unsigned Idx = 0; Idx < I.getNumIncomingValues(); ++Idx) {
+    Value *Inc = I.getIncomingValue(Idx);
+
+    IRBuilder<> B(I.getIncomingBlock(Idx)->getTerminator());
+    unsigned NameSuffix = 0;
+    for (VectorSlice &S : Slices) {
+      const auto ValName =
+          "largephi.extractslice" + std::to_string(NameSuffix++);
+      if (S.NumElts > 1) {
+        SmallVector<int, 4> Mask;
+        for (unsigned K = S.Idx; K < (S.Idx + S.NumElts); ++K)
+          Mask.push_back(K);
+        S.IncomingValues.push_back(B.CreateShuffleVector(Inc, Mask, ValName));
+      } else
+        S.IncomingValues.push_back(B.CreateExtractElement(Inc, S.Idx, ValName));
+    }
+  }
+
+  // Now create one PHI per vector piece.
+  IRBuilder<> B(I.getParent()->getFirstNonPHI());
+  for (VectorSlice &S : Slices) {
+    S.NewPHI = B.CreatePHI(S.Ty, I.getNumIncomingValues());
+    for (auto &[Idx, BB] : enumerate(I.blocks()))
+      S.NewPHI->addIncoming(S.IncomingValues[Idx], BB);
+  }
+
+  // And replace this PHI with a vector of all the previous PHI values.
+  Value *Vec = PoisonValue::get(FVT);
+  unsigned NameSuffix = 0;
+  for (VectorSlice &S : Slices) {
+    const auto ValName = "largephi.insertslice" + std::to_string(NameSuffix++);
+    if (S.NumElts > 1)
+      Vec =
+          B.CreateInsertVector(FVT, Vec, S.NewPHI, B.getInt64(S.Idx), ValName);
+    else
+      Vec = B.CreateInsertElement(Vec, S.NewPHI, S.Idx, ValName);
+  }
+
+  I.replaceAllUsesWith(Vec);
+  I.eraseFromParent();
+  return true;
+}
+
 bool AMDGPUCodeGenPrepare::visitIntrinsicInst(IntrinsicInst &I) {
   switch (I.getIntrinsicID()) {
   case Intrinsic::bitreverse:
     return visitBitreverseIntrinsicInst(I);
   default:
     return false;
   }
 }

llvm/test/CodeGen/AMDGPU/amdgpu-codegenprepare-break-large-phis.ll

+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt -S -mtriple=amdgcn-- -amdgpu-codegenprepare -amdgpu-codegenprepare-break-large-phis-threshold=64 %s | FileCheck  %s --check-prefixes=OPT
+; RUN: opt -S -mtriple=amdgcn-- -amdgpu-codegenprepare -global-isel %s | FileCheck  %s --check-prefixes=NOOPT
+; RUN: opt -S -mtriple=amdgcn-- -amdgpu-codegenprepare -amdgpu-codegenprepare-break-large-phis=0 %s | FileCheck  %s --check-prefixes=NOOPT
+
+; NOTE: This test uses a lower threshold than the default value to increase the amount of cases we can test.
+
+define amdgpu_kernel void @phi_v5f64(<5 x double> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v5f64(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <5 x double> [[IN:%.*]], double 3.140000e+00, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = extractelement <5 x double> [[X]], i64 0
+; OPT-NEXT:    [[TMP1:%.*]] = extractelement <5 x double> [[X]], i64 1
+; OPT-NEXT:    [[TMP2:%.*]] = extractelement <5 x double> [[X]], i64 2
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <5 x double> [[X]], i64 3
+; OPT-NEXT:    [[TMP4:%.*]] = extractelement <5 x double> [[X]], i64 4
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <5 x double> [[IN]], double 9.140000e+00, i32 2
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <5 x double> [[Y]], i64 0
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <5 x double> [[Y]], i64 1
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <5 x double> [[Y]], i64 2
+; OPT-NEXT:    [[TMP8:%.*]] = extractelement <5 x double> [[Y]], i64 3
+; OPT-NEXT:    [[TMP9:%.*]] = extractelement <5 x double> [[Y]], i64 4
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP10:%.*]] = phi double [ [[TMP0]], [[THEN]] ], [ [[TMP5]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP11:%.*]] = phi double [ [[TMP1]], [[THEN]] ], [ [[TMP6]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP12:%.*]] = phi double [ [[TMP2]], [[THEN]] ], [ [[TMP7]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP13:%.*]] = phi double [ [[TMP3]], [[THEN]] ], [ [[TMP8]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP14:%.*]] = phi double [ [[TMP4]], [[THEN]] ], [ [[TMP9]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP15:%.*]] = insertelement <5 x double> poison, double [[TMP10]], i64 0
+; OPT-NEXT:    [[TMP16:%.*]] = insertelement <5 x double> [[TMP15]], double [[TMP11]], i64 1
+; OPT-NEXT:    [[TMP17:%.*]] = insertelement <5 x double> [[TMP16]], double [[TMP12]], i64 2
+; OPT-NEXT:    [[TMP18:%.*]] = insertelement <5 x double> [[TMP17]], double [[TMP13]], i64 3
+; OPT-NEXT:    [[TMP19:%.*]] = insertelement <5 x double> [[TMP18]], double [[TMP14]], i64 4
+; OPT-NEXT:    store <5 x double> [[TMP19]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v5f64(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <5 x double> [[IN:%.*]], double 3.140000e+00, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <5 x double> [[IN]], double 9.140000e+00, i32 2
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <5 x double> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <5 x double> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <5 x double> %in, double 3.14, i32 3
+  br label %finally
+else:
+  %y = insertelement <5 x double> %in, double 9.14, i32 2
+  br label %finally
+finally:
+  %val = phi <5 x double> [%x, %then], [%y, %else]
+  store <5 x double> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v7f64(<7 x double> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v7f64(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <7 x double> [[IN:%.*]], double 3.140000e+00, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = extractelement <7 x double> [[X]], i64 0
+; OPT-NEXT:    [[TMP1:%.*]] = extractelement <7 x double> [[X]], i64 1
+; OPT-NEXT:    [[TMP2:%.*]] = extractelement <7 x double> [[X]], i64 2
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <7 x double> [[X]], i64 3
+; OPT-NEXT:    [[TMP4:%.*]] = extractelement <7 x double> [[X]], i64 4
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <7 x double> [[X]], i64 5
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <7 x double> [[X]], i64 6
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <7 x double> [[IN]], double 9.140000e+00, i32 6
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <7 x double> [[Y]], i64 0
+; OPT-NEXT:    [[TMP8:%.*]] = extractelement <7 x double> [[Y]], i64 1
+; OPT-NEXT:    [[TMP9:%.*]] = extractelement <7 x double> [[Y]], i64 2
+; OPT-NEXT:    [[TMP10:%.*]] = extractelement <7 x double> [[Y]], i64 3
+; OPT-NEXT:    [[TMP11:%.*]] = extractelement <7 x double> [[Y]], i64 4
+; OPT-NEXT:    [[TMP12:%.*]] = extractelement <7 x double> [[Y]], i64 5
+; OPT-NEXT:    [[TMP13:%.*]] = extractelement <7 x double> [[Y]], i64 6
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP14:%.*]] = phi double [ [[TMP0]], [[THEN]] ], [ [[TMP7]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP15:%.*]] = phi double [ [[TMP1]], [[THEN]] ], [ [[TMP8]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP16:%.*]] = phi double [ [[TMP2]], [[THEN]] ], [ [[TMP9]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP17:%.*]] = phi double [ [[TMP3]], [[THEN]] ], [ [[TMP10]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP18:%.*]] = phi double [ [[TMP4]], [[THEN]] ], [ [[TMP11]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP19:%.*]] = phi double [ [[TMP5]], [[THEN]] ], [ [[TMP12]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP20:%.*]] = phi double [ [[TMP6]], [[THEN]] ], [ [[TMP13]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP21:%.*]] = insertelement <7 x double> poison, double [[TMP14]], i64 0
+; OPT-NEXT:    [[TMP22:%.*]] = insertelement <7 x double> [[TMP21]], double [[TMP15]], i64 1
+; OPT-NEXT:    [[TMP23:%.*]] = insertelement <7 x double> [[TMP22]], double [[TMP16]], i64 2
+; OPT-NEXT:    [[TMP24:%.*]] = insertelement <7 x double> [[TMP23]], double [[TMP17]], i64 3
+; OPT-NEXT:    [[TMP25:%.*]] = insertelement <7 x double> [[TMP24]], double [[TMP18]], i64 4
+; OPT-NEXT:    [[TMP26:%.*]] = insertelement <7 x double> [[TMP25]], double [[TMP19]], i64 5
+; OPT-NEXT:    [[TMP27:%.*]] = insertelement <7 x double> [[TMP26]], double [[TMP20]], i64 6
+; OPT-NEXT:    store <7 x double> [[TMP27]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v7f64(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <7 x double> [[IN:%.*]], double 3.140000e+00, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <7 x double> [[IN]], double 9.140000e+00, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <7 x double> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <7 x double> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <7 x double> %in, double 3.14, i32 3
+  br label %finally
+else:
+  %y = insertelement <7 x double> %in, double 9.14, i32 6
+  br label %finally
+finally:
+  %val = phi <7 x double> [%x, %then], [%y, %else]
+  store <7 x double> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v11f64(<11 x double> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v11f64(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <11 x double> [[IN:%.*]], double 3.140000e+00, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = extractelement <11 x double> [[X]], i64 0
+; OPT-NEXT:    [[TMP1:%.*]] = extractelement <11 x double> [[X]], i64 1
+; OPT-NEXT:    [[TMP2:%.*]] = extractelement <11 x double> [[X]], i64 2
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <11 x double> [[X]], i64 3
+; OPT-NEXT:    [[TMP4:%.*]] = extractelement <11 x double> [[X]], i64 4
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <11 x double> [[X]], i64 5
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <11 x double> [[X]], i64 6
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <11 x double> [[X]], i64 7
+; OPT-NEXT:    [[TMP8:%.*]] = extractelement <11 x double> [[X]], i64 8
+; OPT-NEXT:    [[TMP9:%.*]] = extractelement <11 x double> [[X]], i64 9
+; OPT-NEXT:    [[TMP10:%.*]] = extractelement <11 x double> [[X]], i64 10
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <11 x double> [[IN]], double 9.140000e+00, i32 6
+; OPT-NEXT:    [[TMP11:%.*]] = extractelement <11 x double> [[Y]], i64 0
+; OPT-NEXT:    [[TMP12:%.*]] = extractelement <11 x double> [[Y]], i64 1
+; OPT-NEXT:    [[TMP13:%.*]] = extractelement <11 x double> [[Y]], i64 2
+; OPT-NEXT:    [[TMP14:%.*]] = extractelement <11 x double> [[Y]], i64 3
+; OPT-NEXT:    [[TMP15:%.*]] = extractelement <11 x double> [[Y]], i64 4
+; OPT-NEXT:    [[TMP16:%.*]] = extractelement <11 x double> [[Y]], i64 5
+; OPT-NEXT:    [[TMP17:%.*]] = extractelement <11 x double> [[Y]], i64 6
+; OPT-NEXT:    [[TMP18:%.*]] = extractelement <11 x double> [[Y]], i64 7
+; OPT-NEXT:    [[TMP19:%.*]] = extractelement <11 x double> [[Y]], i64 8
+; OPT-NEXT:    [[TMP20:%.*]] = extractelement <11 x double> [[Y]], i64 9
+; OPT-NEXT:    [[TMP21:%.*]] = extractelement <11 x double> [[Y]], i64 10
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP22:%.*]] = phi double [ [[TMP0]], [[THEN]] ], [ [[TMP11]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP23:%.*]] = phi double [ [[TMP1]], [[THEN]] ], [ [[TMP12]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP24:%.*]] = phi double [ [[TMP2]], [[THEN]] ], [ [[TMP13]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP25:%.*]] = phi double [ [[TMP3]], [[THEN]] ], [ [[TMP14]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP26:%.*]] = phi double [ [[TMP4]], [[THEN]] ], [ [[TMP15]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP27:%.*]] = phi double [ [[TMP5]], [[THEN]] ], [ [[TMP16]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP28:%.*]] = phi double [ [[TMP6]], [[THEN]] ], [ [[TMP17]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP29:%.*]] = phi double [ [[TMP7]], [[THEN]] ], [ [[TMP18]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP30:%.*]] = phi double [ [[TMP8]], [[THEN]] ], [ [[TMP19]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP31:%.*]] = phi double [ [[TMP9]], [[THEN]] ], [ [[TMP20]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP32:%.*]] = phi double [ [[TMP10]], [[THEN]] ], [ [[TMP21]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP33:%.*]] = insertelement <11 x double> poison, double [[TMP22]], i64 0
+; OPT-NEXT:    [[TMP34:%.*]] = insertelement <11 x double> [[TMP33]], double [[TMP23]], i64 1
+; OPT-NEXT:    [[TMP35:%.*]] = insertelement <11 x double> [[TMP34]], double [[TMP24]], i64 2
+; OPT-NEXT:    [[TMP36:%.*]] = insertelement <11 x double> [[TMP35]], double [[TMP25]], i64 3
+; OPT-NEXT:    [[TMP37:%.*]] = insertelement <11 x double> [[TMP36]], double [[TMP26]], i64 4
+; OPT-NEXT:    [[TMP38:%.*]] = insertelement <11 x double> [[TMP37]], double [[TMP27]], i64 5
+; OPT-NEXT:    [[TMP39:%.*]] = insertelement <11 x double> [[TMP38]], double [[TMP28]], i64 6
+; OPT-NEXT:    [[TMP40:%.*]] = insertelement <11 x double> [[TMP39]], double [[TMP29]], i64 7
+; OPT-NEXT:    [[TMP41:%.*]] = insertelement <11 x double> [[TMP40]], double [[TMP30]], i64 8
+; OPT-NEXT:    [[TMP42:%.*]] = insertelement <11 x double> [[TMP41]], double [[TMP31]], i64 9
+; OPT-NEXT:    [[TMP43:%.*]] = insertelement <11 x double> [[TMP42]], double [[TMP32]], i64 10
+; OPT-NEXT:    store <11 x double> [[TMP43]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v11f64(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <11 x double> [[IN:%.*]], double 3.140000e+00, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <11 x double> [[IN]], double 9.140000e+00, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <11 x double> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <11 x double> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <11 x double> %in, double 3.14, i32 3
+  br label %finally
+else:
+  %y = insertelement <11 x double> %in, double 9.14, i32 6
+  br label %finally
+finally:
+  %val = phi <11 x double> [%x, %then], [%y, %else]
+  store <11 x double> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v11f64_cst(<11 x double> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v11f64_cst(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[FINALLY:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <11 x double> [[IN:%.*]], double 3.140000e+00, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = extractelement <11 x double> [[X]], i64 0
+; OPT-NEXT:    [[TMP1:%.*]] = extractelement <11 x double> [[X]], i64 1
+; OPT-NEXT:    [[TMP2:%.*]] = extractelement <11 x double> [[X]], i64 2
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <11 x double> [[X]], i64 3
+; OPT-NEXT:    [[TMP4:%.*]] = extractelement <11 x double> [[X]], i64 4
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <11 x double> [[X]], i64 5
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <11 x double> [[X]], i64 6
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <11 x double> [[X]], i64 7
+; OPT-NEXT:    [[TMP8:%.*]] = extractelement <11 x double> [[X]], i64 8
+; OPT-NEXT:    [[TMP9:%.*]] = extractelement <11 x double> [[X]], i64 9
+; OPT-NEXT:    [[TMP10:%.*]] = extractelement <11 x double> [[X]], i64 10
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP11:%.*]] = phi double [ [[TMP0]], [[THEN]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+; OPT-NEXT:    [[TMP12:%.*]] = phi double [ [[TMP1]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP13:%.*]] = phi double [ [[TMP2]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP14:%.*]] = phi double [ [[TMP3]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP15:%.*]] = phi double [ [[TMP4]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP16:%.*]] = phi double [ [[TMP5]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP17:%.*]] = phi double [ [[TMP6]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP18:%.*]] = phi double [ [[TMP7]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP19:%.*]] = phi double [ [[TMP8]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP20:%.*]] = phi double [ [[TMP9]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP21:%.*]] = phi double [ [[TMP10]], [[THEN]] ], [ 0.000000e+00, [[ENTRY]] ]
+; OPT-NEXT:    [[TMP22:%.*]] = insertelement <11 x double> poison, double [[TMP11]], i64 0
+; OPT-NEXT:    [[TMP23:%.*]] = insertelement <11 x double> [[TMP22]], double [[TMP12]], i64 1
+; OPT-NEXT:    [[TMP24:%.*]] = insertelement <11 x double> [[TMP23]], double [[TMP13]], i64 2
+; OPT-NEXT:    [[TMP25:%.*]] = insertelement <11 x double> [[TMP24]], double [[TMP14]], i64 3
+; OPT-NEXT:    [[TMP26:%.*]] = insertelement <11 x double> [[TMP25]], double [[TMP15]], i64 4
+; OPT-NEXT:    [[TMP27:%.*]] = insertelement <11 x double> [[TMP26]], double [[TMP16]], i64 5
+; OPT-NEXT:    [[TMP28:%.*]] = insertelement <11 x double> [[TMP27]], double [[TMP17]], i64 6
+; OPT-NEXT:    [[TMP29:%.*]] = insertelement <11 x double> [[TMP28]], double [[TMP18]], i64 7
+; OPT-NEXT:    [[TMP30:%.*]] = insertelement <11 x double> [[TMP29]], double [[TMP19]], i64 8
+; OPT-NEXT:    [[TMP31:%.*]] = insertelement <11 x double> [[TMP30]], double [[TMP20]], i64 9
+; OPT-NEXT:    [[TMP32:%.*]] = insertelement <11 x double> [[TMP31]], double [[TMP21]], i64 10
+; OPT-NEXT:    store <11 x double> [[TMP32]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v11f64_cst(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[FINALLY:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <11 x double> [[IN:%.*]], double 3.140000e+00, i32 3
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <11 x double> [ [[X]], [[THEN]] ], [ zeroinitializer, [[ENTRY:%.*]] ]
+; NOOPT-NEXT:    store <11 x double> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %finally
+then:
+  %x = insertelement <11 x double> %in, double 3.14, i32 3
+  br label %finally
+finally:
+  %val = phi <11 x double> [%x, %then], [zeroinitializer, %entry]
+  store <11 x double> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v15i64(<15 x i64> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v15i64(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <15 x i64> [[IN:%.*]], i64 42, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = extractelement <15 x i64> [[X]], i64 0
+; OPT-NEXT:    [[TMP1:%.*]] = extractelement <15 x i64> [[X]], i64 1
+; OPT-NEXT:    [[TMP2:%.*]] = extractelement <15 x i64> [[X]], i64 2
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <15 x i64> [[X]], i64 3
+; OPT-NEXT:    [[TMP4:%.*]] = extractelement <15 x i64> [[X]], i64 4
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <15 x i64> [[X]], i64 5
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <15 x i64> [[X]], i64 6
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <15 x i64> [[X]], i64 7
+; OPT-NEXT:    [[TMP8:%.*]] = extractelement <15 x i64> [[X]], i64 8
+; OPT-NEXT:    [[TMP9:%.*]] = extractelement <15 x i64> [[X]], i64 9
+; OPT-NEXT:    [[TMP10:%.*]] = extractelement <15 x i64> [[X]], i64 10
+; OPT-NEXT:    [[TMP11:%.*]] = extractelement <15 x i64> [[X]], i64 11
+; OPT-NEXT:    [[TMP12:%.*]] = extractelement <15 x i64> [[X]], i64 12
+; OPT-NEXT:    [[TMP13:%.*]] = extractelement <15 x i64> [[X]], i64 13
+; OPT-NEXT:    [[TMP14:%.*]] = extractelement <15 x i64> [[X]], i64 14
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <15 x i64> [[IN]], i64 64, i32 6
+; OPT-NEXT:    [[TMP15:%.*]] = extractelement <15 x i64> [[Y]], i64 0
+; OPT-NEXT:    [[TMP16:%.*]] = extractelement <15 x i64> [[Y]], i64 1
+; OPT-NEXT:    [[TMP17:%.*]] = extractelement <15 x i64> [[Y]], i64 2
+; OPT-NEXT:    [[TMP18:%.*]] = extractelement <15 x i64> [[Y]], i64 3
+; OPT-NEXT:    [[TMP19:%.*]] = extractelement <15 x i64> [[Y]], i64 4
+; OPT-NEXT:    [[TMP20:%.*]] = extractelement <15 x i64> [[Y]], i64 5
+; OPT-NEXT:    [[TMP21:%.*]] = extractelement <15 x i64> [[Y]], i64 6
+; OPT-NEXT:    [[TMP22:%.*]] = extractelement <15 x i64> [[Y]], i64 7
+; OPT-NEXT:    [[TMP23:%.*]] = extractelement <15 x i64> [[Y]], i64 8
+; OPT-NEXT:    [[TMP24:%.*]] = extractelement <15 x i64> [[Y]], i64 9
+; OPT-NEXT:    [[TMP25:%.*]] = extractelement <15 x i64> [[Y]], i64 10
+; OPT-NEXT:    [[TMP26:%.*]] = extractelement <15 x i64> [[Y]], i64 11
+; OPT-NEXT:    [[TMP27:%.*]] = extractelement <15 x i64> [[Y]], i64 12
+; OPT-NEXT:    [[TMP28:%.*]] = extractelement <15 x i64> [[Y]], i64 13
+; OPT-NEXT:    [[TMP29:%.*]] = extractelement <15 x i64> [[Y]], i64 14
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP30:%.*]] = phi i64 [ [[TMP0]], [[THEN]] ], [ [[TMP15]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP31:%.*]] = phi i64 [ [[TMP1]], [[THEN]] ], [ [[TMP16]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP32:%.*]] = phi i64 [ [[TMP2]], [[THEN]] ], [ [[TMP17]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP33:%.*]] = phi i64 [ [[TMP3]], [[THEN]] ], [ [[TMP18]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP34:%.*]] = phi i64 [ [[TMP4]], [[THEN]] ], [ [[TMP19]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP35:%.*]] = phi i64 [ [[TMP5]], [[THEN]] ], [ [[TMP20]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP36:%.*]] = phi i64 [ [[TMP6]], [[THEN]] ], [ [[TMP21]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP37:%.*]] = phi i64 [ [[TMP7]], [[THEN]] ], [ [[TMP22]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP38:%.*]] = phi i64 [ [[TMP8]], [[THEN]] ], [ [[TMP23]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP39:%.*]] = phi i64 [ [[TMP9]], [[THEN]] ], [ [[TMP24]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP40:%.*]] = phi i64 [ [[TMP10]], [[THEN]] ], [ [[TMP25]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP41:%.*]] = phi i64 [ [[TMP11]], [[THEN]] ], [ [[TMP26]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP42:%.*]] = phi i64 [ [[TMP12]], [[THEN]] ], [ [[TMP27]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP43:%.*]] = phi i64 [ [[TMP13]], [[THEN]] ], [ [[TMP28]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP44:%.*]] = phi i64 [ [[TMP14]], [[THEN]] ], [ [[TMP29]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP45:%.*]] = insertelement <15 x i64> poison, i64 [[TMP30]], i64 0
+; OPT-NEXT:    [[TMP46:%.*]] = insertelement <15 x i64> [[TMP45]], i64 [[TMP31]], i64 1
+; OPT-NEXT:    [[TMP47:%.*]] = insertelement <15 x i64> [[TMP46]], i64 [[TMP32]], i64 2
+; OPT-NEXT:    [[TMP48:%.*]] = insertelement <15 x i64> [[TMP47]], i64 [[TMP33]], i64 3
+; OPT-NEXT:    [[TMP49:%.*]] = insertelement <15 x i64> [[TMP48]], i64 [[TMP34]], i64 4
+; OPT-NEXT:    [[TMP50:%.*]] = insertelement <15 x i64> [[TMP49]], i64 [[TMP35]], i64 5
+; OPT-NEXT:    [[TMP51:%.*]] = insertelement <15 x i64> [[TMP50]], i64 [[TMP36]], i64 6
+; OPT-NEXT:    [[TMP52:%.*]] = insertelement <15 x i64> [[TMP51]], i64 [[TMP37]], i64 7
+; OPT-NEXT:    [[TMP53:%.*]] = insertelement <15 x i64> [[TMP52]], i64 [[TMP38]], i64 8
+; OPT-NEXT:    [[TMP54:%.*]] = insertelement <15 x i64> [[TMP53]], i64 [[TMP39]], i64 9
+; OPT-NEXT:    [[TMP55:%.*]] = insertelement <15 x i64> [[TMP54]], i64 [[TMP40]], i64 10
+; OPT-NEXT:    [[TMP56:%.*]] = insertelement <15 x i64> [[TMP55]], i64 [[TMP41]], i64 11
+; OPT-NEXT:    [[TMP57:%.*]] = insertelement <15 x i64> [[TMP56]], i64 [[TMP42]], i64 12
+; OPT-NEXT:    [[TMP58:%.*]] = insertelement <15 x i64> [[TMP57]], i64 [[TMP43]], i64 13
+; OPT-NEXT:    [[TMP59:%.*]] = insertelement <15 x i64> [[TMP58]], i64 [[TMP44]], i64 14
+; OPT-NEXT:    store <15 x i64> [[TMP59]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v15i64(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <15 x i64> [[IN:%.*]], i64 42, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <15 x i64> [[IN]], i64 64, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <15 x i64> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <15 x i64> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <15 x i64> %in, i64 42, i32 3
+  br label %finally
+else:
+  %y = insertelement <15 x i64> %in, i64 64, i32 6
+  br label %finally
+finally:
+  %val = phi <15 x i64> [%x, %then], [%y, %else]
+  store <15 x i64> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v27i16(<27 x i16> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v27i16(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <27 x i16> [[IN:%.*]], i16 42, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 0, i32 1>
+; OPT-NEXT:    [[TMP1:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 2, i32 3>
+; OPT-NEXT:    [[TMP2:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 4, i32 5>
+; OPT-NEXT:    [[TMP3:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 6, i32 7>
+; OPT-NEXT:    [[TMP4:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 8, i32 9>
+; OPT-NEXT:    [[TMP5:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 10, i32 11>
+; OPT-NEXT:    [[TMP6:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 12, i32 13>
+; OPT-NEXT:    [[TMP7:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 14, i32 15>
+; OPT-NEXT:    [[TMP8:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 16, i32 17>
+; OPT-NEXT:    [[TMP9:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 18, i32 19>
+; OPT-NEXT:    [[TMP10:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 20, i32 21>
+; OPT-NEXT:    [[TMP11:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 22, i32 23>
+; OPT-NEXT:    [[TMP12:%.*]] = shufflevector <27 x i16> [[X]], <27 x i16> poison, <2 x i32> <i32 24, i32 25>
+; OPT-NEXT:    [[TMP13:%.*]] = extractelement <27 x i16> [[X]], i64 26
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <27 x i16> [[IN]], i16 64, i32 6
+; OPT-NEXT:    [[TMP14:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 0, i32 1>
+; OPT-NEXT:    [[TMP15:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 2, i32 3>
+; OPT-NEXT:    [[TMP16:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 4, i32 5>
+; OPT-NEXT:    [[TMP17:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 6, i32 7>
+; OPT-NEXT:    [[TMP18:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 8, i32 9>
+; OPT-NEXT:    [[TMP19:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 10, i32 11>
+; OPT-NEXT:    [[TMP20:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 12, i32 13>
+; OPT-NEXT:    [[TMP21:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 14, i32 15>
+; OPT-NEXT:    [[TMP22:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 16, i32 17>
+; OPT-NEXT:    [[TMP23:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 18, i32 19>
+; OPT-NEXT:    [[TMP24:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 20, i32 21>
+; OPT-NEXT:    [[TMP25:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 22, i32 23>
+; OPT-NEXT:    [[TMP26:%.*]] = shufflevector <27 x i16> [[Y]], <27 x i16> poison, <2 x i32> <i32 24, i32 25>
+; OPT-NEXT:    [[TMP27:%.*]] = extractelement <27 x i16> [[Y]], i64 26
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP28:%.*]] = phi <2 x i16> [ [[TMP0]], [[THEN]] ], [ [[TMP14]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP29:%.*]] = phi <2 x i16> [ [[TMP1]], [[THEN]] ], [ [[TMP15]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP30:%.*]] = phi <2 x i16> [ [[TMP2]], [[THEN]] ], [ [[TMP16]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP31:%.*]] = phi <2 x i16> [ [[TMP3]], [[THEN]] ], [ [[TMP17]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP32:%.*]] = phi <2 x i16> [ [[TMP4]], [[THEN]] ], [ [[TMP18]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP33:%.*]] = phi <2 x i16> [ [[TMP5]], [[THEN]] ], [ [[TMP19]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP34:%.*]] = phi <2 x i16> [ [[TMP6]], [[THEN]] ], [ [[TMP20]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP35:%.*]] = phi <2 x i16> [ [[TMP7]], [[THEN]] ], [ [[TMP21]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP36:%.*]] = phi <2 x i16> [ [[TMP8]], [[THEN]] ], [ [[TMP22]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP37:%.*]] = phi <2 x i16> [ [[TMP9]], [[THEN]] ], [ [[TMP23]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP38:%.*]] = phi <2 x i16> [ [[TMP10]], [[THEN]] ], [ [[TMP24]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP39:%.*]] = phi <2 x i16> [ [[TMP11]], [[THEN]] ], [ [[TMP25]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP40:%.*]] = phi <2 x i16> [ [[TMP12]], [[THEN]] ], [ [[TMP26]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP41:%.*]] = phi i16 [ [[TMP13]], [[THEN]] ], [ [[TMP27]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP42:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> poison, <2 x i16> [[TMP28]], i64 0)
+; OPT-NEXT:    [[TMP43:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP42]], <2 x i16> [[TMP29]], i64 2)
+; OPT-NEXT:    [[TMP44:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP43]], <2 x i16> [[TMP30]], i64 4)
+; OPT-NEXT:    [[TMP45:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP44]], <2 x i16> [[TMP31]], i64 6)
+; OPT-NEXT:    [[TMP46:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP45]], <2 x i16> [[TMP32]], i64 8)
+; OPT-NEXT:    [[TMP47:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP46]], <2 x i16> [[TMP33]], i64 10)
+; OPT-NEXT:    [[TMP48:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP47]], <2 x i16> [[TMP34]], i64 12)
+; OPT-NEXT:    [[TMP49:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP48]], <2 x i16> [[TMP35]], i64 14)
+; OPT-NEXT:    [[TMP50:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP49]], <2 x i16> [[TMP36]], i64 16)
+; OPT-NEXT:    [[TMP51:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP50]], <2 x i16> [[TMP37]], i64 18)
+; OPT-NEXT:    [[TMP52:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP51]], <2 x i16> [[TMP38]], i64 20)
+; OPT-NEXT:    [[TMP53:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP52]], <2 x i16> [[TMP39]], i64 22)
+; OPT-NEXT:    [[TMP54:%.*]] = call <27 x i16> @llvm.vector.insert.v27i16.v2i16(<27 x i16> [[TMP53]], <2 x i16> [[TMP40]], i64 24)
+; OPT-NEXT:    [[TMP55:%.*]] = insertelement <27 x i16> [[TMP54]], i16 [[TMP41]], i64 26
+; OPT-NEXT:    store <27 x i16> [[TMP55]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v27i16(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <27 x i16> [[IN:%.*]], i16 42, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <27 x i16> [[IN]], i16 64, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <27 x i16> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <27 x i16> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <27 x i16> %in, i16 42, i32 3
+  br label %finally
+else:
+  %y = insertelement <27 x i16> %in, i16 64, i32 6
+  br label %finally
+finally:
+  %val = phi <27 x i16> [%x, %then], [%y, %else]
+  store <27 x i16> %val, ptr %out, align 1
+  ret void
+}
+
+
+define amdgpu_kernel void @phi_v23i8(<23 x i8> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v23i8(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <23 x i8> [[IN:%.*]], i8 42, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = shufflevector <23 x i8> [[X]], <23 x i8> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3>
+; OPT-NEXT:    [[TMP1:%.*]] = shufflevector <23 x i8> [[X]], <23 x i8> poison, <4 x i32> <i32 4, i32 5, i32 6, i32 7>
+; OPT-NEXT:    [[TMP2:%.*]] = shufflevector <23 x i8> [[X]], <23 x i8> poison, <4 x i32> <i32 8, i32 9, i32 10, i32 11>
+; OPT-NEXT:    [[TMP3:%.*]] = shufflevector <23 x i8> [[X]], <23 x i8> poison, <4 x i32> <i32 12, i32 13, i32 14, i32 15>
+; OPT-NEXT:    [[TMP4:%.*]] = shufflevector <23 x i8> [[X]], <23 x i8> poison, <4 x i32> <i32 16, i32 17, i32 18, i32 19>
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <23 x i8> [[X]], i64 20
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <23 x i8> [[X]], i64 21
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <23 x i8> [[X]], i64 22
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <23 x i8> [[IN]], i8 64, i32 6
+; OPT-NEXT:    [[TMP8:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3>
+; OPT-NEXT:    [[TMP9:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 4, i32 5, i32 6, i32 7>
+; OPT-NEXT:    [[TMP10:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 8, i32 9, i32 10, i32 11>
+; OPT-NEXT:    [[TMP11:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 12, i32 13, i32 14, i32 15>
+; OPT-NEXT:    [[TMP12:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 16, i32 17, i32 18, i32 19>
+; OPT-NEXT:    [[TMP13:%.*]] = extractelement <23 x i8> [[Y]], i64 20
+; OPT-NEXT:    [[TMP14:%.*]] = extractelement <23 x i8> [[Y]], i64 21
+; OPT-NEXT:    [[TMP15:%.*]] = extractelement <23 x i8> [[Y]], i64 22
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP16:%.*]] = phi <4 x i8> [ [[TMP0]], [[THEN]] ], [ [[TMP8]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP17:%.*]] = phi <4 x i8> [ [[TMP1]], [[THEN]] ], [ [[TMP9]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP18:%.*]] = phi <4 x i8> [ [[TMP2]], [[THEN]] ], [ [[TMP10]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP19:%.*]] = phi <4 x i8> [ [[TMP3]], [[THEN]] ], [ [[TMP11]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP20:%.*]] = phi <4 x i8> [ [[TMP4]], [[THEN]] ], [ [[TMP12]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP21:%.*]] = phi i8 [ [[TMP5]], [[THEN]] ], [ [[TMP13]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP22:%.*]] = phi i8 [ [[TMP6]], [[THEN]] ], [ [[TMP14]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP23:%.*]] = phi i8 [ [[TMP7]], [[THEN]] ], [ [[TMP15]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP24:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> poison, <4 x i8> [[TMP16]], i64 0)
+; OPT-NEXT:    [[TMP25:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> [[TMP24]], <4 x i8> [[TMP17]], i64 4)
+; OPT-NEXT:    [[TMP26:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> [[TMP25]], <4 x i8> [[TMP18]], i64 8)
+; OPT-NEXT:    [[TMP27:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> [[TMP26]], <4 x i8> [[TMP19]], i64 12)
+; OPT-NEXT:    [[TMP28:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> [[TMP27]], <4 x i8> [[TMP20]], i64 16)
+; OPT-NEXT:    [[TMP29:%.*]] = insertelement <23 x i8> [[TMP28]], i8 [[TMP21]], i64 20
+; OPT-NEXT:    [[TMP30:%.*]] = insertelement <23 x i8> [[TMP29]], i8 [[TMP22]], i64 21
+; OPT-NEXT:    [[TMP31:%.*]] = insertelement <23 x i8> [[TMP30]], i8 [[TMP23]], i64 22
+; OPT-NEXT:    store <23 x i8> [[TMP31]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v23i8(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <23 x i8> [[IN:%.*]], i8 42, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <23 x i8> [[IN]], i8 64, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <23 x i8> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <23 x i8> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <23 x i8> %in, i8 42, i32 3
+  br label %finally
+else:
+  %y = insertelement <23 x i8> %in, i8 64, i32 6
+  br label %finally
+finally:
+  %val = phi <23 x i8> [%x, %then], [%y, %else]
+  store <23 x i8> %val, ptr %out, align 1
+  ret void
+}
+
+
+define amdgpu_kernel void @phi_v23i8_zeroinit(<23 x i8> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v23i8_zeroinit(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <23 x i8> [[IN:%.*]], i8 64, i32 6
+; OPT-NEXT:    [[TMP0:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3>
+; OPT-NEXT:    [[TMP1:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 4, i32 5, i32 6, i32 7>
+; OPT-NEXT:    [[TMP2:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 8, i32 9, i32 10, i32 11>
+; OPT-NEXT:    [[TMP3:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 12, i32 13, i32 14, i32 15>
+; OPT-NEXT:    [[TMP4:%.*]] = shufflevector <23 x i8> [[Y]], <23 x i8> poison, <4 x i32> <i32 16, i32 17, i32 18, i32 19>
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <23 x i8> [[Y]], i64 20
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <23 x i8> [[Y]], i64 21
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <23 x i8> [[Y]], i64 22
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP8:%.*]] = phi <4 x i8> [ zeroinitializer, [[THEN]] ], [ [[TMP0]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP9:%.*]] = phi <4 x i8> [ zeroinitializer, [[THEN]] ], [ [[TMP1]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP10:%.*]] = phi <4 x i8> [ zeroinitializer, [[THEN]] ], [ [[TMP2]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP11:%.*]] = phi <4 x i8> [ zeroinitializer, [[THEN]] ], [ [[TMP3]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP12:%.*]] = phi <4 x i8> [ zeroinitializer, [[THEN]] ], [ [[TMP4]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP13:%.*]] = phi i8 [ 0, [[THEN]] ], [ [[TMP5]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP14:%.*]] = phi i8 [ 0, [[THEN]] ], [ [[TMP6]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP15:%.*]] = phi i8 [ 0, [[THEN]] ], [ [[TMP7]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP16:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> poison, <4 x i8> [[TMP8]], i64 0)
+; OPT-NEXT:    [[TMP17:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> [[TMP16]], <4 x i8> [[TMP9]], i64 4)
+; OPT-NEXT:    [[TMP18:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> [[TMP17]], <4 x i8> [[TMP10]], i64 8)
+; OPT-NEXT:    [[TMP19:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> [[TMP18]], <4 x i8> [[TMP11]], i64 12)
+; OPT-NEXT:    [[TMP20:%.*]] = call <23 x i8> @llvm.vector.insert.v23i8.v4i8(<23 x i8> [[TMP19]], <4 x i8> [[TMP12]], i64 16)
+; OPT-NEXT:    [[TMP21:%.*]] = insertelement <23 x i8> [[TMP20]], i8 [[TMP13]], i64 20
+; OPT-NEXT:    [[TMP22:%.*]] = insertelement <23 x i8> [[TMP21]], i8 [[TMP14]], i64 21
+; OPT-NEXT:    [[TMP23:%.*]] = insertelement <23 x i8> [[TMP22]], i8 [[TMP15]], i64 22
+; OPT-NEXT:    store <23 x i8> [[TMP23]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v23i8_zeroinit(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <23 x i8> [[IN:%.*]], i8 64, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <23 x i8> [ zeroinitializer, [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <23 x i8> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  br label %finally
+else:
+  %y = insertelement <23 x i8> %in, i8 64, i32 6
+  br label %finally
+finally:
+  %val = phi <23 x i8> [zeroinitializer, %then], [%y, %else]
+  store <23 x i8> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v15i8_random_constant_init(<15 x i8> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v15i8_random_constant_init(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <15 x i8> [[IN:%.*]], i8 64, i32 6
+; OPT-NEXT:    [[TMP0:%.*]] = shufflevector <15 x i8> [[Y]], <15 x i8> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3>
+; OPT-NEXT:    [[TMP1:%.*]] = shufflevector <15 x i8> [[Y]], <15 x i8> poison, <4 x i32> <i32 4, i32 5, i32 6, i32 7>
+; OPT-NEXT:    [[TMP2:%.*]] = shufflevector <15 x i8> [[Y]], <15 x i8> poison, <4 x i32> <i32 8, i32 9, i32 10, i32 11>
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <15 x i8> [[Y]], i64 12
+; OPT-NEXT:    [[TMP4:%.*]] = extractelement <15 x i8> [[Y]], i64 13
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <15 x i8> [[Y]], i64 14
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP6:%.*]] = phi <4 x i8> [ <i8 poison, i8 1, i8 2, i8 3>, [[THEN]] ], [ [[TMP0]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP7:%.*]] = phi <4 x i8> [ <i8 4, i8 undef, i8 6, i8 7>, [[THEN]] ], [ [[TMP1]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP8:%.*]] = phi <4 x i8> [ <i8 9, i8 10, i8 11, i8 12>, [[THEN]] ], [ [[TMP2]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP9:%.*]] = phi i8 [ 13, [[THEN]] ], [ [[TMP3]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP10:%.*]] = phi i8 [ 14, [[THEN]] ], [ [[TMP4]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP11:%.*]] = phi i8 [ undef, [[THEN]] ], [ [[TMP5]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP12:%.*]] = call <15 x i8> @llvm.vector.insert.v15i8.v4i8(<15 x i8> poison, <4 x i8> [[TMP6]], i64 0)
+; OPT-NEXT:    [[TMP13:%.*]] = call <15 x i8> @llvm.vector.insert.v15i8.v4i8(<15 x i8> [[TMP12]], <4 x i8> [[TMP7]], i64 4)
+; OPT-NEXT:    [[TMP14:%.*]] = call <15 x i8> @llvm.vector.insert.v15i8.v4i8(<15 x i8> [[TMP13]], <4 x i8> [[TMP8]], i64 8)
+; OPT-NEXT:    [[TMP15:%.*]] = insertelement <15 x i8> [[TMP14]], i8 [[TMP9]], i64 12
+; OPT-NEXT:    [[TMP16:%.*]] = insertelement <15 x i8> [[TMP15]], i8 [[TMP10]], i64 13
+; OPT-NEXT:    [[TMP17:%.*]] = insertelement <15 x i8> [[TMP16]], i8 [[TMP11]], i64 14
+; OPT-NEXT:    store <15 x i8> [[TMP17]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v15i8_random_constant_init(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <15 x i8> [[IN:%.*]], i8 64, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <15 x i8> [ <i8 poison, i8 1, i8 2, i8 3, i8 4, i8 undef, i8 6, i8 7, i8 9, i8 10, i8 11, i8 12, i8 13, i8 14, i8 undef>, [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <15 x i8> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  br label %finally
+else:
+  %y = insertelement <15 x i8> %in, i8 64, i32 6
+  br label %finally
+finally:
+  %val = phi <15 x i8> [<i8 poison, i8 1, i8 2, i8 3, i8 4, i8 undef, i8 6, i8 7, i8 9, i8 10, i8 11, i8 12, i8 13, i8 14, i8 undef>, %then], [%y, %else]
+  store <15 x i8> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v23i32(<23 x i32> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v23i32(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <23 x i32> [[IN:%.*]], i32 42, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = extractelement <23 x i32> [[X]], i64 0
+; OPT-NEXT:    [[TMP1:%.*]] = extractelement <23 x i32> [[X]], i64 1
+; OPT-NEXT:    [[TMP2:%.*]] = extractelement <23 x i32> [[X]], i64 2
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <23 x i32> [[X]], i64 3
+; OPT-NEXT:    [[TMP4:%.*]] = extractelement <23 x i32> [[X]], i64 4
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <23 x i32> [[X]], i64 5
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <23 x i32> [[X]], i64 6
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <23 x i32> [[X]], i64 7
+; OPT-NEXT:    [[TMP8:%.*]] = extractelement <23 x i32> [[X]], i64 8
+; OPT-NEXT:    [[TMP9:%.*]] = extractelement <23 x i32> [[X]], i64 9
+; OPT-NEXT:    [[TMP10:%.*]] = extractelement <23 x i32> [[X]], i64 10
+; OPT-NEXT:    [[TMP11:%.*]] = extractelement <23 x i32> [[X]], i64 11
+; OPT-NEXT:    [[TMP12:%.*]] = extractelement <23 x i32> [[X]], i64 12
+; OPT-NEXT:    [[TMP13:%.*]] = extractelement <23 x i32> [[X]], i64 13
+; OPT-NEXT:    [[TMP14:%.*]] = extractelement <23 x i32> [[X]], i64 14
+; OPT-NEXT:    [[TMP15:%.*]] = extractelement <23 x i32> [[X]], i64 15
+; OPT-NEXT:    [[TMP16:%.*]] = extractelement <23 x i32> [[X]], i64 16
+; OPT-NEXT:    [[TMP17:%.*]] = extractelement <23 x i32> [[X]], i64 17
+; OPT-NEXT:    [[TMP18:%.*]] = extractelement <23 x i32> [[X]], i64 18
+; OPT-NEXT:    [[TMP19:%.*]] = extractelement <23 x i32> [[X]], i64 19
+; OPT-NEXT:    [[TMP20:%.*]] = extractelement <23 x i32> [[X]], i64 20
+; OPT-NEXT:    [[TMP21:%.*]] = extractelement <23 x i32> [[X]], i64 21
+; OPT-NEXT:    [[TMP22:%.*]] = extractelement <23 x i32> [[X]], i64 22
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <23 x i32> [[IN]], i32 64, i32 6
+; OPT-NEXT:    [[TMP23:%.*]] = extractelement <23 x i32> [[Y]], i64 0
+; OPT-NEXT:    [[TMP24:%.*]] = extractelement <23 x i32> [[Y]], i64 1
+; OPT-NEXT:    [[TMP25:%.*]] = extractelement <23 x i32> [[Y]], i64 2
+; OPT-NEXT:    [[TMP26:%.*]] = extractelement <23 x i32> [[Y]], i64 3
+; OPT-NEXT:    [[TMP27:%.*]] = extractelement <23 x i32> [[Y]], i64 4
+; OPT-NEXT:    [[TMP28:%.*]] = extractelement <23 x i32> [[Y]], i64 5
+; OPT-NEXT:    [[TMP29:%.*]] = extractelement <23 x i32> [[Y]], i64 6
+; OPT-NEXT:    [[TMP30:%.*]] = extractelement <23 x i32> [[Y]], i64 7
+; OPT-NEXT:    [[TMP31:%.*]] = extractelement <23 x i32> [[Y]], i64 8
+; OPT-NEXT:    [[TMP32:%.*]] = extractelement <23 x i32> [[Y]], i64 9
+; OPT-NEXT:    [[TMP33:%.*]] = extractelement <23 x i32> [[Y]], i64 10
+; OPT-NEXT:    [[TMP34:%.*]] = extractelement <23 x i32> [[Y]], i64 11
+; OPT-NEXT:    [[TMP35:%.*]] = extractelement <23 x i32> [[Y]], i64 12
+; OPT-NEXT:    [[TMP36:%.*]] = extractelement <23 x i32> [[Y]], i64 13
+; OPT-NEXT:    [[TMP37:%.*]] = extractelement <23 x i32> [[Y]], i64 14
+; OPT-NEXT:    [[TMP38:%.*]] = extractelement <23 x i32> [[Y]], i64 15
+; OPT-NEXT:    [[TMP39:%.*]] = extractelement <23 x i32> [[Y]], i64 16
+; OPT-NEXT:    [[TMP40:%.*]] = extractelement <23 x i32> [[Y]], i64 17
+; OPT-NEXT:    [[TMP41:%.*]] = extractelement <23 x i32> [[Y]], i64 18
+; OPT-NEXT:    [[TMP42:%.*]] = extractelement <23 x i32> [[Y]], i64 19
+; OPT-NEXT:    [[TMP43:%.*]] = extractelement <23 x i32> [[Y]], i64 20
+; OPT-NEXT:    [[TMP44:%.*]] = extractelement <23 x i32> [[Y]], i64 21
+; OPT-NEXT:    [[TMP45:%.*]] = extractelement <23 x i32> [[Y]], i64 22
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP46:%.*]] = phi i32 [ [[TMP0]], [[THEN]] ], [ [[TMP23]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP47:%.*]] = phi i32 [ [[TMP1]], [[THEN]] ], [ [[TMP24]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP48:%.*]] = phi i32 [ [[TMP2]], [[THEN]] ], [ [[TMP25]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP49:%.*]] = phi i32 [ [[TMP3]], [[THEN]] ], [ [[TMP26]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP50:%.*]] = phi i32 [ [[TMP4]], [[THEN]] ], [ [[TMP27]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP51:%.*]] = phi i32 [ [[TMP5]], [[THEN]] ], [ [[TMP28]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP52:%.*]] = phi i32 [ [[TMP6]], [[THEN]] ], [ [[TMP29]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP53:%.*]] = phi i32 [ [[TMP7]], [[THEN]] ], [ [[TMP30]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP54:%.*]] = phi i32 [ [[TMP8]], [[THEN]] ], [ [[TMP31]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP55:%.*]] = phi i32 [ [[TMP9]], [[THEN]] ], [ [[TMP32]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP56:%.*]] = phi i32 [ [[TMP10]], [[THEN]] ], [ [[TMP33]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP57:%.*]] = phi i32 [ [[TMP11]], [[THEN]] ], [ [[TMP34]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP58:%.*]] = phi i32 [ [[TMP12]], [[THEN]] ], [ [[TMP35]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP59:%.*]] = phi i32 [ [[TMP13]], [[THEN]] ], [ [[TMP36]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP60:%.*]] = phi i32 [ [[TMP14]], [[THEN]] ], [ [[TMP37]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP61:%.*]] = phi i32 [ [[TMP15]], [[THEN]] ], [ [[TMP38]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP62:%.*]] = phi i32 [ [[TMP16]], [[THEN]] ], [ [[TMP39]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP63:%.*]] = phi i32 [ [[TMP17]], [[THEN]] ], [ [[TMP40]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP64:%.*]] = phi i32 [ [[TMP18]], [[THEN]] ], [ [[TMP41]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP65:%.*]] = phi i32 [ [[TMP19]], [[THEN]] ], [ [[TMP42]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP66:%.*]] = phi i32 [ [[TMP20]], [[THEN]] ], [ [[TMP43]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP67:%.*]] = phi i32 [ [[TMP21]], [[THEN]] ], [ [[TMP44]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP68:%.*]] = phi i32 [ [[TMP22]], [[THEN]] ], [ [[TMP45]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP69:%.*]] = insertelement <23 x i32> poison, i32 [[TMP46]], i64 0
+; OPT-NEXT:    [[TMP70:%.*]] = insertelement <23 x i32> [[TMP69]], i32 [[TMP47]], i64 1
+; OPT-NEXT:    [[TMP71:%.*]] = insertelement <23 x i32> [[TMP70]], i32 [[TMP48]], i64 2
+; OPT-NEXT:    [[TMP72:%.*]] = insertelement <23 x i32> [[TMP71]], i32 [[TMP49]], i64 3
+; OPT-NEXT:    [[TMP73:%.*]] = insertelement <23 x i32> [[TMP72]], i32 [[TMP50]], i64 4
+; OPT-NEXT:    [[TMP74:%.*]] = insertelement <23 x i32> [[TMP73]], i32 [[TMP51]], i64 5
+; OPT-NEXT:    [[TMP75:%.*]] = insertelement <23 x i32> [[TMP74]], i32 [[TMP52]], i64 6
+; OPT-NEXT:    [[TMP76:%.*]] = insertelement <23 x i32> [[TMP75]], i32 [[TMP53]], i64 7
+; OPT-NEXT:    [[TMP77:%.*]] = insertelement <23 x i32> [[TMP76]], i32 [[TMP54]], i64 8
+; OPT-NEXT:    [[TMP78:%.*]] = insertelement <23 x i32> [[TMP77]], i32 [[TMP55]], i64 9
+; OPT-NEXT:    [[TMP79:%.*]] = insertelement <23 x i32> [[TMP78]], i32 [[TMP56]], i64 10
+; OPT-NEXT:    [[TMP80:%.*]] = insertelement <23 x i32> [[TMP79]], i32 [[TMP57]], i64 11
+; OPT-NEXT:    [[TMP81:%.*]] = insertelement <23 x i32> [[TMP80]], i32 [[TMP58]], i64 12
+; OPT-NEXT:    [[TMP82:%.*]] = insertelement <23 x i32> [[TMP81]], i32 [[TMP59]], i64 13
+; OPT-NEXT:    [[TMP83:%.*]] = insertelement <23 x i32> [[TMP82]], i32 [[TMP60]], i64 14
+; OPT-NEXT:    [[TMP84:%.*]] = insertelement <23 x i32> [[TMP83]], i32 [[TMP61]], i64 15
+; OPT-NEXT:    [[TMP85:%.*]] = insertelement <23 x i32> [[TMP84]], i32 [[TMP62]], i64 16
+; OPT-NEXT:    [[TMP86:%.*]] = insertelement <23 x i32> [[TMP85]], i32 [[TMP63]], i64 17
+; OPT-NEXT:    [[TMP87:%.*]] = insertelement <23 x i32> [[TMP86]], i32 [[TMP64]], i64 18
+; OPT-NEXT:    [[TMP88:%.*]] = insertelement <23 x i32> [[TMP87]], i32 [[TMP65]], i64 19
+; OPT-NEXT:    [[TMP89:%.*]] = insertelement <23 x i32> [[TMP88]], i32 [[TMP66]], i64 20
+; OPT-NEXT:    [[TMP90:%.*]] = insertelement <23 x i32> [[TMP89]], i32 [[TMP67]], i64 21
+; OPT-NEXT:    [[TMP91:%.*]] = insertelement <23 x i32> [[TMP90]], i32 [[TMP68]], i64 22
+; OPT-NEXT:    store <23 x i32> [[TMP91]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v23i32(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <23 x i32> [[IN:%.*]], i32 42, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <23 x i32> [[IN]], i32 64, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <23 x i32> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <23 x i32> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <23 x i32> %in, i32 42, i32 3
+  br label %finally
+else:
+  %y = insertelement <23 x i32> %in, i32 64, i32 6
+  br label %finally
+finally:
+  %val = phi <23 x i32> [%x, %then], [%y, %else]
+  store <23 x i32> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v16i64(<16 x i64> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v16i64(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <16 x i64> [[IN:%.*]], i64 42, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = extractelement <16 x i64> [[X]], i64 0
+; OPT-NEXT:    [[TMP1:%.*]] = extractelement <16 x i64> [[X]], i64 1
+; OPT-NEXT:    [[TMP2:%.*]] = extractelement <16 x i64> [[X]], i64 2
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <16 x i64> [[X]], i64 3
+; OPT-NEXT:    [[TMP4:%.*]] = extractelement <16 x i64> [[X]], i64 4
+; OPT-NEXT:    [[TMP5:%.*]] = extractelement <16 x i64> [[X]], i64 5
+; OPT-NEXT:    [[TMP6:%.*]] = extractelement <16 x i64> [[X]], i64 6
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <16 x i64> [[X]], i64 7
+; OPT-NEXT:    [[TMP8:%.*]] = extractelement <16 x i64> [[X]], i64 8
+; OPT-NEXT:    [[TMP9:%.*]] = extractelement <16 x i64> [[X]], i64 9
+; OPT-NEXT:    [[TMP10:%.*]] = extractelement <16 x i64> [[X]], i64 10
+; OPT-NEXT:    [[TMP11:%.*]] = extractelement <16 x i64> [[X]], i64 11
+; OPT-NEXT:    [[TMP12:%.*]] = extractelement <16 x i64> [[X]], i64 12
+; OPT-NEXT:    [[TMP13:%.*]] = extractelement <16 x i64> [[X]], i64 13
+; OPT-NEXT:    [[TMP14:%.*]] = extractelement <16 x i64> [[X]], i64 14
+; OPT-NEXT:    [[TMP15:%.*]] = extractelement <16 x i64> [[X]], i64 15
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <16 x i64> [[IN]], i64 64, i32 6
+; OPT-NEXT:    [[TMP16:%.*]] = extractelement <16 x i64> [[Y]], i64 0
+; OPT-NEXT:    [[TMP17:%.*]] = extractelement <16 x i64> [[Y]], i64 1
+; OPT-NEXT:    [[TMP18:%.*]] = extractelement <16 x i64> [[Y]], i64 2
+; OPT-NEXT:    [[TMP19:%.*]] = extractelement <16 x i64> [[Y]], i64 3
+; OPT-NEXT:    [[TMP20:%.*]] = extractelement <16 x i64> [[Y]], i64 4
+; OPT-NEXT:    [[TMP21:%.*]] = extractelement <16 x i64> [[Y]], i64 5
+; OPT-NEXT:    [[TMP22:%.*]] = extractelement <16 x i64> [[Y]], i64 6
+; OPT-NEXT:    [[TMP23:%.*]] = extractelement <16 x i64> [[Y]], i64 7
+; OPT-NEXT:    [[TMP24:%.*]] = extractelement <16 x i64> [[Y]], i64 8
+; OPT-NEXT:    [[TMP25:%.*]] = extractelement <16 x i64> [[Y]], i64 9
+; OPT-NEXT:    [[TMP26:%.*]] = extractelement <16 x i64> [[Y]], i64 10
+; OPT-NEXT:    [[TMP27:%.*]] = extractelement <16 x i64> [[Y]], i64 11
+; OPT-NEXT:    [[TMP28:%.*]] = extractelement <16 x i64> [[Y]], i64 12
+; OPT-NEXT:    [[TMP29:%.*]] = extractelement <16 x i64> [[Y]], i64 13
+; OPT-NEXT:    [[TMP30:%.*]] = extractelement <16 x i64> [[Y]], i64 14
+; OPT-NEXT:    [[TMP31:%.*]] = extractelement <16 x i64> [[Y]], i64 15
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP32:%.*]] = phi i64 [ [[TMP0]], [[THEN]] ], [ [[TMP16]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP33:%.*]] = phi i64 [ [[TMP1]], [[THEN]] ], [ [[TMP17]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP34:%.*]] = phi i64 [ [[TMP2]], [[THEN]] ], [ [[TMP18]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP35:%.*]] = phi i64 [ [[TMP3]], [[THEN]] ], [ [[TMP19]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP36:%.*]] = phi i64 [ [[TMP4]], [[THEN]] ], [ [[TMP20]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP37:%.*]] = phi i64 [ [[TMP5]], [[THEN]] ], [ [[TMP21]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP38:%.*]] = phi i64 [ [[TMP6]], [[THEN]] ], [ [[TMP22]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP39:%.*]] = phi i64 [ [[TMP7]], [[THEN]] ], [ [[TMP23]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP40:%.*]] = phi i64 [ [[TMP8]], [[THEN]] ], [ [[TMP24]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP41:%.*]] = phi i64 [ [[TMP9]], [[THEN]] ], [ [[TMP25]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP42:%.*]] = phi i64 [ [[TMP10]], [[THEN]] ], [ [[TMP26]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP43:%.*]] = phi i64 [ [[TMP11]], [[THEN]] ], [ [[TMP27]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP44:%.*]] = phi i64 [ [[TMP12]], [[THEN]] ], [ [[TMP28]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP45:%.*]] = phi i64 [ [[TMP13]], [[THEN]] ], [ [[TMP29]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP46:%.*]] = phi i64 [ [[TMP14]], [[THEN]] ], [ [[TMP30]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP47:%.*]] = phi i64 [ [[TMP15]], [[THEN]] ], [ [[TMP31]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP48:%.*]] = insertelement <16 x i64> poison, i64 [[TMP32]], i64 0
+; OPT-NEXT:    [[TMP49:%.*]] = insertelement <16 x i64> [[TMP48]], i64 [[TMP33]], i64 1
+; OPT-NEXT:    [[TMP50:%.*]] = insertelement <16 x i64> [[TMP49]], i64 [[TMP34]], i64 2
+; OPT-NEXT:    [[TMP51:%.*]] = insertelement <16 x i64> [[TMP50]], i64 [[TMP35]], i64 3
+; OPT-NEXT:    [[TMP52:%.*]] = insertelement <16 x i64> [[TMP51]], i64 [[TMP36]], i64 4
+; OPT-NEXT:    [[TMP53:%.*]] = insertelement <16 x i64> [[TMP52]], i64 [[TMP37]], i64 5
+; OPT-NEXT:    [[TMP54:%.*]] = insertelement <16 x i64> [[TMP53]], i64 [[TMP38]], i64 6
+; OPT-NEXT:    [[TMP55:%.*]] = insertelement <16 x i64> [[TMP54]], i64 [[TMP39]], i64 7
+; OPT-NEXT:    [[TMP56:%.*]] = insertelement <16 x i64> [[TMP55]], i64 [[TMP40]], i64 8
+; OPT-NEXT:    [[TMP57:%.*]] = insertelement <16 x i64> [[TMP56]], i64 [[TMP41]], i64 9
+; OPT-NEXT:    [[TMP58:%.*]] = insertelement <16 x i64> [[TMP57]], i64 [[TMP42]], i64 10
+; OPT-NEXT:    [[TMP59:%.*]] = insertelement <16 x i64> [[TMP58]], i64 [[TMP43]], i64 11
+; OPT-NEXT:    [[TMP60:%.*]] = insertelement <16 x i64> [[TMP59]], i64 [[TMP44]], i64 12
+; OPT-NEXT:    [[TMP61:%.*]] = insertelement <16 x i64> [[TMP60]], i64 [[TMP45]], i64 13
+; OPT-NEXT:    [[TMP62:%.*]] = insertelement <16 x i64> [[TMP61]], i64 [[TMP46]], i64 14
+; OPT-NEXT:    [[TMP63:%.*]] = insertelement <16 x i64> [[TMP62]], i64 [[TMP47]], i64 15
+; OPT-NEXT:    store <16 x i64> [[TMP63]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v16i64(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <16 x i64> [[IN:%.*]], i64 42, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <16 x i64> [[IN]], i64 64, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <16 x i64> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <16 x i64> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <16 x i64> %in, i64 42, i32 3
+  br label %finally
+else:
+  %y = insertelement <16 x i64> %in, i64 64, i32 6
+  br label %finally
+finally:
+  %val = phi <16 x i64> [%x, %then], [%y, %else]
+  store <16 x i64> %val, ptr %out, align 1
+  ret void
+}
+
+define amdgpu_kernel void @phi_v7i16(<7 x i16> %in, ptr %out, i1 %cond) {
+; OPT-LABEL: @phi_v7i16(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; OPT:       then:
+; OPT-NEXT:    [[X:%.*]] = insertelement <7 x i16> [[IN:%.*]], i16 3, i32 3
+; OPT-NEXT:    [[TMP0:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 0, i32 1>
+; OPT-NEXT:    [[TMP1:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 2, i32 3>
+; OPT-NEXT:    [[TMP2:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 4, i32 5>
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <7 x i16> [[X]], i64 6
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <7 x i16> [[IN]], i16 9, i32 6
+; OPT-NEXT:    [[TMP4:%.*]] = shufflevector <7 x i16> [[Y]], <7 x i16> poison, <2 x i32> <i32 0, i32 1>
+; OPT-NEXT:    [[TMP5:%.*]] = shufflevector <7 x i16> [[Y]], <7 x i16> poison, <2 x i32> <i32 2, i32 3>
+; OPT-NEXT:    [[TMP6:%.*]] = shufflevector <7 x i16> [[Y]], <7 x i16> poison, <2 x i32> <i32 4, i32 5>
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <7 x i16> [[Y]], i64 6
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP8:%.*]] = phi <2 x i16> [ [[TMP0]], [[THEN]] ], [ [[TMP4]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP9:%.*]] = phi <2 x i16> [ [[TMP1]], [[THEN]] ], [ [[TMP5]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP10:%.*]] = phi <2 x i16> [ [[TMP2]], [[THEN]] ], [ [[TMP6]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP11:%.*]] = phi i16 [ [[TMP3]], [[THEN]] ], [ [[TMP7]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP12:%.*]] = call <7 x i16> @llvm.vector.insert.v7i16.v2i16(<7 x i16> poison, <2 x i16> [[TMP8]], i64 0)
+; OPT-NEXT:    [[TMP13:%.*]] = call <7 x i16> @llvm.vector.insert.v7i16.v2i16(<7 x i16> [[TMP12]], <2 x i16> [[TMP9]], i64 2)
+; OPT-NEXT:    [[TMP14:%.*]] = call <7 x i16> @llvm.vector.insert.v7i16.v2i16(<7 x i16> [[TMP13]], <2 x i16> [[TMP10]], i64 4)
+; OPT-NEXT:    [[TMP15:%.*]] = insertelement <7 x i16> [[TMP14]], i16 [[TMP11]], i64 6
+; OPT-NEXT:    store <7 x i16> [[TMP15]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v7i16(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    br i1 [[COND:%.*]], label [[THEN:%.*]], label [[ELSE:%.*]]
+; NOOPT:       then:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <7 x i16> [[IN:%.*]], i16 3, i32 3
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <7 x i16> [[IN]], i16 9, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <7 x i16> [ [[X]], [[THEN]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <7 x i16> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  br i1 %cond, label %then, label %else
+then:
+  %x = insertelement <7 x i16> %in, i16 3, i32 3
+  br label %finally
+else:
+  %y = insertelement <7 x i16> %in, i16 9, i32 6
+  br label %finally
+finally:
+  %val = phi <7 x i16> [%x, %then], [%y, %else]
+  store <7 x i16> %val, ptr %out, align 1
+  ret void
+}
+
+
+define amdgpu_kernel void @phi_v7i16_switch(<7 x i16> %in, ptr %out, i8 %cond) {
+; OPT-LABEL: @phi_v7i16_switch(
+; OPT-NEXT:  entry:
+; OPT-NEXT:    [[X:%.*]] = insertelement <7 x i16> [[IN:%.*]], i16 3, i32 3
+; OPT-NEXT:    switch i8 [[COND:%.*]], label [[ELSE:%.*]] [
+; OPT-NEXT:    i8 0, label [[THEN_1:%.*]]
+; OPT-NEXT:    i8 3, label [[THEN_2:%.*]]
+; OPT-NEXT:    ]
+; OPT:       then.1:
+; OPT-NEXT:    [[TMP0:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 0, i32 1>
+; OPT-NEXT:    [[TMP1:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 2, i32 3>
+; OPT-NEXT:    [[TMP2:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 4, i32 5>
+; OPT-NEXT:    [[TMP3:%.*]] = extractelement <7 x i16> [[X]], i64 6
+; OPT-NEXT:    br label [[FINALLY:%.*]]
+; OPT:       then.2:
+; OPT-NEXT:    [[TMP4:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 0, i32 1>
+; OPT-NEXT:    [[TMP5:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 2, i32 3>
+; OPT-NEXT:    [[TMP6:%.*]] = shufflevector <7 x i16> [[X]], <7 x i16> poison, <2 x i32> <i32 4, i32 5>
+; OPT-NEXT:    [[TMP7:%.*]] = extractelement <7 x i16> [[X]], i64 6
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       else:
+; OPT-NEXT:    [[Y:%.*]] = insertelement <7 x i16> [[IN]], i16 9, i32 6
+; OPT-NEXT:    [[TMP8:%.*]] = shufflevector <7 x i16> [[Y]], <7 x i16> poison, <2 x i32> <i32 0, i32 1>
+; OPT-NEXT:    [[TMP9:%.*]] = shufflevector <7 x i16> [[Y]], <7 x i16> poison, <2 x i32> <i32 2, i32 3>
+; OPT-NEXT:    [[TMP10:%.*]] = shufflevector <7 x i16> [[Y]], <7 x i16> poison, <2 x i32> <i32 4, i32 5>
+; OPT-NEXT:    [[TMP11:%.*]] = extractelement <7 x i16> [[Y]], i64 6
+; OPT-NEXT:    br label [[FINALLY]]
+; OPT:       finally:
+; OPT-NEXT:    [[TMP12:%.*]] = phi <2 x i16> [ [[TMP0]], [[THEN_1]] ], [ [[TMP4]], [[THEN_2]] ], [ [[TMP8]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP13:%.*]] = phi <2 x i16> [ [[TMP1]], [[THEN_1]] ], [ [[TMP5]], [[THEN_2]] ], [ [[TMP9]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP14:%.*]] = phi <2 x i16> [ [[TMP2]], [[THEN_1]] ], [ [[TMP6]], [[THEN_2]] ], [ [[TMP10]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP15:%.*]] = phi i16 [ [[TMP3]], [[THEN_1]] ], [ [[TMP7]], [[THEN_2]] ], [ [[TMP11]], [[ELSE]] ]
+; OPT-NEXT:    [[TMP16:%.*]] = call <7 x i16> @llvm.vector.insert.v7i16.v2i16(<7 x i16> poison, <2 x i16> [[TMP12]], i64 0)
+; OPT-NEXT:    [[TMP17:%.*]] = call <7 x i16> @llvm.vector.insert.v7i16.v2i16(<7 x i16> [[TMP16]], <2 x i16> [[TMP13]], i64 2)
+; OPT-NEXT:    [[TMP18:%.*]] = call <7 x i16> @llvm.vector.insert.v7i16.v2i16(<7 x i16> [[TMP17]], <2 x i16> [[TMP14]], i64 4)
+; OPT-NEXT:    [[TMP19:%.*]] = insertelement <7 x i16> [[TMP18]], i16 [[TMP15]], i64 6
+; OPT-NEXT:    store <7 x i16> [[TMP19]], ptr [[OUT:%.*]], align 1
+; OPT-NEXT:    ret void
+;
+; NOOPT-LABEL: @phi_v7i16_switch(
+; NOOPT-NEXT:  entry:
+; NOOPT-NEXT:    [[X:%.*]] = insertelement <7 x i16> [[IN:%.*]], i16 3, i32 3
+; NOOPT-NEXT:    switch i8 [[COND:%.*]], label [[ELSE:%.*]] [
+; NOOPT-NEXT:    i8 0, label [[THEN_1:%.*]]
+; NOOPT-NEXT:    i8 3, label [[THEN_2:%.*]]
+; NOOPT-NEXT:    ]
+; NOOPT:       then.1:
+; NOOPT-NEXT:    br label [[FINALLY:%.*]]
+; NOOPT:       then.2:
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       else:
+; NOOPT-NEXT:    [[Y:%.*]] = insertelement <7 x i16> [[IN]], i16 9, i32 6
+; NOOPT-NEXT:    br label [[FINALLY]]
+; NOOPT:       finally:
+; NOOPT-NEXT:    [[VAL:%.*]] = phi <7 x i16> [ [[X]], [[THEN_1]] ], [ [[X]], [[THEN_2]] ], [ [[Y]], [[ELSE]] ]
+; NOOPT-NEXT:    store <7 x i16> [[VAL]], ptr [[OUT:%.*]], align 1
+; NOOPT-NEXT:    ret void
+;
+entry:
+  %x = insertelement <7 x i16> %in, i16 3, i32 3
+  switch i8 %cond, label %else [
+  i8 0, label %then.1
+  i8 3, label %then.2
+  ]
+then.1:
+  br label %finally
+then.2:
+  br label %finally
+else:
+  %y = insertelement <7 x i16> %in, i16 9, i32 6
+  br label %finally
+finally:
+  %val = phi <7 x i16> [%x, %then.1], [%x, %then.2], [%y, %else]
+  store <7 x i16> %val, ptr %out, align 1
+  ret void
+}

llvm/test/CodeGen/AMDGPU/early-if-convert.ll

-; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=verde -amdgpu-early-ifcvt=1 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=verde -amdgpu-early-ifcvt=1 -amdgpu-codegenprepare-break-large-phis=0 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 ; XUN: llc -march=amdgcn -mcpu=tonga -amdgpu-early-ifcvt=1 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
+; Note: breaking up large PHIs is disabled to prevent some testcases from becoming
+;  branchless.
+
 ; FIXME: This leaves behind a now unnecessary and with exec
 
 ; GCN-LABEL: {{^}}test_vccnz_ifcvt_triangle:
 ; GCN: buffer_load_dword [[VAL:v[0-9]+]]
 ; GCN: v_cmp_neq_f32_e32 vcc, 1.0, [[VAL]]
 ; GCN: v_add_f32_e32 [[ADD:v[0-9]+]], [[VAL]], [[VAL]]
 ; GCN: v_cndmask_b32_e32 [[RESULT:v[0-9]+]], [[ADD]], [[VAL]], vcc
 ; GCN: buffer_store_dword [[RESULT]]

llvm/test/CodeGen/AMDGPU/extract-subvector-16bit.ll

 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:24 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:26 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:28 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v2, v[2:3], s[4:7], 0 addr64 offset:30 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    v_lshlrev_b32_e32 v2, 16, v7
-; SI-NEXT:    v_lshlrev_b32_e32 v3, 16, v4
-; SI-NEXT:    v_or_b32_e32 v2, v6, v2
-; SI-NEXT:    v_or_b32_e32 v3, v5, v3
+; SI-NEXT:    v_lshlrev_b32_e32 v3, 16, v7
+; SI-NEXT:    v_lshlrev_b32_e32 v2, 16, v4
+; SI-NEXT:    v_or_b32_e32 v2, v5, v2
+; SI-NEXT:    v_or_b32_e32 v3, v6, v3
 ; SI-NEXT:    s_mov_b64 vcc, exec
 ; SI-NEXT:    s_cbranch_execz .LBB3_3
 ; SI-NEXT:    s_branch .LBB3_4
 ; SI-NEXT:  .LBB3_2:
 ; SI-NEXT:    ; implicit-def: $vgpr3
-; SI-NEXT:    ; implicit-def: $vgpr4
 ; SI-NEXT:    ; implicit-def: $vgpr2
+; SI-NEXT:    ; implicit-def: $vgpr4
 ; SI-NEXT:    s_mov_b64 vcc, 0
 ; SI-NEXT:  .LBB3_3: ; %T
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000
 ; SI-NEXT:    s_mov_b32 s4, s6
 ; SI-NEXT:    s_mov_b32 s5, s6
-; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 glc
+; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:2 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 offset:4 glc
+; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:4 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v5, v[0:1], s[4:7], 0 addr64 offset:6 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:8 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:10 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:12 glc

 ; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:26 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:28 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v0, v[0:1], s[4:7], 0 addr64 offset:30 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    v_lshlrev_b32_e32 v0, 16, v5
 ; SI-NEXT:    v_lshlrev_b32_e32 v1, 16, v4
-; SI-NEXT:    v_or_b32_e32 v2, v2, v0
-; SI-NEXT:    v_or_b32_e32 v3, v3, v1
+; SI-NEXT:    v_or_b32_e32 v2, v2, v1
+; SI-NEXT:    v_or_b32_e32 v3, v3, v0
 ; SI-NEXT:  .LBB3_4: ; %exit
 ; SI-NEXT:    v_bfe_i32 v0, v3, 0, 16
-; SI-NEXT:    v_bfe_i32 v1, v4, 0, 16
-; SI-NEXT:    v_bfe_i32 v2, v2, 0, 16
+; SI-NEXT:    v_bfe_i32 v1, v2, 0, 16
+; SI-NEXT:    v_bfe_i32 v2, v4, 0, 16
 ; SI-NEXT:    v_mov_b32_e32 v3, 0xffff
 ; SI-NEXT:    v_mov_b32_e32 v4, 0x8000
 ; SI-NEXT:    v_mov_b32_e32 v5, 0xffff0000
 ; SI-NEXT:    v_bfrev_b32_e32 v6, 1
 ; SI-NEXT:    v_mov_b32_e32 v7, 0xffff8000
-; SI-NEXT:    v_cmp_lt_i32_e32 vcc, -1, v0
-; SI-NEXT:    v_cndmask_b32_e32 v0, v3, v4, vcc
 ; SI-NEXT:    v_cmp_lt_i32_e32 vcc, -1, v1
-; SI-NEXT:    v_cndmask_b32_e32 v1, v5, v6, vcc
+; SI-NEXT:    v_cndmask_b32_e32 v1, v3, v4, vcc
 ; SI-NEXT:    v_cmp_lt_i32_e32 vcc, -1, v2
+; SI-NEXT:    v_cndmask_b32_e32 v4, v5, v6, vcc
+; SI-NEXT:    v_cmp_lt_i32_e32 vcc, -1, v0
 ; SI-NEXT:    v_cndmask_b32_e32 v2, -1, v7, vcc
-; SI-NEXT:    v_or_b32_e32 v0, v0, v1
-; SI-NEXT:    v_lshlrev_b32_e32 v4, 16, v2
+; SI-NEXT:    v_or_b32_e32 v0, v1, v4
+; SI-NEXT:    v_lshlrev_b32_e32 v1, 16, v2
 ; SI-NEXT:    v_and_b32_e32 v3, 0xffff, v2
-; SI-NEXT:    v_or_b32_e32 v2, v3, v4
-; SI-NEXT:    v_alignbit_b32 v1, v2, v1, 16
+; SI-NEXT:    v_or_b32_e32 v2, v3, v1
+; SI-NEXT:    v_alignbit_b32 v1, v2, v4, 16
 ; SI-NEXT:    s_setpc_b64 s[30:31]
 ;
 ; GFX9-LABEL: vec_16xi16_extract_4xi16:
 ; GFX9:       ; %bb.0:
 ; GFX9-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX9-NEXT:    s_cbranch_scc0 .LBB3_2
 ; GFX9-NEXT:  ; %bb.1: ; %F
-; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[2:3], off offset:16 glc
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[2:3], off glc
 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
+; GFX9-NEXT:    global_load_dwordx4 v[6:9], v[2:3], off offset:16 glc
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    ; kill: killed $vgpr2 killed $vgpr3
 ; GFX9-NEXT:    s_cbranch_execz .LBB3_3
 ; GFX9-NEXT:    s_branch .LBB3_4
 ; GFX9-NEXT:  .LBB3_2:
-; GFX9-NEXT:    ; implicit-def: $vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11
+; GFX9-NEXT:    ; implicit-def: $vgpr5
 ; GFX9-NEXT:  .LBB3_3: ; %T
-; GFX9-NEXT:    global_load_dwordx4 v[2:5], v[0:1], off offset:16 glc
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[0:1], off glc
 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
+; GFX9-NEXT:    global_load_dwordx4 v[6:9], v[0:1], off offset:16 glc
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    ; kill: killed $vgpr0 killed $vgpr1
 ; GFX9-NEXT:  .LBB3_4: ; %exit
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
-; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v5 op_sel_hi:[0,0]
+; GFX9-NEXT:    s_mov_b32 s4, 0xffff
+; GFX9-NEXT:    v_bfi_b32 v0, s4, v4, v4
+; GFX9-NEXT:    v_bfi_b32 v0, s4, v4, v0
+; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v0 op_sel_hi:[0,1]
 ; GFX9-NEXT:    s_movk_i32 s4, 0x8000
+; GFX9-NEXT:    v_pk_ashrrev_i16 v2, 15, v5 op_sel_hi:[0,0]
 ; GFX9-NEXT:    v_or_b32_e32 v1, 0xffff8000, v0
-; GFX9-NEXT:    v_or_b32_sdwa v2, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
-; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v4 op_sel_hi:[0,1]
-; GFX9-NEXT:    v_or_b32_e32 v3, 0xffff8000, v0
 ; GFX9-NEXT:    v_or_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
+; GFX9-NEXT:    v_or_b32_e32 v3, 0xffff8000, v2
+; GFX9-NEXT:    v_or_b32_sdwa v2, v2, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    s_mov_b32 s4, 0x5040100
-; GFX9-NEXT:    v_perm_b32 v0, v0, v3, s4
-; GFX9-NEXT:    v_perm_b32 v1, v2, v1, s4
+; GFX9-NEXT:    v_perm_b32 v0, v0, v1, s4
+; GFX9-NEXT:    v_perm_b32 v1, v2, v3, s4
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
   br i1 undef, label %T, label %F
 
 T:
   %t = load volatile <16 x i16>, ptr addrspace(1) %p0
   br label %exit
 
 F:

 ; SI-NEXT:    buffer_load_ushort v4, v[2:3], s[4:7], 0 addr64 offset:2 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v4, v[2:3], s[4:7], 0 addr64 offset:4 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v4, v[2:3], s[4:7], 0 addr64 offset:6 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v6, v[2:3], s[4:7], 0 addr64 offset:8 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v4, v[2:3], s[4:7], 0 addr64 offset:10 glc
+; SI-NEXT:    buffer_load_ushort v5, v[2:3], s[4:7], 0 addr64 offset:10 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:12 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v5, v[2:3], s[4:7], 0 addr64 offset:14 glc
+; SI-NEXT:    buffer_load_ushort v4, v[2:3], s[4:7], 0 addr64 offset:14 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:16 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:18 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:20 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:22 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:24 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:26 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:28 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v2, v[2:3], s[4:7], 0 addr64 offset:30 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    v_lshlrev_b32_e32 v2, 16, v5
-; SI-NEXT:    v_lshlrev_b32_e32 v3, 16, v4
-; SI-NEXT:    v_or_b32_e32 v2, v7, v2
+; SI-NEXT:    v_lshlrev_b32_e32 v2, 16, v4
+; SI-NEXT:    v_lshlrev_b32_e32 v3, 16, v5
 ; SI-NEXT:    v_or_b32_e32 v3, v6, v3
+; SI-NEXT:    v_or_b32_e32 v2, v7, v2
 ; SI-NEXT:    s_mov_b64 vcc, exec
 ; SI-NEXT:    s_cbranch_execz .LBB4_3
 ; SI-NEXT:    s_branch .LBB4_4
 ; SI-NEXT:  .LBB4_2:
-; SI-NEXT:    ; implicit-def: $vgpr3
-; SI-NEXT:    ; implicit-def: $vgpr4
 ; SI-NEXT:    ; implicit-def: $vgpr2
+; SI-NEXT:    ; implicit-def: $vgpr4
+; SI-NEXT:    ; implicit-def: $vgpr3
 ; SI-NEXT:    ; implicit-def: $vgpr5
 ; SI-NEXT:    s_mov_b64 vcc, 0
 ; SI-NEXT:  .LBB4_3: ; %T
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000
 ; SI-NEXT:    s_mov_b32 s4, s6
 ; SI-NEXT:    s_mov_b32 s5, s6
 ; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 offset:2 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 offset:4 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 offset:6 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:8 glc
+; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 offset:8 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:10 glc
+; SI-NEXT:    buffer_load_ushort v5, v[0:1], s[4:7], 0 addr64 offset:10 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 offset:12 glc
+; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:12 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v5, v[0:1], s[4:7], 0 addr64 offset:14 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:14 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:16 glc
+; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:16 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:18 glc
+; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:18 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:20 glc
+; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:20 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:22 glc
+; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:22 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:24 glc
+; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:24 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:26 glc
+; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:26 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:28 glc
+; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:28 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v0, v[0:1], s[4:7], 0 addr64 offset:30 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    v_lshlrev_b32_e32 v0, 16, v5
-; SI-NEXT:    v_lshlrev_b32_e32 v1, 16, v4
-; SI-NEXT:    v_or_b32_e32 v2, v2, v0
-; SI-NEXT:    v_or_b32_e32 v3, v3, v1
+; SI-NEXT:    v_lshlrev_b32_e32 v0, 16, v4
+; SI-NEXT:    v_lshlrev_b32_e32 v1, 16, v5
+; SI-NEXT:    v_or_b32_e32 v3, v2, v1
+; SI-NEXT:    v_or_b32_e32 v2, v6, v0
 ; SI-NEXT:  .LBB4_4: ; %exit
 ; SI-NEXT:    v_bfe_i32 v0, v3, 0, 16
-; SI-NEXT:    v_bfe_i32 v1, v4, 0, 16
+; SI-NEXT:    v_bfe_i32 v1, v5, 0, 16
 ; SI-NEXT:    v_bfe_i32 v2, v2, 0, 16
-; SI-NEXT:    v_bfe_i32 v3, v5, 0, 16
+; SI-NEXT:    v_bfe_i32 v3, v4, 0, 16
 ; SI-NEXT:    v_mov_b32_e32 v4, 0xffff
 ; SI-NEXT:    v_mov_b32_e32 v5, 0x8000
 ; SI-NEXT:    v_mov_b32_e32 v6, 0xffff0000
 ; SI-NEXT:    v_bfrev_b32_e32 v7, 1
 ; SI-NEXT:    v_cmp_lt_i32_e32 vcc, -1, v0
 ; SI-NEXT:    v_cndmask_b32_e32 v0, v4, v5, vcc
 ; SI-NEXT:    v_cmp_lt_i32_e32 vcc, -1, v1
 ; SI-NEXT:    v_cndmask_b32_e32 v1, v6, v7, vcc
 ; SI-NEXT:    v_cmp_lt_i32_e32 vcc, -1, v2
 ; SI-NEXT:    v_cndmask_b32_e32 v2, v4, v5, vcc
 ; SI-NEXT:    v_cmp_lt_i32_e32 vcc, -1, v3
 ; SI-NEXT:    v_cndmask_b32_e32 v3, v6, v7, vcc
 ; SI-NEXT:    v_or_b32_e32 v0, v0, v1
 ; SI-NEXT:    v_or_b32_e32 v2, v2, v3
 ; SI-NEXT:    v_alignbit_b32 v1, v2, v1, 16
 ; SI-NEXT:    v_lshrrev_b32_e32 v3, 16, v3
 ; SI-NEXT:    s_setpc_b64 s[30:31]
 ;
 ; GFX9-LABEL: vec_16xi16_extract_4xi16_2:
 ; GFX9:       ; %bb.0:
 ; GFX9-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX9-NEXT:    s_cbranch_scc0 .LBB4_2
 ; GFX9-NEXT:  ; %bb.1: ; %F
-; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[2:3], off offset:16 glc
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[2:3], off glc
 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
-; GFX9-NEXT:    ; kill: killed $vgpr2 killed $vgpr3
+; GFX9-NEXT:    global_load_dwordx4 v[8:11], v[2:3], off offset:16 glc
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    s_cbranch_execz .LBB4_3
 ; GFX9-NEXT:    s_branch .LBB4_4
 ; GFX9-NEXT:  .LBB4_2:
-; GFX9-NEXT:    ; implicit-def: $vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11
+; GFX9-NEXT:    ; implicit-def: $vgpr7
 ; GFX9-NEXT:  .LBB4_3: ; %T
-; GFX9-NEXT:    global_load_dwordx4 v[2:5], v[0:1], off offset:16 glc
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[0:1], off glc
 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
+; GFX9-NEXT:    global_load_dwordx4 v[2:5], v[0:1], off offset:16 glc
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    ; kill: killed $vgpr0 killed $vgpr1
 ; GFX9-NEXT:  .LBB4_4: ; %exit
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
-; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v7 op_sel_hi:[0,1]
+; GFX9-NEXT:    s_mov_b32 s4, 0xffff
+; GFX9-NEXT:    v_bfi_b32 v0, s4, v6, v6
+; GFX9-NEXT:    v_bfi_b32 v0, s4, v6, v0
+; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v0 op_sel_hi:[0,1]
 ; GFX9-NEXT:    s_movk_i32 s4, 0x8000
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
+; GFX9-NEXT:    v_pk_ashrrev_i16 v2, 15, v7 op_sel_hi:[0,1]
 ; GFX9-NEXT:    v_or_b32_e32 v1, 0xffff8000, v0
-; GFX9-NEXT:    v_or_b32_sdwa v2, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
-; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v6 op_sel_hi:[0,1]
-; GFX9-NEXT:    v_or_b32_e32 v3, 0xffff8000, v0
 ; GFX9-NEXT:    v_or_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
+; GFX9-NEXT:    v_or_b32_e32 v3, 0xffff8000, v2
+; GFX9-NEXT:    v_or_b32_sdwa v2, v2, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    s_mov_b32 s4, 0x5040100
-; GFX9-NEXT:    v_perm_b32 v0, v0, v3, s4
-; GFX9-NEXT:    v_perm_b32 v1, v2, v1, s4
+; GFX9-NEXT:    v_perm_b32 v0, v0, v1, s4
+; GFX9-NEXT:    v_perm_b32 v1, v2, v3, s4
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
   br i1 undef, label %T, label %F
 
 T:
   %t = load volatile <16 x i16>, ptr addrspace(1) %p0
   br label %exit
 
 F:

 ; SI-NEXT:    buffer_load_ushort v4, v[2:3], s[4:7], 0 addr64 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v5, v[2:3], s[4:7], 0 addr64 offset:2 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v6, v[2:3], s[4:7], 0 addr64 offset:4 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:6 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:8 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:8 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:10 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:10 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:12 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:12 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:14 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:14 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:16 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:16 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:18 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:18 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:20 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:20 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:22 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:22 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:24 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:24 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:26 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:26 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v8, v[2:3], s[4:7], 0 addr64 offset:28 glc
+; SI-NEXT:    buffer_load_ushort v7, v[2:3], s[4:7], 0 addr64 offset:28 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v2, v[2:3], s[4:7], 0 addr64 offset:30 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    v_lshlrev_b32_e32 v2, 16, v7
-; SI-NEXT:    v_lshlrev_b32_e32 v7, 16, v5
 ; SI-NEXT:    v_cvt_f32_f16_e32 v3, v5
-; SI-NEXT:    v_or_b32_e32 v2, v6, v2
-; SI-NEXT:    v_or_b32_e32 v4, v4, v7
-; SI-NEXT:    v_cvt_f32_f16_e32 v2, v2
 ; SI-NEXT:    v_cvt_f32_f16_e32 v4, v4
+; SI-NEXT:    v_cvt_f32_f16_e32 v2, v6
 ; SI-NEXT:    s_mov_b64 vcc, exec
 ; SI-NEXT:    s_cbranch_execz .LBB5_3
 ; SI-NEXT:    s_branch .LBB5_4
 ; SI-NEXT:  .LBB5_2:
+; SI-NEXT:    ; implicit-def: $vgpr2
 ; SI-NEXT:    ; implicit-def: $vgpr4
 ; SI-NEXT:    ; implicit-def: $vgpr3
-; SI-NEXT:    ; implicit-def: $vgpr2
 ; SI-NEXT:    s_mov_b64 vcc, 0
 ; SI-NEXT:  .LBB5_3: ; %T
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000
 ; SI-NEXT:    s_mov_b32 s4, s6
 ; SI-NEXT:    s_mov_b32 s5, s6
 ; SI-NEXT:    buffer_load_ushort v2, v[0:1], s[4:7], 0 addr64 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v3, v[0:1], s[4:7], 0 addr64 offset:2 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:4 glc
+; SI-NEXT:    buffer_load_ushort v5, v[0:1], s[4:7], 0 addr64 offset:4 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v5, v[0:1], s[4:7], 0 addr64 offset:6 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:6 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:8 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:8 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:10 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:10 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:12 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:12 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:14 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:14 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:16 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:16 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:18 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:18 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:20 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:20 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:22 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:22 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:24 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:24 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:26 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:26 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    buffer_load_ushort v6, v[0:1], s[4:7], 0 addr64 offset:28 glc
+; SI-NEXT:    buffer_load_ushort v4, v[0:1], s[4:7], 0 addr64 offset:28 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    buffer_load_ushort v0, v[0:1], s[4:7], 0 addr64 offset:30 glc
 ; SI-NEXT:    s_waitcnt vmcnt(0)
-; SI-NEXT:    v_lshlrev_b32_e32 v0, 16, v5
-; SI-NEXT:    v_lshlrev_b32_e32 v1, 16, v3
-; SI-NEXT:    v_or_b32_e32 v0, v4, v0
-; SI-NEXT:    v_or_b32_e32 v1, v2, v1
-; SI-NEXT:    v_cvt_f32_f16_e32 v2, v0
-; SI-NEXT:    v_cvt_f32_f16_e32 v4, v1
 ; SI-NEXT:    v_cvt_f32_f16_e32 v3, v3
+; SI-NEXT:    v_cvt_f32_f16_e32 v4, v2
+; SI-NEXT:    v_cvt_f32_f16_e32 v2, v5
 ; SI-NEXT:  .LBB5_4: ; %exit
 ; SI-NEXT:    v_cvt_f16_f32_e32 v0, v4
 ; SI-NEXT:    v_cvt_f16_f32_e32 v1, v3
 ; SI-NEXT:    v_cvt_f16_f32_e32 v2, v2
 ; SI-NEXT:    v_mov_b32_e32 v3, 0x3fa00000
 ; SI-NEXT:    v_mov_b32_e32 v4, 0x3f200000
 ; SI-NEXT:    v_cvt_f32_f16_e32 v0, v0
 ; SI-NEXT:    v_cvt_f32_f16_e32 v1, v1
 ; SI-NEXT:    v_cvt_f32_f16_e32 v2, v2
 ; SI-NEXT:    v_cmp_nge_f32_e32 vcc, 0.5, v0
 ; SI-NEXT:    v_cndmask_b32_e32 v0, v3, v4, vcc
-; SI-NEXT:    v_cmp_nge_f32_e32 vcc, 0.5, v1
-; SI-NEXT:    v_cndmask_b32_e32 v1, v3, v4, vcc
 ; SI-NEXT:    v_cmp_nge_f32_e32 vcc, 0.5, v2
 ; SI-NEXT:    v_cndmask_b32_e32 v2, v3, v4, vcc
+; SI-NEXT:    v_cmp_nge_f32_e32 vcc, 0.5, v1
+; SI-NEXT:    v_cndmask_b32_e32 v1, v3, v4, vcc
 ; SI-NEXT:    v_mov_b32_e32 v3, v2
 ; SI-NEXT:    s_setpc_b64 s[30:31]
 ;
 ; GFX9-LABEL: vec_16xf16_extract_4xf16:
 ; GFX9:       ; %bb.0:
 ; GFX9-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX9-NEXT:    s_cbranch_scc0 .LBB5_2
 ; GFX9-NEXT:  ; %bb.1: ; %F
-; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[2:3], off offset:16 glc
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[2:3], off glc
 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
+; GFX9-NEXT:    global_load_dwordx4 v[6:9], v[2:3], off offset:16 glc
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    ; kill: killed $vgpr2 killed $vgpr3
 ; GFX9-NEXT:    s_cbranch_execz .LBB5_3
 ; GFX9-NEXT:    s_branch .LBB5_4
 ; GFX9-NEXT:  .LBB5_2:
-; GFX9-NEXT:    ; implicit-def: $vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11
+; GFX9-NEXT:    ; implicit-def: $vgpr5
 ; GFX9-NEXT:  .LBB5_3: ; %T
-; GFX9-NEXT:    global_load_dwordx4 v[2:5], v[0:1], off offset:16 glc
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    global_load_dwordx4 v[4:7], v[0:1], off glc
 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
+; GFX9-NEXT:    global_load_dwordx4 v[6:9], v[0:1], off offset:16 glc
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    ; kill: killed $vgpr0 killed $vgpr1
 ; GFX9-NEXT:  .LBB5_4: ; %exit
+; GFX9-NEXT:    s_mov_b32 s4, 0xffff
+; GFX9-NEXT:    v_bfi_b32 v0, s4, v4, v4
+; GFX9-NEXT:    v_bfi_b32 v0, s4, v4, v0
 ; GFX9-NEXT:    s_mov_b32 s4, 0x5040100
-; GFX9-NEXT:    s_waitcnt vmcnt(0)
-; GFX9-NEXT:    v_perm_b32 v0, v5, v5, s4
-; GFX9-NEXT:    v_mov_b32_e32 v1, 0x3800
-; GFX9-NEXT:    v_mov_b32_e32 v2, 0x3900
-; GFX9-NEXT:    v_mov_b32_e32 v3, 0x3d00
+; GFX9-NEXT:    v_mov_b32_e32 v2, 0x3800
+; GFX9-NEXT:    v_mov_b32_e32 v3, 0x3900
+; GFX9-NEXT:    v_mov_b32_e32 v4, 0x3d00
 ; GFX9-NEXT:    v_cmp_ge_f16_e32 vcc, 0.5, v0
-; GFX9-NEXT:    v_cndmask_b32_e32 v5, v2, v3, vcc
-; GFX9-NEXT:    v_cmp_nle_f16_sdwa vcc, v0, v1 src0_sel:WORD_1 src1_sel:DWORD
-; GFX9-NEXT:    v_cndmask_b32_e32 v6, v3, v2, vcc
-; GFX9-NEXT:    v_cmp_ge_f16_e32 vcc, 0.5, v4
-; GFX9-NEXT:    v_cndmask_b32_e32 v0, v2, v3, vcc
-; GFX9-NEXT:    v_cmp_le_f16_sdwa vcc, v4, v1 src0_sel:WORD_1 src1_sel:DWORD
-; GFX9-NEXT:    v_cndmask_b32_e32 v1, v2, v3, vcc
-; GFX9-NEXT:    v_pack_b32_f16 v0, v0, v1
-; GFX9-NEXT:    v_pack_b32_f16 v1, v5, v6
+; GFX9-NEXT:    v_perm_b32 v1, v5, v5, s4
+; GFX9-NEXT:    v_cndmask_b32_e32 v5, v3, v4, vcc
+; GFX9-NEXT:    v_cmp_le_f16_sdwa vcc, v0, v2 src0_sel:WORD_1 src1_sel:DWORD
+; GFX9-NEXT:    v_cndmask_b32_e32 v0, v3, v4, vcc
+; GFX9-NEXT:    v_cmp_ge_f16_e32 vcc, 0.5, v1
+; GFX9-NEXT:    s_waitcnt vmcnt(0)
+; GFX9-NEXT:    v_cndmask_b32_e32 v6, v3, v4, vcc
+; GFX9-NEXT:    v_cmp_nle_f16_sdwa vcc, v1, v2 src0_sel:WORD_1 src1_sel:DWORD
+; GFX9-NEXT:    v_cndmask_b32_e32 v1, v4, v3, vcc
+; GFX9-NEXT:    v_pack_b32_f16 v0, v5, v0
+; GFX9-NEXT:    v_pack_b32_f16 v1, v6, v1
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
   br i1 undef, label %T, label %F
 
 T:
   %t = load volatile <16 x half>, ptr addrspace(1) %p0
   br label %exit
 
 F:

llvm/test/CodeGen/AMDGPU/mfma-loop.ll

 ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX908,GFX908_A %s
 ; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX90A,GFX908_A,GFX940_A %s
 ; RUN: llc -march=amdgcn -mcpu=gfx940 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX940,GFX940_A %s
 
 ; GCN-LABEL: {{^}}test_mfma_loop_zeroinit:
 
-; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
-; GFX90A:          v_accvgpr_write_b32 [[LEAD:a[0-9]+]], 0
-; GFX90A-COUNT-31: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]]
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 
 ; Check that we do not copy agprs to vgprs and back inside the loop.
 
 ; GCN: [[LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GFX908_A: v_mfma_f32_32x32x1f32
 ; GFX940:   v_mfma_f32_32x32x1_2b_f32
 ; GCN-NOT:  v_accvgpr

 }
 
 ; GCN-LABEL: {{^}}test_mfma_loop_unfoldable_splat:
 
 ; Check that we do not use 32 temp vgprs, but rotate 3 vgprs only.
 ; 3 vgprs are needed to avoid wait states between writes.
 ; Check that we do not use 32 temp sgprs as well.
 
-; GFX908:          v_mov_b32_e32 [[TMP:v[0-9]+]], 0x42f60000
-; GFX940_A:        s_mov_b32 [[TMP:s[0-9]+]], 0x42f60000
-; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
-; GFX90A:          v_accvgpr_write_b32 [[LEAD:a[0-9]+]], [[TMP]]
-; GFX90A-COUNT-31: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]]
+; GCN:          v_mov_b32_e32 [[TMP:v[0-9]+]], 0x42f60000
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
 
 ; GCN: [[LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GFX908_A: v_mfma_f32_32x32x1f32
 ; GFX940:   v_mfma_f32_32x32x1_2b_f32
 ; GCN-NOT:  v_accvgpr
 ; GCN:      s_cbranch_scc1 [[LOOP]]
 

 
 exit:
   store <32 x float> %mai.1, ptr addrspace(1) %arg
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_loop_non_splat:
 
-; GCN:             v_accvgpr_write_b32 [[LEAD:a[0-9]+]], 0{{$}}
+; GCN-COUNT-31:    v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 ; GCN:             v_accvgpr_write_b32 a{{[0-9]+}}, 1.0{{$}}
-; GFX908-COUNT-30: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
-; GFX90A-COUNT-30: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]]{{$}}
 
 ; GCN: [[LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GFX908_A: v_mfma_f32_32x32x1f32
 ; GFX940:   v_mfma_f32_32x32x1_2b_f32
 ; GCN-NOT:  v_accvgpr
 ; GCN:      s_cbranch_scc1 [[LOOP]]
 
exit:
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_loop_unfoldable_seq:
 
 ; Check that we do not use 32 temp vgprs, but rotate 3 vgprs only.
 ; 3 vgprs are needed to avoid wait states between writes.
 
-; GFX908: v_mov_b32_e32 [[TMP1:v[0-9]+]], 0x42f60000
-; GFX908: v_mov_b32_e32 [[TMP2:v[0-9]+]], 0x42f80000
-; GFX908: v_mov_b32_e32 [[TMP3:v[0-9]+]], 0x42fe0000
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
-; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
-; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x42f80000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x42fa0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x42fc0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x42fe0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43000000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43010000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43020000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43030000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43040000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43050000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43060000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43070000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43080000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43090000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x430a0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x430b0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x430c0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x430d0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x430e0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x430f0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43100000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43110000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43120000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43130000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43140000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43150000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43160000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43170000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43180000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x43190000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+; GFX-908: v_mov_b32_e32 v0, 0x431a0000
+; GFX-908: s_nop 1
+; GFX-908: v_accvgpr_write_b32 {{[0-9]+}}, v0
+
+; FIXME: Constant is now in VGPR instead of SGPR.
 
-; GFX940_A-COUNT-32: s_mov_b32 s{{[0-9]+}}, 0x4{{[0-9a-f]+}}
-; GFX940_A-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; GFX940_A: v_mov_b32_e32 v{{[0-9]+}}, 0x4{{[0-9a-f]+}}
+; GFX940_A-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
 
 ; GCN: [[LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GFX908_A: v_mfma_f32_32x32x1f32
 ; GFX940:   v_mfma_f32_32x32x1_2b_f32
 ; GCN-NOT:  v_accvgpr
 ; GCN:      s_cbranch_scc1 [[LOOP]]
 
exit:
   store <32 x float> %mai.1, ptr addrspace(1) %arg
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_loop_sgpr_init:
 
 ; GFX908:          v_mov_b32_e32 [[TMP:v[0-9]+]], s{{[0-9]+}}
 ; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
-; GFX940_A:        v_accvgpr_write_b32 [[LEAD:a[0-9]+]], s{{[0-9]+}}
-; GFX90A-COUNT-31: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]]
+; GFX940_A-COUNT-32:        v_accvgpr_write_b32 [[LEAD:a[0-9]+]], s{{[0-9]+}}
 
 ; GCN: [[LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GFX908_A: v_mfma_f32_32x32x1f32
 ; GFX940:   v_mfma_f32_32x32x1_2b_f32
 ; GCN-NOT:  v_accvgpr
 ; GCN:      s_cbranch_scc1 [[LOOP]]
 

 ; GFX908-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 ; GFX908-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 ; GFX908-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 ; GFX908-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 ; GFX908-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 ; GFX908-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 
 ; GFX90A-DAG:      v_accvgpr_write_b32 [[LEAD:a[0-9]+]], 0
-; GFX90A-COUNT-28: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]]
+; GFX90A-COUNT-28: v_accvgpr_write_b32 a{{[0-9]+}}, 0
 
 ; GCN: [[LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GFX908_A: v_mfma_f32_32x32x1f32
 ; GFX940:   v_mfma_f32_32x32x1_2b_f32
 ; GCN-NOT:  v_accvgpr
 ; GCN:      s_cbranch_scc1 [[LOOP]]
 

 ; GFX908:          v_mfma_f32_32x32x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9:]+}}]
 ; GFX90A-NOT:      v_accvgpr
 ; GFX90A:          v_mfma_f32_32x32x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, 0{{$}}
 ; GFX90A-NOT:      v_accvgpr
 ; GFX940:          v_mfma_f32_32x32x1_2b_f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, 0{{$}}
 
 ; Check that we are using only one tmp VGPR.
 
-; GCN:             v_accvgpr_read_b32 [[TMP:v[0-9]+]], a{{[0-9]+}}
-; GFX908-COUNT-31: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]{{$}}
-; GFX90A:          v_accvgpr_write_b32 [[LEAD:a[0-9]+]], [[TMP]]{{$}}
-; GFX90A-COUNT-29: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]]
+; GFX908:             v_accvgpr_read_b32 [[TMP:v[0-9]+]], a{{[0-9]+}}
+; GFX940_A-COUNT-31:  v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
 
 ; GCN: [[LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GFX908_A: v_mfma_f32_32x32x1f32
 ; GFX940:   v_mfma_f32_32x32x1_2b_f32
 ; GCN-NOT:  v_accvgpr
 ; GCN:      s_cbranch_scc1 [[LOOP]]
 

 
 exit:
   store <32 x float> %mai.1, ptr addrspace(1) %arg
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_nested_loop_zeroinit:
 
-; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
-; GFX90A:          v_accvgpr_write_b32 [[LEAD:a[0-9]+]], 0
-; GFX90A-COUNT-31: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]]
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
 
 ; Check that we do not copy agprs to vgprs and back in an outer loop.
 
 ; GCN: [[OUTER_LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GCN: [[INNER_LOOP:.LBB[0-9_]+]]:
 ; GCN-NOT:  v_accvgpr
 ; GFX908_A: v_mfma_f32_32x32x1f32

llvm/test/CodeGen/AMDGPU/tuple-allocation-failure.ll

 bb:
   tail call void @wobble()
   unreachable
 }
 
 define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i32 %tmp5.i.i, i32 %tmp427.i, i1 %tmp438.i, double %tmp27.i, i1 %tmp48.i) {
 ; GLOBALNESS1-LABEL: kernel:
 ; GLOBALNESS1:       ; %bb.0: ; %bb
-; GLOBALNESS1-NEXT:    s_mov_b64 s[54:55], s[6:7]
-; GLOBALNESS1-NEXT:    s_load_dwordx4 s[36:39], s[8:9], 0x0
+; GLOBALNESS1-NEXT:    s_mov_b64 s[36:37], s[6:7]
+; GLOBALNESS1-NEXT:    s_load_dwordx4 s[76:79], s[8:9], 0x0
 ; GLOBALNESS1-NEXT:    s_load_dword s6, s[8:9], 0x14
-; GLOBALNESS1-NEXT:    v_mov_b32_e32 v43, v0
+; GLOBALNESS1-NEXT:    v_mov_b32_e32 v42, v0
 ; GLOBALNESS1-NEXT:    v_mov_b32_e32 v40, 0
 ; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[0:1], 0, 0
 ; GLOBALNESS1-NEXT:    global_store_dword v[0:1], v40, off
 ; GLOBALNESS1-NEXT:    s_waitcnt lgkmcnt(0)
-; GLOBALNESS1-NEXT:    global_load_dword v0, v40, s[36:37]
-; GLOBALNESS1-NEXT:    s_add_u32 flat_scratch_lo, s12, s17
-; GLOBALNESS1-NEXT:    s_mov_b64 s[64:65], s[4:5]
+; GLOBALNESS1-NEXT:    global_load_dword v0, v40, s[76:77]
+; GLOBALNESS1-NEXT:    s_mov_b64 s[40:41], s[4:5]
 ; GLOBALNESS1-NEXT:    s_load_dwordx2 s[4:5], s[8:9], 0x18
 ; GLOBALNESS1-NEXT:    s_load_dword s7, s[8:9], 0x20
+; GLOBALNESS1-NEXT:    s_add_u32 flat_scratch_lo, s12, s17
 ; GLOBALNESS1-NEXT:    s_addc_u32 flat_scratch_hi, s13, 0
 ; GLOBALNESS1-NEXT:    s_add_u32 s0, s0, s17
 ; GLOBALNESS1-NEXT:    s_addc_u32 s1, s1, 0
 ; GLOBALNESS1-NEXT:    v_mov_b32_e32 v41, 0x40994400
-; GLOBALNESS1-NEXT:    s_bitcmp1_b32 s38, 0
+; GLOBALNESS1-NEXT:    s_bitcmp1_b32 s78, 0
 ; GLOBALNESS1-NEXT:    s_waitcnt lgkmcnt(0)
-; GLOBALNESS1-NEXT:    v_cmp_ngt_f64_e64 s[36:37], s[4:5], v[40:41]
-; GLOBALNESS1-NEXT:    v_cmp_ngt_f64_e64 s[40:41], s[4:5], 0
+; GLOBALNESS1-NEXT:    v_cmp_ngt_f64_e32 vcc, s[4:5], v[40:41]
+; GLOBALNESS1-NEXT:    v_cmp_ngt_f64_e64 s[4:5], s[4:5], 0
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[4:5]
 ; GLOBALNESS1-NEXT:    s_cselect_b64 s[4:5], -1, 0
-; GLOBALNESS1-NEXT:    s_xor_b64 s[94:95], s[4:5], -1
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v3, 0, 1, s[4:5]
+; GLOBALNESS1-NEXT:    s_xor_b64 s[4:5], s[4:5], -1
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v1, 0, 1, vcc
 ; GLOBALNESS1-NEXT:    s_bitcmp1_b32 s6, 0
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[42:43], 1, v1
 ; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v1, 0, 1, s[4:5]
 ; GLOBALNESS1-NEXT:    s_cselect_b64 s[4:5], -1, 0
-; GLOBALNESS1-NEXT:    s_xor_b64 s[88:89], s[4:5], -1
+; GLOBALNESS1-NEXT:    s_xor_b64 s[4:5], s[4:5], -1
 ; GLOBALNESS1-NEXT:    s_bitcmp1_b32 s7, 0
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[48:49], 1, v1
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v1, 0, 1, s[4:5]
 ; GLOBALNESS1-NEXT:    s_cselect_b64 s[4:5], -1, 0
 ; GLOBALNESS1-NEXT:    s_getpc_b64 s[6:7]
 ; GLOBALNESS1-NEXT:    s_add_u32 s6, s6, wobble@gotpcrel32@lo+4
 ; GLOBALNESS1-NEXT:    s_addc_u32 s7, s7, wobble@gotpcrel32@hi+12
-; GLOBALNESS1-NEXT:    s_xor_b64 s[86:87], s[4:5], -1
-; GLOBALNESS1-NEXT:    s_load_dwordx2 s[66:67], s[6:7], 0x0
-; GLOBALNESS1-NEXT:    s_mov_b32 s98, s16
-; GLOBALNESS1-NEXT:    s_mov_b64 s[62:63], s[8:9]
-; GLOBALNESS1-NEXT:    s_mov_b32 s99, s15
-; GLOBALNESS1-NEXT:    s_mov_b32 s100, s14
+; GLOBALNESS1-NEXT:    s_xor_b64 s[4:5], s[4:5], -1
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[50:51], 1, v1
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v1, 0, 1, s[4:5]
+; GLOBALNESS1-NEXT:    s_load_dwordx2 s[76:77], s[6:7], 0x0
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[52:53], 1, v1
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[44:45], 1, v2
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[46:47], 1, v3
+; GLOBALNESS1-NEXT:    s_mov_b32 s70, s16
+; GLOBALNESS1-NEXT:    s_mov_b64 s[38:39], s[8:9]
+; GLOBALNESS1-NEXT:    s_mov_b32 s71, s15
+; GLOBALNESS1-NEXT:    s_mov_b32 s72, s14
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[34:35], s[10:11]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[92:93], 0x80
-; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[42:43], 1, v1
-; GLOBALNESS1-NEXT:    s_mov_b32 s69, 0x3ff00000
+; GLOBALNESS1-NEXT:    s_mov_b64 s[74:75], 0x80
 ; GLOBALNESS1-NEXT:    s_mov_b32 s32, 0
-; GLOBALNESS1-NEXT:    ; implicit-def: $agpr32_agpr33_agpr34_agpr35_agpr36_agpr37_agpr38_agpr39_agpr40_agpr41_agpr42_agpr43_agpr44_agpr45_agpr46_agpr47_agpr48_agpr49_agpr50_agpr51_agpr52_agpr53_agpr54_agpr55_agpr56_agpr57_agpr58_agpr59_agpr60_agpr61_agpr62_agpr63
+; GLOBALNESS1-NEXT:    ; implicit-def: $vgpr44_vgpr45
 ; GLOBALNESS1-NEXT:    s_waitcnt vmcnt(0)
-; GLOBALNESS1-NEXT:    v_cmp_gt_i32_e64 s[4:5], 0, v0
-; GLOBALNESS1-NEXT:    v_writelane_b32 v42, s4, 0
-; GLOBALNESS1-NEXT:    v_writelane_b32 v42, s5, 1
-; GLOBALNESS1-NEXT:    v_cmp_eq_u32_e64 s[4:5], 1, v0
-; GLOBALNESS1-NEXT:    v_writelane_b32 v42, s4, 2
-; GLOBALNESS1-NEXT:    v_writelane_b32 v42, s5, 3
-; GLOBALNESS1-NEXT:    v_cmp_eq_u32_e64 s[4:5], 0, v0
-; GLOBALNESS1-NEXT:    v_writelane_b32 v42, s4, 4
-; GLOBALNESS1-NEXT:    v_cmp_gt_i32_e64 s[90:91], 1, v0
-; GLOBALNESS1-NEXT:    v_writelane_b32 v42, s5, 5
+; GLOBALNESS1-NEXT:    v_cmp_gt_i32_e32 vcc, 0, v0
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v1, 0, 1, vcc
+; GLOBALNESS1-NEXT:    v_cmp_gt_i32_e32 vcc, 1, v0
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v2, 0, 1, vcc
+; GLOBALNESS1-NEXT:    v_cmp_eq_u32_e32 vcc, 1, v0
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v3, 0, 1, vcc
+; GLOBALNESS1-NEXT:    v_cmp_eq_u32_e32 vcc, 0, v0
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e64 v0, 0, 1, vcc
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[54:55], 1, v1
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[56:57], 1, v2
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[58:59], 1, v3
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e64 s[60:61], 1, v0
 ; GLOBALNESS1-NEXT:    s_branch .LBB1_4
 ; GLOBALNESS1-NEXT:  .LBB1_1: ; %bb70.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    v_readlane_b32 s6, v42, 4
-; GLOBALNESS1-NEXT:    v_readlane_b32 s7, v42, 5
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[6:7]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[60:61]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccz .LBB1_29
-; GLOBALNESS1-NEXT:  .LBB1_2: ; %Flow6
+; GLOBALNESS1-NEXT:  .LBB1_2: ; %Flow15
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS1-NEXT:    s_or_b64 exec, exec, s[4:5]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], 0
 ; GLOBALNESS1-NEXT:    ; implicit-def: $sgpr4_sgpr5
-; GLOBALNESS1-NEXT:  .LBB1_3: ; %Flow19
+; GLOBALNESS1-NEXT:  .LBB1_3: ; %Flow28
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a63, v31
 ; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[6:7]
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a62, v30
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a61, v29
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a60, v28
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a59, v27
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a58, v26
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a57, v25
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a56, v24
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a55, v23
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a54, v22
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a53, v21
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a52, v20
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a51, v19
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a50, v18
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a49, v17
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a48, v16
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a47, v15
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a46, v14
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a45, v13
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a44, v12
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a43, v11
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a42, v10
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a41, v9
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a40, v8
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a39, v7
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a38, v6
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a37, v5
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a36, v4
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a35, v3
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a34, v2
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a33, v1
-; GLOBALNESS1-NEXT:    v_accvgpr_write_b32 a32, v0
+; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[44:45], v[0:1], v[0:1] op_sel:[0,1]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_30
 ; GLOBALNESS1-NEXT:  .LBB1_4: ; %bb5
 ; GLOBALNESS1-NEXT:    ; =>This Loop Header: Depth=1
 ; GLOBALNESS1-NEXT:    ; Child Loop BB1_15 Depth 2
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[0:1], s[92:93], s[92:93] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    flat_load_dword v44, v[0:1]
-; GLOBALNESS1-NEXT:    s_add_u32 s8, s62, 40
+; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[0:1], s[74:75], s[74:75] op_sel:[0,1]
+; GLOBALNESS1-NEXT:    flat_load_dword v43, v[0:1]
+; GLOBALNESS1-NEXT:    s_add_u32 s8, s38, 40
 ; GLOBALNESS1-NEXT:    buffer_store_dword v40, off, s[0:3], 0
-; GLOBALNESS1-NEXT:    flat_load_dword v45, v[0:1]
-; GLOBALNESS1-NEXT:    s_addc_u32 s9, s63, 0
-; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[64:65]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[54:55]
+; GLOBALNESS1-NEXT:    flat_load_dword v46, v[0:1]
+; GLOBALNESS1-NEXT:    s_addc_u32 s9, s39, 0
+; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[36:37]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS1-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS1-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS1-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v43
+; GLOBALNESS1-NEXT:    s_mov_b32 s12, s72
+; GLOBALNESS1-NEXT:    s_mov_b32 s13, s71
+; GLOBALNESS1-NEXT:    s_mov_b32 s14, s70
+; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v42
 ; GLOBALNESS1-NEXT:    s_waitcnt lgkmcnt(0)
-; GLOBALNESS1-NEXT:    s_swappc_b64 s[30:31], s[66:67]
-; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[42:43]
+; GLOBALNESS1-NEXT:    s_swappc_b64 s[30:31], s[76:77]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[46:47]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], -1
 ; GLOBALNESS1-NEXT:    ; implicit-def: $sgpr4_sgpr5
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_8
 ; GLOBALNESS1-NEXT:  ; %bb.5: ; %NodeBlock
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    s_cmp_lt_i32 s39, 1
+; GLOBALNESS1-NEXT:    s_cmp_lt_i32 s79, 1
 ; GLOBALNESS1-NEXT:    s_cbranch_scc1 .LBB1_7
-; GLOBALNESS1-NEXT:  ; %bb.6: ; %LeafBlock3
+; GLOBALNESS1-NEXT:  ; %bb.6: ; %LeafBlock12
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    s_cmp_lg_u32 s39, 1
+; GLOBALNESS1-NEXT:    s_cmp_lg_u32 s79, 1
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], -1
 ; GLOBALNESS1-NEXT:    s_cselect_b64 s[6:7], -1, 0
 ; GLOBALNESS1-NEXT:    s_cbranch_execnz .LBB1_8
 ; GLOBALNESS1-NEXT:    s_branch .LBB1_23
 ; GLOBALNESS1-NEXT:  .LBB1_7: ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], 0
 ; GLOBALNESS1-NEXT:    ; implicit-def: $sgpr4_sgpr5
 ; GLOBALNESS1-NEXT:    s_branch .LBB1_23
-; GLOBALNESS1-NEXT:  .LBB1_8: ; %Flow16
+; GLOBALNESS1-NEXT:  .LBB1_8: ; %Flow25
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[6:7]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccz .LBB1_24
 ; GLOBALNESS1-NEXT:  .LBB1_9: ; %baz.exit.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[32:33], 0, 0
-; GLOBALNESS1-NEXT:    flat_load_dword v0, v[32:33]
-; GLOBALNESS1-NEXT:    s_mov_b32 s68, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s70, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s71, s69
-; GLOBALNESS1-NEXT:    s_mov_b32 s72, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s73, s69
-; GLOBALNESS1-NEXT:    s_mov_b32 s74, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s75, s69
-; GLOBALNESS1-NEXT:    s_mov_b32 s76, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s77, s69
-; GLOBALNESS1-NEXT:    s_mov_b32 s78, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s79, s69
-; GLOBALNESS1-NEXT:    s_mov_b32 s80, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s81, s69
-; GLOBALNESS1-NEXT:    s_mov_b32 s82, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s83, s69
-; GLOBALNESS1-NEXT:    s_mov_b32 s84, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s85, s69
+; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[2:3], 0, 0
+; GLOBALNESS1-NEXT:    flat_load_dword v0, v[2:3]
 ; GLOBALNESS1-NEXT:    s_waitcnt vmcnt(0) lgkmcnt(0)
-; GLOBALNESS1-NEXT:    v_cmp_gt_i32_e64 s[96:97], 0, v0
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[0:1], s[68:69], s[68:69] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[2:3], s[70:71], s[70:71] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[4:5], s[72:73], s[72:73] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[6:7], s[74:75], s[74:75] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[8:9], s[76:77], s[76:77] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[10:11], s[78:79], s[78:79] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[12:13], s[80:81], s[80:81] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[14:15], s[82:83], s[82:83] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[16:17], s[84:85], s[84:85] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[18:19], s[86:87], s[86:87] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[20:21], s[88:89], s[88:89] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[22:23], s[90:91], s[90:91] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[24:25], s[92:93], s[92:93] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[26:27], s[94:95], s[94:95] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[28:29], s[96:97], s[96:97] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[30:31], s[98:99], s[98:99] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    s_and_saveexec_b64 s[70:71], s[96:97]
+; GLOBALNESS1-NEXT:    v_cmp_gt_i32_e64 s[62:63], 0, v0
+; GLOBALNESS1-NEXT:    v_mov_b32_e32 v0, 0
+; GLOBALNESS1-NEXT:    v_mov_b32_e32 v1, 0x3ff00000
+; GLOBALNESS1-NEXT:    s_and_saveexec_b64 s[80:81], s[62:63]
 ; GLOBALNESS1-NEXT:    s_cbranch_execz .LBB1_26
 ; GLOBALNESS1-NEXT:  ; %bb.10: ; %bb33.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    global_load_dwordx2 v[0:1], v[32:33], off
-; GLOBALNESS1-NEXT:    v_readlane_b32 s4, v42, 0
-; GLOBALNESS1-NEXT:    v_readlane_b32 s5, v42, 1
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[4:5]
+; GLOBALNESS1-NEXT:    global_load_dwordx2 v[0:1], v[2:3], off
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[54:55]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_12
 ; GLOBALNESS1-NEXT:  ; %bb.11: ; %bb39.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS1-NEXT:    v_mov_b32_e32 v41, v40
 ; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[2:3], 0, 0
 ; GLOBALNESS1-NEXT:    global_store_dwordx2 v[2:3], v[40:41], off
 ; GLOBALNESS1-NEXT:  .LBB1_12: ; %bb44.lr.ph.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e32 vcc, 0, v45
-; GLOBALNESS1-NEXT:    v_cndmask_b32_e32 v2, 0, v44, vcc
-; GLOBALNESS1-NEXT:    s_mov_b64 s[72:73], s[42:43]
-; GLOBALNESS1-NEXT:    s_mov_b32 s75, s39
+; GLOBALNESS1-NEXT:    v_cmp_ne_u32_e32 vcc, 0, v46
+; GLOBALNESS1-NEXT:    v_cndmask_b32_e32 v2, 0, v43, vcc
 ; GLOBALNESS1-NEXT:    s_waitcnt vmcnt(0)
-; GLOBALNESS1-NEXT:    v_cmp_nlt_f64_e64 s[56:57], 0, v[0:1]
-; GLOBALNESS1-NEXT:    v_cmp_eq_u32_e64 s[58:59], 0, v2
+; GLOBALNESS1-NEXT:    v_cmp_nlt_f64_e64 s[64:65], 0, v[0:1]
+; GLOBALNESS1-NEXT:    v_cmp_eq_u32_e64 s[66:67], 0, v2
 ; GLOBALNESS1-NEXT:    s_branch .LBB1_15
-; GLOBALNESS1-NEXT:  .LBB1_13: ; %Flow7
+; GLOBALNESS1-NEXT:  .LBB1_13: ; %Flow16
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
 ; GLOBALNESS1-NEXT:    s_or_b64 exec, exec, s[4:5]
 ; GLOBALNESS1-NEXT:  .LBB1_14: ; %bb63.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[86:87]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[52:53]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccz .LBB1_25
 ; GLOBALNESS1-NEXT:  .LBB1_15: ; %bb44.i
 ; GLOBALNESS1-NEXT:    ; Parent Loop BB1_4 Depth=1
 ; GLOBALNESS1-NEXT:    ; => This Inner Loop Header: Depth=2
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[94:95]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[48:49]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_14
 ; GLOBALNESS1-NEXT:  ; %bb.16: ; %bb46.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[88:89]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[50:51]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_14
 ; GLOBALNESS1-NEXT:  ; %bb.17: ; %bb50.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[36:37]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[42:43]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_20
 ; GLOBALNESS1-NEXT:  ; %bb.18: ; %bb3.i.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[40:41]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[44:45]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_20
 ; GLOBALNESS1-NEXT:  ; %bb.19: ; %bb6.i.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[56:57]
+; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[64:65]
 ; GLOBALNESS1-NEXT:  .LBB1_20: ; %spam.exit.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[90:91]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[56:57]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_14
 ; GLOBALNESS1-NEXT:  ; %bb.21: ; %bb55.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS1-NEXT:    s_add_u32 s60, s62, 40
-; GLOBALNESS1-NEXT:    s_addc_u32 s61, s63, 0
-; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[64:65]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[54:55]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[8:9], s[60:61]
+; GLOBALNESS1-NEXT:    s_add_u32 s68, s38, 40
+; GLOBALNESS1-NEXT:    s_addc_u32 s69, s39, 0
+; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[36:37]
+; GLOBALNESS1-NEXT:    s_mov_b64 s[8:9], s[68:69]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS1-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS1-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS1-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v43
-; GLOBALNESS1-NEXT:    s_swappc_b64 s[30:31], s[66:67]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[44:45], 0, 0
-; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[64:65]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[54:55]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[8:9], s[60:61]
+; GLOBALNESS1-NEXT:    s_mov_b32 s12, s72
+; GLOBALNESS1-NEXT:    s_mov_b32 s13, s71
+; GLOBALNESS1-NEXT:    s_mov_b32 s14, s70
+; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v42
+; GLOBALNESS1-NEXT:    s_swappc_b64 s[30:31], s[76:77]
+; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[46:47], 0, 0
+; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[36:37]
+; GLOBALNESS1-NEXT:    s_mov_b64 s[8:9], s[68:69]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS1-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS1-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS1-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v43
-; GLOBALNESS1-NEXT:    global_store_dwordx2 v[44:45], a[32:33], off
-; GLOBALNESS1-NEXT:    s_swappc_b64 s[30:31], s[66:67]
-; GLOBALNESS1-NEXT:    s_and_saveexec_b64 s[4:5], s[58:59]
+; GLOBALNESS1-NEXT:    s_mov_b32 s12, s72
+; GLOBALNESS1-NEXT:    s_mov_b32 s13, s71
+; GLOBALNESS1-NEXT:    s_mov_b32 s14, s70
+; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v42
+; GLOBALNESS1-NEXT:    global_store_dwordx2 v[46:47], v[44:45], off
+; GLOBALNESS1-NEXT:    s_swappc_b64 s[30:31], s[76:77]
+; GLOBALNESS1-NEXT:    s_and_saveexec_b64 s[4:5], s[66:67]
 ; GLOBALNESS1-NEXT:    s_cbranch_execz .LBB1_13
 ; GLOBALNESS1-NEXT:  ; %bb.22: ; %bb62.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_15 Depth=2
 ; GLOBALNESS1-NEXT:    v_mov_b32_e32 v41, v40
-; GLOBALNESS1-NEXT:    global_store_dwordx2 v[44:45], v[40:41], off
+; GLOBALNESS1-NEXT:    global_store_dwordx2 v[46:47], v[40:41], off
 ; GLOBALNESS1-NEXT:    s_branch .LBB1_13
 ; GLOBALNESS1-NEXT:  .LBB1_23: ; %LeafBlock
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    s_cmp_lg_u32 s39, 0
+; GLOBALNESS1-NEXT:    s_cmp_lg_u32 s79, 0
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], 0
 ; GLOBALNESS1-NEXT:    s_cselect_b64 s[6:7], -1, 0
 ; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[6:7]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_9
 ; GLOBALNESS1-NEXT:  .LBB1_24: ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], -1
-; GLOBALNESS1-NEXT:    ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31
+; GLOBALNESS1-NEXT:    ; implicit-def: $vgpr0_vgpr1
 ; GLOBALNESS1-NEXT:    s_branch .LBB1_3
-; GLOBALNESS1-NEXT:  .LBB1_25: ; %Flow14
+; GLOBALNESS1-NEXT:  .LBB1_25: ; %Flow23
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[36:37]
-; GLOBALNESS1-NEXT:    s_mov_b32 s36, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s37, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s38, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s39, s93
-; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[40:41]
-; GLOBALNESS1-NEXT:    s_mov_b32 s40, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s41, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s42, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s43, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s44, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s45, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s46, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s47, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s48, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s49, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s50, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s51, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s52, s93
-; GLOBALNESS1-NEXT:    s_mov_b32 s53, s93
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[0:1], s[36:37], s[36:37] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[2:3], s[38:39], s[38:39] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[4:5], s[40:41], s[40:41] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[6:7], s[42:43], s[42:43] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[8:9], s[44:45], s[44:45] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[10:11], s[46:47], s[46:47] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[12:13], s[48:49], s[48:49] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[14:15], s[50:51], s[50:51] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[16:17], s[52:53], s[52:53] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[18:19], s[54:55], s[54:55] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[20:21], s[56:57], s[56:57] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[22:23], s[58:59], s[58:59] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[24:25], s[60:61], s[60:61] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[26:27], s[62:63], s[62:63] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[28:29], s[64:65], s[64:65] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[30:31], s[66:67], s[66:67] op_sel:[0,1]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[40:41], s[6:7]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[36:37], s[4:5]
-; GLOBALNESS1-NEXT:    s_mov_b32 s39, s75
-; GLOBALNESS1-NEXT:    s_mov_b64 s[42:43], s[72:73]
-; GLOBALNESS1-NEXT:  .LBB1_26: ; %Flow15
+; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[0:1], 0, 0
+; GLOBALNESS1-NEXT:  .LBB1_26: ; %Flow24
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    s_or_b64 exec, exec, s[70:71]
-; GLOBALNESS1-NEXT:    s_and_saveexec_b64 s[4:5], s[96:97]
+; GLOBALNESS1-NEXT:    s_or_b64 exec, exec, s[80:81]
+; GLOBALNESS1-NEXT:    s_and_saveexec_b64 s[4:5], s[62:63]
 ; GLOBALNESS1-NEXT:    s_cbranch_execz .LBB1_2
 ; GLOBALNESS1-NEXT:  ; %bb.27: ; %bb67.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS1-NEXT:    v_readlane_b32 s6, v42, 2
-; GLOBALNESS1-NEXT:    v_readlane_b32 s7, v42, 3
-; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[6:7]
+; GLOBALNESS1-NEXT:    s_and_b64 vcc, exec, s[58:59]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_1
 ; GLOBALNESS1-NEXT:  ; %bb.28: ; %bb69.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS1-NEXT:    v_mov_b32_e32 v41, v40
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[32:33], 0, 0
-; GLOBALNESS1-NEXT:    global_store_dwordx2 v[32:33], v[40:41], off
+; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[2:3], 0, 0
+; GLOBALNESS1-NEXT:    global_store_dwordx2 v[2:3], v[40:41], off
 ; GLOBALNESS1-NEXT:    s_branch .LBB1_1
 ; GLOBALNESS1-NEXT:  .LBB1_29: ; %bb73.i
 ; GLOBALNESS1-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS1-NEXT:    v_mov_b32_e32 v41, v40
-; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[32:33], 0, 0
-; GLOBALNESS1-NEXT:    global_store_dwordx2 v[32:33], v[40:41], off
+; GLOBALNESS1-NEXT:    v_pk_mov_b32 v[2:3], 0, 0
+; GLOBALNESS1-NEXT:    global_store_dwordx2 v[2:3], v[40:41], off
 ; GLOBALNESS1-NEXT:    s_branch .LBB1_2
 ; GLOBALNESS1-NEXT:  .LBB1_30: ; %loop.exit.guard
 ; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[4:5]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], -1
 ; GLOBALNESS1-NEXT:    s_cbranch_vccz .LBB1_32
 ; GLOBALNESS1-NEXT:  ; %bb.31: ; %bb7.i.i
-; GLOBALNESS1-NEXT:    s_add_u32 s8, s62, 40
-; GLOBALNESS1-NEXT:    s_addc_u32 s9, s63, 0
-; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[64:65]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[54:55]
+; GLOBALNESS1-NEXT:    s_add_u32 s8, s38, 40
+; GLOBALNESS1-NEXT:    s_addc_u32 s9, s39, 0
+; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[36:37]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS1-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS1-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS1-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v43
+; GLOBALNESS1-NEXT:    s_mov_b32 s12, s72
+; GLOBALNESS1-NEXT:    s_mov_b32 s13, s71
+; GLOBALNESS1-NEXT:    s_mov_b32 s14, s70
+; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v42
 ; GLOBALNESS1-NEXT:    s_getpc_b64 s[16:17]
 ; GLOBALNESS1-NEXT:    s_add_u32 s16, s16, widget@rel32@lo+4
 ; GLOBALNESS1-NEXT:    s_addc_u32 s17, s17, widget@rel32@hi+12
 ; GLOBALNESS1-NEXT:    s_swappc_b64 s[30:31], s[16:17]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], 0
 ; GLOBALNESS1-NEXT:  .LBB1_32: ; %Flow
 ; GLOBALNESS1-NEXT:    s_andn2_b64 vcc, exec, s[4:5]
 ; GLOBALNESS1-NEXT:    s_cbranch_vccnz .LBB1_34
 ; GLOBALNESS1-NEXT:  ; %bb.33: ; %bb11.i.i
-; GLOBALNESS1-NEXT:    s_add_u32 s8, s62, 40
-; GLOBALNESS1-NEXT:    s_addc_u32 s9, s63, 0
-; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[64:65]
-; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[54:55]
+; GLOBALNESS1-NEXT:    s_add_u32 s8, s38, 40
+; GLOBALNESS1-NEXT:    s_addc_u32 s9, s39, 0
+; GLOBALNESS1-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS1-NEXT:    s_mov_b64 s[6:7], s[36:37]
 ; GLOBALNESS1-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS1-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS1-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS1-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v43
+; GLOBALNESS1-NEXT:    s_mov_b32 s12, s72
+; GLOBALNESS1-NEXT:    s_mov_b32 s13, s71
+; GLOBALNESS1-NEXT:    s_mov_b32 s14, s70
+; GLOBALNESS1-NEXT:    v_mov_b32_e32 v31, v42
 ; GLOBALNESS1-NEXT:    s_getpc_b64 s[16:17]
 ; GLOBALNESS1-NEXT:    s_add_u32 s16, s16, widget@rel32@lo+4
 ; GLOBALNESS1-NEXT:    s_addc_u32 s17, s17, widget@rel32@hi+12
 ; GLOBALNESS1-NEXT:    s_swappc_b64 s[30:31], s[16:17]
 ; GLOBALNESS1-NEXT:  .LBB1_34: ; %UnifiedUnreachableBlock
 ;
 ; GLOBALNESS0-LABEL: kernel:
 ; GLOBALNESS0:       ; %bb.0: ; %bb
-; GLOBALNESS0-NEXT:    s_mov_b64 s[54:55], s[6:7]
-; GLOBALNESS0-NEXT:    s_load_dwordx4 s[36:39], s[8:9], 0x0
+; GLOBALNESS0-NEXT:    s_mov_b64 s[36:37], s[6:7]
+; GLOBALNESS0-NEXT:    s_load_dwordx4 s[72:75], s[8:9], 0x0
 ; GLOBALNESS0-NEXT:    s_load_dword s6, s[8:9], 0x14
-; GLOBALNESS0-NEXT:    v_mov_b32_e32 v43, v0
+; GLOBALNESS0-NEXT:    v_mov_b32_e32 v42, v0
 ; GLOBALNESS0-NEXT:    v_mov_b32_e32 v40, 0
 ; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[0:1], 0, 0
 ; GLOBALNESS0-NEXT:    global_store_dword v[0:1], v40, off
 ; GLOBALNESS0-NEXT:    s_waitcnt lgkmcnt(0)
-; GLOBALNESS0-NEXT:    global_load_dword v0, v40, s[36:37]
-; GLOBALNESS0-NEXT:    s_add_u32 flat_scratch_lo, s12, s17
-; GLOBALNESS0-NEXT:    s_mov_b64 s[62:63], s[4:5]
+; GLOBALNESS0-NEXT:    global_load_dword v0, v40, s[72:73]
+; GLOBALNESS0-NEXT:    s_mov_b64 s[40:41], s[4:5]
 ; GLOBALNESS0-NEXT:    s_load_dwordx2 s[4:5], s[8:9], 0x18
 ; GLOBALNESS0-NEXT:    s_load_dword s7, s[8:9], 0x20
+; GLOBALNESS0-NEXT:    s_add_u32 flat_scratch_lo, s12, s17
 ; GLOBALNESS0-NEXT:    s_addc_u32 flat_scratch_hi, s13, 0
 ; GLOBALNESS0-NEXT:    s_add_u32 s0, s0, s17
 ; GLOBALNESS0-NEXT:    s_addc_u32 s1, s1, 0
 ; GLOBALNESS0-NEXT:    v_mov_b32_e32 v41, 0x40994400
-; GLOBALNESS0-NEXT:    s_bitcmp1_b32 s38, 0
+; GLOBALNESS0-NEXT:    s_bitcmp1_b32 s74, 0
 ; GLOBALNESS0-NEXT:    s_waitcnt lgkmcnt(0)
-; GLOBALNESS0-NEXT:    v_cmp_ngt_f64_e64 s[36:37], s[4:5], v[40:41]
-; GLOBALNESS0-NEXT:    v_cmp_ngt_f64_e64 s[40:41], s[4:5], 0
+; GLOBALNESS0-NEXT:    v_cmp_ngt_f64_e32 vcc, s[4:5], v[40:41]
+; GLOBALNESS0-NEXT:    v_cmp_ngt_f64_e64 s[4:5], s[4:5], 0
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[4:5]
 ; GLOBALNESS0-NEXT:    s_cselect_b64 s[4:5], -1, 0
-; GLOBALNESS0-NEXT:    s_xor_b64 s[94:95], s[4:5], -1
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v3, 0, 1, s[4:5]
+; GLOBALNESS0-NEXT:    s_xor_b64 s[4:5], s[4:5], -1
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v1, 0, 1, vcc
 ; GLOBALNESS0-NEXT:    s_bitcmp1_b32 s6, 0
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[42:43], 1, v1
 ; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v1, 0, 1, s[4:5]
 ; GLOBALNESS0-NEXT:    s_cselect_b64 s[4:5], -1, 0
-; GLOBALNESS0-NEXT:    s_xor_b64 s[88:89], s[4:5], -1
+; GLOBALNESS0-NEXT:    s_xor_b64 s[4:5], s[4:5], -1
 ; GLOBALNESS0-NEXT:    s_bitcmp1_b32 s7, 0
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[48:49], 1, v1
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v1, 0, 1, s[4:5]
 ; GLOBALNESS0-NEXT:    s_cselect_b64 s[4:5], -1, 0
 ; GLOBALNESS0-NEXT:    s_getpc_b64 s[6:7]
 ; GLOBALNESS0-NEXT:    s_add_u32 s6, s6, wobble@gotpcrel32@lo+4
 ; GLOBALNESS0-NEXT:    s_addc_u32 s7, s7, wobble@gotpcrel32@hi+12
-; GLOBALNESS0-NEXT:    s_xor_b64 s[86:87], s[4:5], -1
-; GLOBALNESS0-NEXT:    s_load_dwordx2 s[66:67], s[6:7], 0x0
-; GLOBALNESS0-NEXT:    s_mov_b32 s98, s16
-; GLOBALNESS0-NEXT:    s_mov_b64 s[60:61], s[8:9]
-; GLOBALNESS0-NEXT:    s_mov_b32 s99, s15
-; GLOBALNESS0-NEXT:    s_mov_b32 s100, s14
+; GLOBALNESS0-NEXT:    s_xor_b64 s[4:5], s[4:5], -1
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[50:51], 1, v1
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v1, 0, 1, s[4:5]
+; GLOBALNESS0-NEXT:    s_load_dwordx2 s[78:79], s[6:7], 0x0
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[52:53], 1, v1
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[44:45], 1, v2
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[46:47], 1, v3
+; GLOBALNESS0-NEXT:    s_mov_b32 s68, s16
+; GLOBALNESS0-NEXT:    s_mov_b64 s[38:39], s[8:9]
+; GLOBALNESS0-NEXT:    s_mov_b32 s69, s15
+; GLOBALNESS0-NEXT:    s_mov_b32 s70, s14
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[34:35], s[10:11]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[92:93], 0x80
-; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[42:43], 1, v1
-; GLOBALNESS0-NEXT:    s_mov_b32 s69, 0x3ff00000
+; GLOBALNESS0-NEXT:    s_mov_b64 s[76:77], 0x80
 ; GLOBALNESS0-NEXT:    s_mov_b32 s32, 0
-; GLOBALNESS0-NEXT:    ; implicit-def: $agpr32_agpr33_agpr34_agpr35_agpr36_agpr37_agpr38_agpr39_agpr40_agpr41_agpr42_agpr43_agpr44_agpr45_agpr46_agpr47_agpr48_agpr49_agpr50_agpr51_agpr52_agpr53_agpr54_agpr55_agpr56_agpr57_agpr58_agpr59_agpr60_agpr61_agpr62_agpr63
+; GLOBALNESS0-NEXT:    ; implicit-def: $vgpr44_vgpr45
 ; GLOBALNESS0-NEXT:    s_waitcnt vmcnt(0)
-; GLOBALNESS0-NEXT:    v_cmp_gt_i32_e64 s[4:5], 0, v0
-; GLOBALNESS0-NEXT:    v_writelane_b32 v42, s4, 0
-; GLOBALNESS0-NEXT:    v_writelane_b32 v42, s5, 1
-; GLOBALNESS0-NEXT:    v_cmp_eq_u32_e64 s[4:5], 1, v0
-; GLOBALNESS0-NEXT:    v_writelane_b32 v42, s4, 2
-; GLOBALNESS0-NEXT:    v_writelane_b32 v42, s5, 3
-; GLOBALNESS0-NEXT:    v_cmp_eq_u32_e64 s[4:5], 0, v0
-; GLOBALNESS0-NEXT:    v_writelane_b32 v42, s4, 4
-; GLOBALNESS0-NEXT:    v_cmp_gt_i32_e64 s[90:91], 1, v0
-; GLOBALNESS0-NEXT:    v_writelane_b32 v42, s5, 5
+; GLOBALNESS0-NEXT:    v_cmp_gt_i32_e32 vcc, 0, v0
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v1, 0, 1, vcc
+; GLOBALNESS0-NEXT:    v_cmp_gt_i32_e32 vcc, 1, v0
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v2, 0, 1, vcc
+; GLOBALNESS0-NEXT:    v_cmp_eq_u32_e32 vcc, 1, v0
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v3, 0, 1, vcc
+; GLOBALNESS0-NEXT:    v_cmp_eq_u32_e32 vcc, 0, v0
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e64 v0, 0, 1, vcc
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[54:55], 1, v1
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[56:57], 1, v2
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[58:59], 1, v3
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e64 s[60:61], 1, v0
 ; GLOBALNESS0-NEXT:    s_branch .LBB1_4
 ; GLOBALNESS0-NEXT:  .LBB1_1: ; %bb70.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    v_readlane_b32 s6, v42, 4
-; GLOBALNESS0-NEXT:    v_readlane_b32 s7, v42, 5
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[6:7]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[60:61]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccz .LBB1_29
-; GLOBALNESS0-NEXT:  .LBB1_2: ; %Flow6
+; GLOBALNESS0-NEXT:  .LBB1_2: ; %Flow15
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS0-NEXT:    s_or_b64 exec, exec, s[4:5]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], 0
 ; GLOBALNESS0-NEXT:    ; implicit-def: $sgpr4_sgpr5
-; GLOBALNESS0-NEXT:  .LBB1_3: ; %Flow19
+; GLOBALNESS0-NEXT:  .LBB1_3: ; %Flow28
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a63, v31
 ; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[6:7]
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a62, v30
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a61, v29
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a60, v28
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a59, v27
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a58, v26
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a57, v25
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a56, v24
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a55, v23
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a54, v22
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a53, v21
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a52, v20
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a51, v19
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a50, v18
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a49, v17
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a48, v16
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a47, v15
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a46, v14
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a45, v13
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a44, v12
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a43, v11
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a42, v10
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a41, v9
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a40, v8
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a39, v7
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a38, v6
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a37, v5
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a36, v4
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a35, v3
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a34, v2
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a33, v1
-; GLOBALNESS0-NEXT:    v_accvgpr_write_b32 a32, v0
+; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[44:45], v[0:1], v[0:1] op_sel:[0,1]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_30
 ; GLOBALNESS0-NEXT:  .LBB1_4: ; %bb5
 ; GLOBALNESS0-NEXT:    ; =>This Loop Header: Depth=1
 ; GLOBALNESS0-NEXT:    ; Child Loop BB1_15 Depth 2
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[0:1], s[92:93], s[92:93] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    flat_load_dword v44, v[0:1]
-; GLOBALNESS0-NEXT:    s_add_u32 s8, s60, 40
+; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[0:1], s[76:77], s[76:77] op_sel:[0,1]
+; GLOBALNESS0-NEXT:    flat_load_dword v43, v[0:1]
+; GLOBALNESS0-NEXT:    s_add_u32 s8, s38, 40
 ; GLOBALNESS0-NEXT:    buffer_store_dword v40, off, s[0:3], 0
-; GLOBALNESS0-NEXT:    flat_load_dword v45, v[0:1]
-; GLOBALNESS0-NEXT:    s_addc_u32 s9, s61, 0
-; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[62:63]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[54:55]
+; GLOBALNESS0-NEXT:    flat_load_dword v46, v[0:1]
+; GLOBALNESS0-NEXT:    s_addc_u32 s9, s39, 0
+; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[36:37]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS0-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS0-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS0-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v43
+; GLOBALNESS0-NEXT:    s_mov_b32 s12, s70
+; GLOBALNESS0-NEXT:    s_mov_b32 s13, s69
+; GLOBALNESS0-NEXT:    s_mov_b32 s14, s68
+; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v42
 ; GLOBALNESS0-NEXT:    s_waitcnt lgkmcnt(0)
-; GLOBALNESS0-NEXT:    s_swappc_b64 s[30:31], s[66:67]
-; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[42:43]
+; GLOBALNESS0-NEXT:    s_swappc_b64 s[30:31], s[78:79]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[46:47]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], -1
 ; GLOBALNESS0-NEXT:    ; implicit-def: $sgpr4_sgpr5
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_8
 ; GLOBALNESS0-NEXT:  ; %bb.5: ; %NodeBlock
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    s_cmp_lt_i32 s39, 1
+; GLOBALNESS0-NEXT:    s_cmp_lt_i32 s75, 1
 ; GLOBALNESS0-NEXT:    s_cbranch_scc1 .LBB1_7
-; GLOBALNESS0-NEXT:  ; %bb.6: ; %LeafBlock3
+; GLOBALNESS0-NEXT:  ; %bb.6: ; %LeafBlock12
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    s_cmp_lg_u32 s39, 1
+; GLOBALNESS0-NEXT:    s_cmp_lg_u32 s75, 1
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], -1
 ; GLOBALNESS0-NEXT:    s_cselect_b64 s[6:7], -1, 0
 ; GLOBALNESS0-NEXT:    s_cbranch_execnz .LBB1_8
 ; GLOBALNESS0-NEXT:    s_branch .LBB1_23
 ; GLOBALNESS0-NEXT:  .LBB1_7: ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], 0
 ; GLOBALNESS0-NEXT:    ; implicit-def: $sgpr4_sgpr5
 ; GLOBALNESS0-NEXT:    s_branch .LBB1_23
-; GLOBALNESS0-NEXT:  .LBB1_8: ; %Flow16
+; GLOBALNESS0-NEXT:  .LBB1_8: ; %Flow25
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[6:7]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccz .LBB1_24
 ; GLOBALNESS0-NEXT:  .LBB1_9: ; %baz.exit.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[32:33], 0, 0
-; GLOBALNESS0-NEXT:    flat_load_dword v0, v[32:33]
-; GLOBALNESS0-NEXT:    s_mov_b32 s68, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s70, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s71, s69
-; GLOBALNESS0-NEXT:    s_mov_b32 s72, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s73, s69
-; GLOBALNESS0-NEXT:    s_mov_b32 s74, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s75, s69
-; GLOBALNESS0-NEXT:    s_mov_b32 s76, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s77, s69
-; GLOBALNESS0-NEXT:    s_mov_b32 s78, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s79, s69
-; GLOBALNESS0-NEXT:    s_mov_b32 s80, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s81, s69
-; GLOBALNESS0-NEXT:    s_mov_b32 s82, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s83, s69
-; GLOBALNESS0-NEXT:    s_mov_b32 s84, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s85, s69
+; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[2:3], 0, 0
+; GLOBALNESS0-NEXT:    flat_load_dword v0, v[2:3]
 ; GLOBALNESS0-NEXT:    s_waitcnt vmcnt(0) lgkmcnt(0)
-; GLOBALNESS0-NEXT:    v_cmp_gt_i32_e64 s[96:97], 0, v0
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[0:1], s[68:69], s[68:69] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[2:3], s[70:71], s[70:71] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[4:5], s[72:73], s[72:73] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[6:7], s[74:75], s[74:75] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[8:9], s[76:77], s[76:77] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[10:11], s[78:79], s[78:79] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[12:13], s[80:81], s[80:81] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[14:15], s[82:83], s[82:83] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[16:17], s[84:85], s[84:85] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[18:19], s[86:87], s[86:87] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[20:21], s[88:89], s[88:89] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[22:23], s[90:91], s[90:91] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[24:25], s[92:93], s[92:93] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[26:27], s[94:95], s[94:95] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[28:29], s[96:97], s[96:97] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[30:31], s[98:99], s[98:99] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    s_and_saveexec_b64 s[70:71], s[96:97]
+; GLOBALNESS0-NEXT:    v_cmp_gt_i32_e64 s[62:63], 0, v0
+; GLOBALNESS0-NEXT:    v_mov_b32_e32 v0, 0
+; GLOBALNESS0-NEXT:    v_mov_b32_e32 v1, 0x3ff00000
+; GLOBALNESS0-NEXT:    s_and_saveexec_b64 s[80:81], s[62:63]
 ; GLOBALNESS0-NEXT:    s_cbranch_execz .LBB1_26
 ; GLOBALNESS0-NEXT:  ; %bb.10: ; %bb33.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    global_load_dwordx2 v[0:1], v[32:33], off
-; GLOBALNESS0-NEXT:    v_readlane_b32 s4, v42, 0
-; GLOBALNESS0-NEXT:    v_readlane_b32 s5, v42, 1
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[4:5]
+; GLOBALNESS0-NEXT:    global_load_dwordx2 v[0:1], v[2:3], off
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[54:55]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_12
 ; GLOBALNESS0-NEXT:  ; %bb.11: ; %bb39.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS0-NEXT:    v_mov_b32_e32 v41, v40
 ; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[2:3], 0, 0
 ; GLOBALNESS0-NEXT:    global_store_dwordx2 v[2:3], v[40:41], off
 ; GLOBALNESS0-NEXT:  .LBB1_12: ; %bb44.lr.ph.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e32 vcc, 0, v45
-; GLOBALNESS0-NEXT:    v_cndmask_b32_e32 v2, 0, v44, vcc
-; GLOBALNESS0-NEXT:    s_mov_b64 s[72:73], s[42:43]
-; GLOBALNESS0-NEXT:    s_mov_b32 s75, s39
+; GLOBALNESS0-NEXT:    v_cmp_ne_u32_e32 vcc, 0, v46
+; GLOBALNESS0-NEXT:    v_cndmask_b32_e32 v2, 0, v43, vcc
 ; GLOBALNESS0-NEXT:    s_waitcnt vmcnt(0)
-; GLOBALNESS0-NEXT:    v_cmp_nlt_f64_e64 s[56:57], 0, v[0:1]
-; GLOBALNESS0-NEXT:    v_cmp_eq_u32_e64 s[58:59], 0, v2
+; GLOBALNESS0-NEXT:    v_cmp_nlt_f64_e64 s[64:65], 0, v[0:1]
+; GLOBALNESS0-NEXT:    v_cmp_eq_u32_e64 s[66:67], 0, v2
 ; GLOBALNESS0-NEXT:    s_branch .LBB1_15
-; GLOBALNESS0-NEXT:  .LBB1_13: ; %Flow7
+; GLOBALNESS0-NEXT:  .LBB1_13: ; %Flow16
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
 ; GLOBALNESS0-NEXT:    s_or_b64 exec, exec, s[4:5]
 ; GLOBALNESS0-NEXT:  .LBB1_14: ; %bb63.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[86:87]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[52:53]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccz .LBB1_25
 ; GLOBALNESS0-NEXT:  .LBB1_15: ; %bb44.i
 ; GLOBALNESS0-NEXT:    ; Parent Loop BB1_4 Depth=1
 ; GLOBALNESS0-NEXT:    ; => This Inner Loop Header: Depth=2
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[94:95]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[48:49]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_14
 ; GLOBALNESS0-NEXT:  ; %bb.16: ; %bb46.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[88:89]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[50:51]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_14
 ; GLOBALNESS0-NEXT:  ; %bb.17: ; %bb50.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[36:37]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[42:43]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_20
 ; GLOBALNESS0-NEXT:  ; %bb.18: ; %bb3.i.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[40:41]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[44:45]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_20
 ; GLOBALNESS0-NEXT:  ; %bb.19: ; %bb6.i.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[56:57]
+; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[64:65]
 ; GLOBALNESS0-NEXT:  .LBB1_20: ; %spam.exit.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[90:91]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[56:57]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_14
 ; GLOBALNESS0-NEXT:  ; %bb.21: ; %bb55.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
-; GLOBALNESS0-NEXT:    s_add_u32 s64, s60, 40
-; GLOBALNESS0-NEXT:    s_addc_u32 s65, s61, 0
-; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[62:63]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[54:55]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[8:9], s[64:65]
+; GLOBALNESS0-NEXT:    s_add_u32 s72, s38, 40
+; GLOBALNESS0-NEXT:    s_addc_u32 s73, s39, 0
+; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[36:37]
+; GLOBALNESS0-NEXT:    s_mov_b64 s[8:9], s[72:73]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS0-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS0-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS0-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v43
-; GLOBALNESS0-NEXT:    s_swappc_b64 s[30:31], s[66:67]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[44:45], 0, 0
-; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[62:63]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[54:55]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[8:9], s[64:65]
+; GLOBALNESS0-NEXT:    s_mov_b32 s12, s70
+; GLOBALNESS0-NEXT:    s_mov_b32 s13, s69
+; GLOBALNESS0-NEXT:    s_mov_b32 s14, s68
+; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v42
+; GLOBALNESS0-NEXT:    s_swappc_b64 s[30:31], s[78:79]
+; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[46:47], 0, 0
+; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[36:37]
+; GLOBALNESS0-NEXT:    s_mov_b64 s[8:9], s[72:73]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS0-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS0-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS0-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v43
-; GLOBALNESS0-NEXT:    global_store_dwordx2 v[44:45], a[32:33], off
-; GLOBALNESS0-NEXT:    s_swappc_b64 s[30:31], s[66:67]
-; GLOBALNESS0-NEXT:    s_and_saveexec_b64 s[4:5], s[58:59]
+; GLOBALNESS0-NEXT:    s_mov_b32 s12, s70
+; GLOBALNESS0-NEXT:    s_mov_b32 s13, s69
+; GLOBALNESS0-NEXT:    s_mov_b32 s14, s68
+; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v42
+; GLOBALNESS0-NEXT:    global_store_dwordx2 v[46:47], v[44:45], off
+; GLOBALNESS0-NEXT:    s_swappc_b64 s[30:31], s[78:79]
+; GLOBALNESS0-NEXT:    s_and_saveexec_b64 s[4:5], s[66:67]
 ; GLOBALNESS0-NEXT:    s_cbranch_execz .LBB1_13
 ; GLOBALNESS0-NEXT:  ; %bb.22: ; %bb62.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_15 Depth=2
 ; GLOBALNESS0-NEXT:    v_mov_b32_e32 v41, v40
-; GLOBALNESS0-NEXT:    global_store_dwordx2 v[44:45], v[40:41], off
+; GLOBALNESS0-NEXT:    global_store_dwordx2 v[46:47], v[40:41], off
 ; GLOBALNESS0-NEXT:    s_branch .LBB1_13
 ; GLOBALNESS0-NEXT:  .LBB1_23: ; %LeafBlock
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    s_cmp_lg_u32 s39, 0
+; GLOBALNESS0-NEXT:    s_cmp_lg_u32 s75, 0
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], 0
 ; GLOBALNESS0-NEXT:    s_cselect_b64 s[6:7], -1, 0
 ; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[6:7]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_9
 ; GLOBALNESS0-NEXT:  .LBB1_24: ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], -1
-; GLOBALNESS0-NEXT:    ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31
+; GLOBALNESS0-NEXT:    ; implicit-def: $vgpr0_vgpr1
 ; GLOBALNESS0-NEXT:    s_branch .LBB1_3
-; GLOBALNESS0-NEXT:  .LBB1_25: ; %Flow14
+; GLOBALNESS0-NEXT:  .LBB1_25: ; %Flow23
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[36:37]
-; GLOBALNESS0-NEXT:    s_mov_b32 s36, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s37, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s38, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s39, s93
-; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[40:41]
-; GLOBALNESS0-NEXT:    s_mov_b32 s40, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s41, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s42, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s43, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s44, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s45, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s46, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s47, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s48, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s49, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s50, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s51, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s52, s93
-; GLOBALNESS0-NEXT:    s_mov_b32 s53, s93
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[0:1], s[36:37], s[36:37] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[2:3], s[38:39], s[38:39] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[4:5], s[40:41], s[40:41] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[6:7], s[42:43], s[42:43] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[8:9], s[44:45], s[44:45] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[10:11], s[46:47], s[46:47] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[12:13], s[48:49], s[48:49] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[14:15], s[50:51], s[50:51] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[16:17], s[52:53], s[52:53] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[18:19], s[54:55], s[54:55] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[20:21], s[56:57], s[56:57] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[22:23], s[58:59], s[58:59] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[24:25], s[60:61], s[60:61] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[26:27], s[62:63], s[62:63] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[28:29], s[64:65], s[64:65] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[30:31], s[66:67], s[66:67] op_sel:[0,1]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[40:41], s[6:7]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[36:37], s[4:5]
-; GLOBALNESS0-NEXT:    s_mov_b32 s39, s75
-; GLOBALNESS0-NEXT:    s_mov_b64 s[42:43], s[72:73]
-; GLOBALNESS0-NEXT:  .LBB1_26: ; %Flow15
+; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[0:1], 0, 0
+; GLOBALNESS0-NEXT:  .LBB1_26: ; %Flow24
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    s_or_b64 exec, exec, s[70:71]
-; GLOBALNESS0-NEXT:    s_and_saveexec_b64 s[4:5], s[96:97]
+; GLOBALNESS0-NEXT:    s_or_b64 exec, exec, s[80:81]
+; GLOBALNESS0-NEXT:    s_and_saveexec_b64 s[4:5], s[62:63]
 ; GLOBALNESS0-NEXT:    s_cbranch_execz .LBB1_2
 ; GLOBALNESS0-NEXT:  ; %bb.27: ; %bb67.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
-; GLOBALNESS0-NEXT:    v_readlane_b32 s6, v42, 2
-; GLOBALNESS0-NEXT:    v_readlane_b32 s7, v42, 3
-; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[6:7]
+; GLOBALNESS0-NEXT:    s_and_b64 vcc, exec, s[58:59]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_1
 ; GLOBALNESS0-NEXT:  ; %bb.28: ; %bb69.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS0-NEXT:    v_mov_b32_e32 v41, v40
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[32:33], 0, 0
-; GLOBALNESS0-NEXT:    global_store_dwordx2 v[32:33], v[40:41], off
+; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[2:3], 0, 0
+; GLOBALNESS0-NEXT:    global_store_dwordx2 v[2:3], v[40:41], off
 ; GLOBALNESS0-NEXT:    s_branch .LBB1_1
 ; GLOBALNESS0-NEXT:  .LBB1_29: ; %bb73.i
 ; GLOBALNESS0-NEXT:    ; in Loop: Header=BB1_4 Depth=1
 ; GLOBALNESS0-NEXT:    v_mov_b32_e32 v41, v40
-; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[32:33], 0, 0
-; GLOBALNESS0-NEXT:    global_store_dwordx2 v[32:33], v[40:41], off
+; GLOBALNESS0-NEXT:    v_pk_mov_b32 v[2:3], 0, 0
+; GLOBALNESS0-NEXT:    global_store_dwordx2 v[2:3], v[40:41], off
 ; GLOBALNESS0-NEXT:    s_branch .LBB1_2
 ; GLOBALNESS0-NEXT:  .LBB1_30: ; %loop.exit.guard
 ; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[4:5]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], -1
 ; GLOBALNESS0-NEXT:    s_cbranch_vccz .LBB1_32
 ; GLOBALNESS0-NEXT:  ; %bb.31: ; %bb7.i.i
-; GLOBALNESS0-NEXT:    s_add_u32 s8, s60, 40
-; GLOBALNESS0-NEXT:    s_addc_u32 s9, s61, 0
-; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[62:63]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[54:55]
+; GLOBALNESS0-NEXT:    s_add_u32 s8, s38, 40
+; GLOBALNESS0-NEXT:    s_addc_u32 s9, s39, 0
+; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[36:37]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS0-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS0-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS0-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v43
+; GLOBALNESS0-NEXT:    s_mov_b32 s12, s70
+; GLOBALNESS0-NEXT:    s_mov_b32 s13, s69
+; GLOBALNESS0-NEXT:    s_mov_b32 s14, s68
+; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v42
 ; GLOBALNESS0-NEXT:    s_getpc_b64 s[16:17]
 ; GLOBALNESS0-NEXT:    s_add_u32 s16, s16, widget@rel32@lo+4
 ; GLOBALNESS0-NEXT:    s_addc_u32 s17, s17, widget@rel32@hi+12
 ; GLOBALNESS0-NEXT:    s_swappc_b64 s[30:31], s[16:17]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], 0
 ; GLOBALNESS0-NEXT:  .LBB1_32: ; %Flow
 ; GLOBALNESS0-NEXT:    s_andn2_b64 vcc, exec, s[4:5]
 ; GLOBALNESS0-NEXT:    s_cbranch_vccnz .LBB1_34
 ; GLOBALNESS0-NEXT:  ; %bb.33: ; %bb11.i.i
-; GLOBALNESS0-NEXT:    s_add_u32 s8, s60, 40
-; GLOBALNESS0-NEXT:    s_addc_u32 s9, s61, 0
-; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[62:63]
-; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[54:55]
+; GLOBALNESS0-NEXT:    s_add_u32 s8, s38, 40
+; GLOBALNESS0-NEXT:    s_addc_u32 s9, s39, 0
+; GLOBALNESS0-NEXT:    s_mov_b64 s[4:5], s[40:41]
+; GLOBALNESS0-NEXT:    s_mov_b64 s[6:7], s[36:37]
 ; GLOBALNESS0-NEXT:    s_mov_b64 s[10:11], s[34:35]
-; GLOBALNESS0-NEXT:    s_mov_b32 s12, s100
-; GLOBALNESS0-NEXT:    s_mov_b32 s13, s99
-; GLOBALNESS0-NEXT:    s_mov_b32 s14, s98
-; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v43
+; GLOBALNESS0-NEXT:    s_mov_b32 s12, s70
+; GLOBALNESS0-NEXT:    s_mov_b32 s13, s69
+; GLOBALNESS0-NEXT:    s_mov_b32 s14, s68
+; GLOBALNESS0-NEXT:    v_mov_b32_e32 v31, v42
 ; GLOBALNESS0-NEXT:    s_getpc_b64 s[16:17]
 ; GLOBALNESS0-NEXT:    s_add_u32 s16, s16, widget@rel32@lo+4
 ; GLOBALNESS0-NEXT:    s_addc_u32 s17, s17, widget@rel32@hi+12
 ; GLOBALNESS0-NEXT:    s_swappc_b64 s[30:31], s[16:17]
 ; GLOBALNESS0-NEXT:  .LBB1_34: ; %UnifiedUnreachableBlock
 bb:
   store i32 0, ptr addrspace(1) null, align 4
   %tmp4 = load i32, ptr addrspace(1) %arg1.global, align 4

llvm/test/CodeGen/AMDGPU/v1024.ll

 ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 ; Check that we do not use AGPRs for v32i32 type
 
 ; GCN-LABEL: {{^}}test_v1024:
 ; GCN-NOT: v_accvgpr
-; GCN-COUNT-32: v_mov_b32_e32
+; GCN-COUNT-8: global_store_dwordx4
 ; GCN-NOT: v_accvgpr
 define amdgpu_kernel void @test_v1024() {
 entry:
   %alloca = alloca <32 x i32>, align 16, addrspace(5)
   call void @llvm.memset.p5.i32(ptr addrspace(5) %alloca, i8 0, i32 128, i1 false)
   br i1 undef, label %if.then.i.i, label %if.else.i
 
 if.then.i.i:                                      ; preds = %entry