Inlining: Run the legacy AlwaysInliner before the regular inliner.

[aemerson]

Summary

We have several situations where it's beneficial for code size to ensure that every
call to always-inline functions are inlined before normal inlining decisions are
made. While the normal inliner runs in a "MandatoryOnly" mode to try to do this,
it only does it on a per-SCC basis, rather than the whole module. Ensuring that
all mandatory inlinings are done before any heuristic based decisions are made
just makes sense.

Despite being referred to the "legacy" AlwaysInliner pass, it's already necessary
for -O0 because the CGSCC inliner is too expensive in compile time to run at -O0.

This also fixes an exponential compile time blow up in
https://github.com/llvm/llvm-project/issues/59126

[aeubanks]

__clang_hip_math.hip is annoying...

We'll need to remove the MandatoryFirst inliner in ModuleInlinerWrapperPass, although not sure if @mtrofin has any issues with that or not

This isn't quite what I had initially thought, but this might be better. (I was thinking that we sort the calls in the inliner to visit alwaysinline calls first, but that might cause more compile time issues since we have to update the call graph after visiting all the calls in a function, but we might be visiting every function twice if we first batch process the alwaysinline calls then all other calls)

llvm/lib/Passes/PassBuilderPipelines.cpp

1082

I think we want to insert lifetime intrinsics when optimizing

[aemerson]

In D143624#4115986, @aeubanks wrote:

__clang_hip_math.hip is annoying...

We'll need to remove the MandatoryFirst inliner in ModuleInlinerWrapperPass, although not sure if @mtrofin has any issues with that or not

This isn't quite what I had initially thought, but this might be better. (I was thinking that we sort the calls in the inliner to visit alwaysinline calls first, but that might cause more compile time issues since we have to update the call graph after visiting all the calls in a function, but we might be visiting every function twice if we first batch process the alwaysinline calls then all other calls)

I think that doesn't actually do the same thing as this, since the Calls vector is populated by visiting the functions in the current SCC. What we're trying to do with this patch is to ensure that all always-inline calls globally are processed first.

[aeubanks]

In D143624#4116080, @aemerson wrote:

In D143624#4115986, @aeubanks wrote:

__clang_hip_math.hip is annoying...

We'll need to remove the MandatoryFirst inliner in ModuleInlinerWrapperPass, although not sure if @mtrofin has any issues with that or not

This isn't quite what I had initially thought, but this might be better. (I was thinking that we sort the calls in the inliner to visit alwaysinline calls first, but that might cause more compile time issues since we have to update the call graph after visiting all the calls in a function, but we might be visiting every function twice if we first batch process the alwaysinline calls then all other calls)

I think that doesn't actually do the same thing as this, since the Calls vector is populated by visiting the functions in the current SCC. What we're trying to do with this patch is to ensure that all always-inline calls globally are processed first.

That's true, but the legacy pass manager where the inliner explosion didn't happen in your case didn't process always-inline calls before other calls. So I don't think it's necessary to process alwaysinline calls globally first to fix your case. However, given that we still do two more rounds of inlining in the inliner pipeline after the alwaysinliner pass you added and your case still doesn't blow up, this solution does seem robust.

[mtrofin]

In D143624#4115986, @aeubanks wrote:

__clang_hip_math.hip is annoying...

We'll need to remove the MandatoryFirst inliner in ModuleInlinerWrapperPass, although not sure if @mtrofin has any issues with that or not

IIRC we had at a point a mandatory , whole-module pass. The idea wast that let's not have N inliners, and that AlwaysInliner had some limitations, and for similar reasons @aemerson pointed out, it'd make sense to first perform the mandatory inlines (D91567). In D94825 we went away from that. I don't remember why. @aeubanks? (it's referencing a patch you started)

This isn't quite what I had initially thought, but this might be better. (I was thinking that we sort the calls in the inliner to visit alwaysinline calls first, but that might cause more compile time issues since we have to update the call graph after visiting all the calls in a function, but we might be visiting every function twice if we first batch process the alwaysinline calls then all other calls)

[aemerson]

In D143624#4116114, @aeubanks wrote:

In D143624#4116080, @aemerson wrote:

In D143624#4115986, @aeubanks wrote:

__clang_hip_math.hip is annoying...

We'll need to remove the MandatoryFirst inliner in ModuleInlinerWrapperPass, although not sure if @mtrofin has any issues with that or not

This isn't quite what I had initially thought, but this might be better. (I was thinking that we sort the calls in the inliner to visit alwaysinline calls first, but that might cause more compile time issues since we have to update the call graph after visiting all the calls in a function, but we might be visiting every function twice if we first batch process the alwaysinline calls then all other calls)

I think that doesn't actually do the same thing as this, since the Calls vector is populated by visiting the functions in the current SCC. What we're trying to do with this patch is to ensure that all always-inline calls globally are processed first.

That's true, but the legacy pass manager where the inliner explosion didn't happen in your case didn't process always-inline calls before other calls. So I don't think it's necessary to process alwaysinline calls globally first to fix your case. However, given that we still do two more rounds of inlining in the inliner pipeline after the alwaysinliner pass you added and your case still doesn't blow up, this solution does seem robust.

Sure, the exponential compile time case is actually just a side benefit here. The motivating reason for this change is actually to improve code size when building codebases that make heavy use of always_inline.

[aeubanks]

In D143624#4116171, @mtrofin wrote:

In D143624#4115986, @aeubanks wrote:

__clang_hip_math.hip is annoying...

We'll need to remove the MandatoryFirst inliner in ModuleInlinerWrapperPass, although not sure if @mtrofin has any issues with that or not

IIRC we had at a point a mandatory , whole-module pass. The idea wast that let's not have N inliners, and that AlwaysInliner had some limitations, and for similar reasons @aemerson pointed out, it'd make sense to first perform the mandatory inlines (D91567). In D94825 we went away from that. I don't remember why. @aeubanks? (it's referencing a patch you started)

This isn't quite what I had initially thought, but this might be better. (I was thinking that we sort the calls in the inliner to visit alwaysinline calls first, but that might cause more compile time issues since we have to update the call graph after visiting all the calls in a function, but we might be visiting every function twice if we first batch process the alwaysinline calls then all other calls)

Without knowing that there were inliner explosion issues, I still think it makes more sense to not visit all functions twice. e.g. it helps with cache locality when compiling if you don't visit functions on two separate walks of the call graph.

But if this solves the issue then I think this patch is good. Getting rid of the mandatory inline advisor is a nice cleanup.

[aeubanks]

In D143624#4116318, @aemerson wrote:

Sure, the exponential compile time case is actually just a side benefit here. The motivating reason for this change is actually to improve code size when building codebases that make heavy use of always_inline.

Ah I didn't see that part of the commit message. Mentioning code size more explicitly in the message would be good.

[aeubanks]

Had a chat offline with @mtrofin, wanted to be clear for future purposes that we do need the separate AlwaysInliner pass because it's used in -O0 and constructing a call graph there is non-trivial in terms of compile time. Originally the mandatory mode of the normal inliner was added to maybe remove the separate AlwaysInliner pass in the future, but that's not going to happen because of what I just said. Given that, we can eventually remove the mandatory mode of the normal inliner after this patch goes through. So this patch should also make mandatory-inlining-first false by default, then we remove it in a separate patch.

[aemerson]

In D143624#4116546, @aeubanks wrote:

Had a chat offline with @mtrofin, wanted to be clear for future purposes that we do need the separate AlwaysInliner pass because it's used in -O0 and constructing a call graph there is non-trivial in terms of compile time. Originally the mandatory mode of the normal inliner was added to maybe remove the separate AlwaysInliner pass in the future, but that's not going to happen because of what I just said. Given that, we can eventually remove the mandatory mode of the normal inliner after this patch goes through. So this patch should also make mandatory-inlining-first false by default, then we remove it in a separate patch.

Ok, sounds good. I'll make the changes.

[mtrofin]

lgtm. Like @aeubanks was saying, let's just give a bit of time (1 month or so?) between when this lands and until we clean up the "mandatory" notion from the advisor, just to make sure nothing breaks/regresses.

[mtrofin]

In D143624#4116546, @aeubanks wrote:

Had a chat offline with @mtrofin, wanted to be clear for future purposes that we do need the separate AlwaysInliner pass because it's used in -O0 and constructing a call graph there is non-trivial in terms of compile time.

Worth maybe spelling that out in the patch description - i.e. why not go the D91567 route again, makes it easier to understand later.

Originally the mandatory mode of the normal inliner was added to maybe remove the separate AlwaysInliner pass in the future, but that's not going to happen because of what I just said. Given that, we can eventually remove the mandatory mode of the normal inliner after this patch goes through. So this patch should also make mandatory-inlining-first false by default, then we remove it in a separate patch.

[aemerson]

Address comments. Disable -mandatory-inlining-first by default and insert lifetime intrinsics if not at -O0.

[aeubanks]

llvm/lib/Passes/PassBuilderPipelines.cpp

1082

this will never be called with Level == OptimizationLevel::O0, true is good enough

llvm/test/Transforms/Inline/always-inline-newpm.ll

1

a better file name is always-inline-phase-ordering, legacy PM is deprecated anyway

was this file exploding before?

[aeubanks]

(my latest comments didn't get addressed in the land)

[aemerson]

llvm/lib/Passes/PassBuilderPipelines.cpp

1082

Ok.

llvm/test/Transforms/Inline/always-inline-newpm.ll

1

Ok I'll rename, this one is demonstrating different/smaller code size with this change.

[dmgreen]

Hello - I had to revert this because of some large regressions we got from routines in CMSIS-DSP.

The llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll test shows the problem - that's why that test exists to ensure that any pipeline changes don't negatively affect these routines. Unfortunately you just changed the test as opposed to showing the problems that this causes. They might be fixable with some other tweaks elsewhere, but the ordering of inlining seems important for getting the correct code that can be vectorized nicely.

There are some other cases around inlining this thing on v6m cores: https://github.com/ARM-software/CMSIS-DSP/blob/809202bf185280a322efc2e2c850a544747f9d79/Include/arm_math_memory.h#L76, but I'm not sure about the details yet. The mult examples were the really large regressions.

[aemerson]

In D143624#4118257, @dmgreen wrote:

Hello - I had to revert this because of some large regressions we got from routines in CMSIS-DSP.

The llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll test shows the problem - that's why that test exists to ensure that any pipeline changes don't negatively affect these routines. Unfortunately you just changed the test as opposed to showing the problems that this causes. They might be fixable with some other tweaks elsewhere, but the ordering of inlining seems important for getting the correct code that can be vectorized nicely.

There are some other cases around inlining this thing on v6m cores: https://github.com/ARM-software/CMSIS-DSP/blob/809202bf185280a322efc2e2c850a544747f9d79/Include/arm_math_memory.h#L76, but I'm not sure about the details yet. The mult examples were the really large regressions.

It’s not clear from the original commit message why the test is related to inlining order? It seems entirely testing vectorization cost model which should be insensitive to these kind of changes, right?

[dmgreen]

It’s not clear from the original commit message why the test is related to inlining order? It seems entirely testing vectorization cost model which should be insensitive to these kind of changes, right?

It's a phase ordering test - it's testing the entire pipeline including all the inlining and simplification that needs to happen :)

You can run update_test_checks of the file to see the differences. I believe the inlining causes differences in the code that then cause different vector factors to be chosen. I can try to add a similar test for the other case that got worse, if they are similar.

[aemerson]

In D143624#4118341, @dmgreen wrote:

It’s not clear from the original commit message why the test is related to inlining order? It seems entirely testing vectorization cost model which should be insensitive to these kind of changes, right?

It's a phase ordering test - it's testing the entire pipeline including all the inlining and simplification that needs to happen :)

You can run update_test_checks of the file to see the differences. I believe the inlining causes differences in the code that then cause different vector factors to be chosen. I can try to add a similar test for the other case that got worse, if they are similar.

I’ll take a look, but this indicates to me that there’s something missing from the vectoriser or later passes rather than a problem with the inliners behaviour.

[dmgreen]

I’ll take a look, but this indicates to me that there’s something missing from the vectoriser or later passes rather than a problem with the inliners behaviour.

Sure. I'm not saying that this patch is wrong. I'm just saying that unfortunately it leads to some pretty large regressions. Hopefully we can figure out why and put fixes in place instead of just bodging the tests. Hopefully it's just some missing fold to get the code into the same form it was before, after all the inlining has happened.

I took a look at one of the other cases, it appears to be a pretty unfortunate case of the load order in loops leading to LSR not recognizing chains of instructions due to them being order with offsets [2,0,6,4,10,8,..], as opposed to the order they were in before [0,2,4,6,8,10...], which was an easier to reason about. https://godbolt.org/z/Grv64xoxW. I'm not sure exactly what the best way to fix that would be, without making other cases worse.

[aemerson]

@dmgreen I've been looking at this test again trying to see what's missing. The problem now is that only a VF of 4 is chosen. In the good case, instcombine/simplifyCFG runs so that it simplifies down to an smin intrinsic. After this change __SSAT() is inlined first. We then have:

target datalayout = "e-m:e-i8:8:32-i16:16:32-i64:64-i128:128-n32:64-S128"
target triple = "aarch64-linux-gnu"

define void @arm_mult_q15(ptr %pSrcA, ptr %pSrcB, ptr noalias %pDst, i32 %blockSize) {
entry:
  br label %while.cond

while.cond:                                       ; preds = %while.body, %entry
  %pSrcB.addr.0 = phi ptr [ %pSrcB, %entry ], [ %incdec.ptr1, %while.body ]
  %pDst.addr.0 = phi ptr [ %pDst, %entry ], [ %incdec.ptr4, %while.body ]
  %pSrcA.addr.0 = phi ptr [ %pSrcA, %entry ], [ %incdec.ptr, %while.body ]
  %blkCnt.0 = phi i32 [ %blockSize, %entry ], [ %dec, %while.body ]
  %cmp.not = icmp eq i32 %blkCnt.0, 0
  br i1 %cmp.not, label %while.end, label %while.body

while.body:                                       ; preds = %while.cond
  %incdec.ptr = getelementptr inbounds i16, ptr %pSrcA.addr.0, i32 1
  %0 = load i16, ptr %pSrcA.addr.0, align 2
  %conv = sext i16 %0 to i32
  %incdec.ptr1 = getelementptr inbounds i16, ptr %pSrcB.addr.0, i32 1
  %1 = load i16, ptr %pSrcB.addr.0, align 2
  %conv2 = sext i16 %1 to i32
  %mul = mul nsw i32 %conv, %conv2
  %shr = ashr i32 %mul, 15
  %cmp4.i = icmp sgt i32 %shr, 32767
  %switch.i = icmp ult i1 %cmp4.i, true
  %spec.select.i = select i1 %switch.i, i32 %shr, i32 32767
  %conv3 = trunc i32 %spec.select.i to i16
  %incdec.ptr4 = getelementptr inbounds i16, ptr %pDst.addr.0, i32 1
  store i16 %conv3, ptr %pDst.addr.0, align 2
  %dec = add i32 %blkCnt.0, -1
  br label %while.cond

while.end:                                        ; preds = %while.cond
  ret void
}

These instructions are from the callee that should now be combined into smin:

%cmp4.i = icmp sgt i32 %shr, 32767
%switch.i = icmp ult i1 %cmp4.i, true
%spec.select.i = select i1 %switch.i, i32 %shr, i32 32767

... except due to the surrounding instructions, the first icmp is optimized into
icmp sgt i32 %mul, 1073741823 by InstCombinerImpl::foldICmpInstWithConstant()

This breaks the smin recognition. I'm not sure what the best approach is to fix this. InstCombine already has this chunk of code to try to avoid messing with compares that might form min/max patterns but it expects further simplification to fire:

// Test if the ICmpInst instruction is used exclusively by a select as
// part of a minimum or maximum operation. If so, refrain from doing
// any other folding. This helps out other analyses which understand
// non-obfuscated minimum and maximum idioms, such as ScalarEvolution
// and CodeGen. And in this case, at least one of the comparison
// operands has at least one user besides the compare (the select),
// which would often largely negate the benefit of folding anyway.
//
// Do the same for the other patterns recognized by matchSelectPattern.
if (I.hasOneUse())
  if (SelectInst *SI = dyn_cast<SelectInst>(I.user_back())) {
    Value *A, *B;
    SelectPatternResult SPR = matchSelectPattern(SI, A, B);
    if (SPR.Flavor != SPF_UNKNOWN)
      return nullptr;
  }

Any ideas? I'd really like to get this inliner change in because it's fundamentally a good change to have.

[dmgreen]

Hello. It sounds like it is really close to being OK. The combine of the shift just seem to make things more difficult.

The icmp ult i1 %cmp4.i, true is just a not, would it help if it was actually an xor? Or if the not(icmp sgt) was changed to a slt earlier?

I was taking a look at the example but I am not super sure what to suggest. Would it be best if the code that detect min/max looked through not's?

[aemerson]

In D143624#4312905, @dmgreen wrote:

Hello. It sounds like it is really close to being OK. The combine of the shift just seem to make things more difficult.

The icmp ult i1 %cmp4.i, true is just a not, would it help if it was actually an xor? Or if the not(icmp sgt) was changed to a slt earlier?

I was taking a look at the example but I am not super sure what to suggest. Would it be best if the code that detect min/max looked through not's?

I posted an attempt at this: https://reviews.llvm.org/D149725

[dmgreen]

It looks like there is quite a lot more optimization that happens to the function being always-inlined (__SSAT) before this change. Through multiple rounds of instcombine, almost to the end of the pass pipeline. The new version runs a lot less before inlining, only running instcombine->simplifycfg and not seeing another instcombine to clean up the results. Is that because the AlwaysInlinePass is a module pass and it now only runs the passes up to that point?

It does look like there might be a chance to undo the transform, as opposed to prevent the transform that blocks it. Something like https://alive2.llvm.org/ce/z/qHtPqz seems to happen at at least one point. Might that be more preferable?

There are some other changes I#m seeing though, from the same function inlined into different routine. This one for example seems to be not longer applying to canonicalizeClampLike, so the ssat doesn't get created. https://godbolt.org/z/qMW44qfz4. That doesn't seem to be easily undoable without knowing the value is positive though https://alive2.llvm.org/ce/z/v9YdaK.

[nikic]

In D143624#4315468, @dmgreen wrote:

It looks like there is quite a lot more optimization that happens to the function being always-inlined (__SSAT) before this change. Through multiple rounds of instcombine, almost to the end of the pass pipeline. The new version runs a lot less before inlining, only running instcombine->simplifycfg and not seeing another instcombine to clean up the results. Is that because the AlwaysInlinePass is a module pass and it now only runs the passes up to that point?

Yes, which is why I personally think this change isn't a good idea. This essentially breaks our invariant that functions get simplified before they are inlined. This significantly alters the way alwaysinline functions will be optimized relative to normally inlined functions.

[mtrofin]

In D143624#4315508, @nikic wrote:

In D143624#4315468, @dmgreen wrote:

It looks like there is quite a lot more optimization that happens to the function being always-inlined (__SSAT) before this change. Through multiple rounds of instcombine, almost to the end of the pass pipeline. The new version runs a lot less before inlining, only running instcombine->simplifycfg and not seeing another instcombine to clean up the results. Is that because the AlwaysInlinePass is a module pass and it now only runs the passes up to that point?

Yes, which is why I personally think this change isn't a good idea. This essentially breaks our invariant that functions get simplified before they are inlined. This significantly alters the way alwaysinline functions will be optimized relative to normally inlined functions.

(Nitpicking just on the invariant part) Not sure if that's always the invariant, because we could be inlining a call site in a SCC where both caller and callee are in that same SCC.

[aeubanks]

In D143624#4315508, @nikic wrote:

In D143624#4315468, @dmgreen wrote:

It looks like there is quite a lot more optimization that happens to the function being always-inlined (__SSAT) before this change. Through multiple rounds of instcombine, almost to the end of the pass pipeline. The new version runs a lot less before inlining, only running instcombine->simplifycfg and not seeing another instcombine to clean up the results. Is that because the AlwaysInlinePass is a module pass and it now only runs the passes up to that point?

Yes, which is why I personally think this change isn't a good idea. This essentially breaks our invariant that functions get simplified before they are inlined. This significantly alters the way alwaysinline functions will be optimized relative to normally inlined functions.

That invariant shouldn't matter if we're not using heuristics to inline. The normal heuristic-based inliner will still work on simplified callees, but now with the additional benefit of seeing the state of an SCC where there may be alwaysinline calls after the inlinings that must happen having happened.

[nikic]

In D143624#4316357, @aeubanks wrote:

In D143624#4315508, @nikic wrote:

In D143624#4315468, @dmgreen wrote:

It looks like there is quite a lot more optimization that happens to the function being always-inlined (__SSAT) before this change. Through multiple rounds of instcombine, almost to the end of the pass pipeline. The new version runs a lot less before inlining, only running instcombine->simplifycfg and not seeing another instcombine to clean up the results. Is that because the AlwaysInlinePass is a module pass and it now only runs the passes up to that point?

Yes, which is why I personally think this change isn't a good idea. This essentially breaks our invariant that functions get simplified before they are inlined. This significantly alters the way alwaysinline functions will be optimized relative to normally inlined functions.

That invariant shouldn't matter if we're not using heuristics to inline. The normal heuristic-based inliner will still work on simplified callees, but now with the additional benefit of seeing the state of an SCC where there may be alwaysinline calls after the inlinings that must happen having happened.

I agree in the sense that for alwaysinling the simplification doesn't matter for cost modelling purposes. However, it can still have other benefits. One is that we don't repeat unnecessary work. Any simplification we do before inlining doesn't have to be repeated for every (potentially transitive) caller it gets inlined into. The other (and in a way, opposite) is that we simplify the function before inlining and then again the callee after inlining, which may paper over phase ordering issues (or the problem discussed above, which is kind of in the same category). Of course, this is not what this is intended for and we should endeavor to fix such issues -- but the practical outcome is still that you'll probably get worse optimization if you use alwaysinline vs inline after this change, which seems kind of unintuitive.

[aemerson]

I have another attempt at fixing this in D149918

clang/test/CodeGen/code-coverage.c

 /// We support coverage versions 3.4, 4.7 and 4.8.
 /// 3.4 redesigns the format and changed .da to .gcda
 /// 4.7 enables cfg_checksum.
 /// 4.8 (default, compatible with gcov 7) emits the exit block the second.
 // RUN: %clang_cc1 -emit-llvm -disable-red-zone -fprofile-arcs -coverage-version='304*' %s -o - | \
 // RUN:   FileCheck --check-prefixes=CHECK,304 %s
 // RUN: %clang_cc1 -emit-llvm -disable-red-zone -fprofile-arcs -coverage-version='407*' %s -o - | \
 // RUN:   FileCheck --check-prefixes=CHECK,407 %s
 // RUN: %clang_cc1 -emit-llvm -disable-red-zone -fprofile-arcs %s -o - | \
 // RUN:   FileCheck --check-prefixes=CHECK,408 %s
 
 // RUN: %clang_cc1 -emit-llvm -disable-red-zone -fprofile-arcs -coverage-notes-file=aaa.gcno -coverage-data-file=bbb.gcda -debug-info-kind=limited -dwarf-version=4 %s -o - | FileCheck %s --check-prefix GCOV_FILE_INFO
 
 // RUN: %clang_cc1 -emit-llvm-bc -o /dev/null -fdebug-pass-manager -fprofile-arcs %s 2>&1 | FileCheck --check-prefix=NEWPM %s
 // RUN: %clang_cc1 -emit-llvm-bc -o /dev/null -fdebug-pass-manager -fprofile-arcs -O3 %s 2>&1 | FileCheck --check-prefix=NEWPM-O3 %s
 
-// NEWPM-NOT: Running pass
 // NEWPM: Running pass: GCOVProfilerPass
 
 // NEWPM-O3-NOT: Running pass
 // NEWPM-O3: Running pass: Annotation2MetadataPass
 // NEWPM-O3: Running pass: ForceFunctionAttrsPass
 // NEWPM-O3: Running pass: GCOVProfilerPass
 
 int test1(int a) {

clang/test/Frontend/optimization-remark-with-hotness-new-pm.c

 
 void bar(int x) {
   // HOTNESS_OFF: 'foo' inlined into 'bar'
   // HOTNESS_OFF-NOT: hotness:
   // THRESHOLD-NOT: inlined
   // THRESHOLD-NOT: hotness
   // NO_PGO: '-fdiagnostics-show-hotness' requires profile-guided optimization information
   // NO_PGO: '-fdiagnostics-hotness-threshold=' requires profile-guided optimization information
-  // expected-remark@+1 {{'foo' inlined into 'bar': always inline attribute at callsite bar:8:10; (hotness:}}
+  // expected-remark@+1 {{'foo' inlined into 'bar' with (cost=always): always inline attribute at callsite bar:8:10; (hotness:}}
   sum += foo(x, x - 2);
 }
 
 int main(int argc, const char *argv[]) {
   for (int i = 0; i < 30; i++)
     // expected-remark@+1 {{'bar' inlined into 'main' with}}
     bar(argc);
   return sum;
 }

clang/test/Headers/__clang_hip_math.hip

 extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
   return __make_mantissa_base16(p);
 }
 
 // CHECK-LABEL: @test___make_mantissa(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = load i8, ptr [[P:%.*]], align 1, !tbaa [[TBAA3]]
 // CHECK-NEXT:    [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT:    br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I33_I:%.*]]
+// CHECK-NEXT:    br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I17_I:%.*]]
 // CHECK:       if.then.i:
 // CHECK-NEXT:    [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[P]], i64 1
 // CHECK-NEXT:    [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    switch i8 [[TMP1]], label [[WHILE_COND_I17_I:%.*]] [
-// CHECK-NEXT:    i8 120, label [[WHILE_COND_I_I_PREHEADER:%.*]]
-// CHECK-NEXT:    i8 88, label [[WHILE_COND_I_I_PREHEADER]]
+// CHECK-NEXT:    switch i8 [[TMP1]], label [[WHILE_COND_I_I:%.*]] [
+// CHECK-NEXT:    i8 120, label [[WHILE_COND_I33_I_PREHEADER:%.*]]
+// CHECK-NEXT:    i8 88, label [[WHILE_COND_I33_I_PREHEADER]]
 // CHECK-NEXT:    ]
-// CHECK:       while.cond.i.i.preheader:
-// CHECK-NEXT:    br label [[WHILE_COND_I_I:%.*]]
-// CHECK:       while.cond.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I_I_PREHEADER]] ]
-// CHECK-NEXT:    [[__R_0_I_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[WHILE_COND_I_I_PREHEADER]] ]
-// CHECK-NEXT:    [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT:    br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I_I:%.*]]
-// CHECK:       while.body.i.i:
+// CHECK:       while.cond.i33.i.preheader:
+// CHECK-NEXT:    br label [[WHILE_COND_I33_I:%.*]]
+// CHECK:       while.cond.i33.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I30_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I:%.*]], [[CLEANUP_I44_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I33_I_PREHEADER]] ]
+// CHECK-NEXT:    [[__R_0_I31_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I44_I]] ], [ 0, [[WHILE_COND_I33_I_PREHEADER]] ]
+// CHECK-NEXT:    [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I32_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[CMP_NOT_I32_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I:%.*]]
+// CHECK:       while.body.i35.i:
 // CHECK-NEXT:    [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT:    [[OR_COND_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT:    br i1 [[OR_COND_I_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]]
+// CHECK-NEXT:    [[OR_COND_I34_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// CHECK-NEXT:    br i1 [[OR_COND_I34_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]]
 // CHECK:       if.else.i.i:
 // CHECK-NEXT:    [[TMP4:%.*]] = add i8 [[TMP2]], -97
 // CHECK-NEXT:    [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
 // CHECK-NEXT:    br i1 [[OR_COND33_I_I]], label [[IF_END31_I_I]], label [[IF_ELSE17_I_I:%.*]]
 // CHECK:       if.else17.i.i:
 // CHECK-NEXT:    [[TMP5:%.*]] = add i8 [[TMP2]], -65
 // CHECK-NEXT:    [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT:    br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I_I]]
+// CHECK-NEXT:    br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I44_I]]
 // CHECK:       if.end31.i.i:
-// CHECK-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT:    [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I_I]], 4
+// CHECK-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT:    [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I31_I]], 4
 // CHECK-NEXT:    [[CONV25_I_I:%.*]] = sext i8 [[TMP2]] to i64
 // CHECK-NEXT:    [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]]
 // CHECK-NEXT:    [[ADD28_I_I:%.*]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]]
+// CHECK-NEXT:    [[INCDEC_PTR_I42_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I]], i64 1
+// CHECK-NEXT:    br label [[CLEANUP_I44_I]]
+// CHECK:       cleanup.i44.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I43_I]] = phi ptr [ [[INCDEC_PTR_I42_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I30_I]], [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT:    [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I31_I]], [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT:    [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT:    br i1 [[COND_I_I]], label [[WHILE_COND_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK:       while.cond.i.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ]
+// CHECK-NEXT:    [[__R_0_I_I:%.*]] = phi i64 [ [[__R_1_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[IF_THEN_I]] ]
+// CHECK-NEXT:    [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT:    br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I_I:%.*]]
+// CHECK:       while.body.i.i:
+// CHECK-NEXT:    [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT:    [[OR_COND_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT:    br i1 [[OR_COND_I_I]], label [[IF_THEN_I_I:%.*]], label [[CLEANUP_I_I]]
+// CHECK:       if.then.i.i:
+// CHECK-NEXT:    [[MUL_I_I:%.*]] = shl i64 [[__R_0_I_I]], 3
+// CHECK-NEXT:    [[CONV5_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT:    [[ADD_I_I:%.*]] = add i64 [[MUL_I_I]], -48
+// CHECK-NEXT:    [[SUB_I_I:%.*]] = add i64 [[ADD_I_I]], [[CONV5_I_I]]
 // CHECK-NEXT:    [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I]], i64 1
 // CHECK-NEXT:    br label [[CLEANUP_I_I]]
 // CHECK:       cleanup.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT:    [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I_I]], [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT:    [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT:    br i1 [[COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[WHILE_BODY_I_I]] ]
+// CHECK-NEXT:    [[__R_1_I_I]] = phi i64 [ [[SUB_I_I]], [[IF_THEN_I_I]] ], [ [[__R_0_I_I]], [[WHILE_BODY_I_I]] ]
+// CHECK-NEXT:    br i1 [[OR_COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP6]]
 // CHECK:       while.cond.i17.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I14_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I:%.*]], [[CLEANUP_I28_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ]
-// CHECK-NEXT:    [[__R_0_I15_I:%.*]] = phi i64 [ [[__R_1_I27_I:%.*]], [[CLEANUP_I28_I]] ], [ 0, [[IF_THEN_I]] ]
-// CHECK-NEXT:    [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I16_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I14_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I:%.*]], [[CLEANUP_I28_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
+// CHECK-NEXT:    [[__R_0_I15_I:%.*]] = phi i64 [ [[__R_1_I27_I:%.*]], [[CLEANUP_I28_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT:    [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I16_I:%.*]] = icmp eq i8 [[TMP8]], 0
 // CHECK-NEXT:    br i1 [[CMP_NOT_I16_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I19_I:%.*]]
 // CHECK:       while.body.i19.i:
-// CHECK-NEXT:    [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT:    [[OR_COND_I18_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT:    [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT:    [[OR_COND_I18_I:%.*]] = icmp ult i8 [[TMP9]], 10
 // CHECK-NEXT:    br i1 [[OR_COND_I18_I]], label [[IF_THEN_I25_I:%.*]], label [[CLEANUP_I28_I]]
 // CHECK:       if.then.i25.i:
-// CHECK-NEXT:    [[MUL_I20_I:%.*]] = shl i64 [[__R_0_I15_I]], 3
-// CHECK-NEXT:    [[CONV5_I21_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT:    [[MUL_I20_I:%.*]] = mul i64 [[__R_0_I15_I]], 10
+// CHECK-NEXT:    [[CONV5_I21_I:%.*]] = sext i8 [[TMP8]] to i64
 // CHECK-NEXT:    [[ADD_I22_I:%.*]] = add i64 [[MUL_I20_I]], -48
 // CHECK-NEXT:    [[SUB_I23_I:%.*]] = add i64 [[ADD_I22_I]], [[CONV5_I21_I]]
 // CHECK-NEXT:    [[INCDEC_PTR_I24_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I]], i64 1
 // CHECK-NEXT:    br label [[CLEANUP_I28_I]]
 // CHECK:       cleanup.i28.i:
 // CHECK-NEXT:    [[__TAGP_ADDR_1_I26_I]] = phi ptr [ [[INCDEC_PTR_I24_I]], [[IF_THEN_I25_I]] ], [ [[__TAGP_ADDR_0_I14_I]], [[WHILE_BODY_I19_I]] ]
 // CHECK-NEXT:    [[__R_1_I27_I]] = phi i64 [ [[SUB_I23_I]], [[IF_THEN_I25_I]] ], [ [[__R_0_I15_I]], [[WHILE_BODY_I19_I]] ]
-// CHECK-NEXT:    br i1 [[OR_COND_I18_I]], label [[WHILE_COND_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP6]]
-// CHECK:       while.cond.i33.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I30_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I:%.*]], [[CLEANUP_I44_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
-// CHECK-NEXT:    [[__R_0_I31_I:%.*]] = phi i64 [ [[__R_1_I43_I:%.*]], [[CLEANUP_I44_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT:    [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I32_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT:    br i1 [[CMP_NOT_I32_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I35_I:%.*]]
-// CHECK:       while.body.i35.i:
-// CHECK-NEXT:    [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT:    [[OR_COND_I34_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT:    br i1 [[OR_COND_I34_I]], label [[IF_THEN_I41_I:%.*]], label [[CLEANUP_I44_I]]
-// CHECK:       if.then.i41.i:
-// CHECK-NEXT:    [[MUL_I36_I:%.*]] = mul i64 [[__R_0_I31_I]], 10
-// CHECK-NEXT:    [[CONV5_I37_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT:    [[ADD_I38_I:%.*]] = add i64 [[MUL_I36_I]], -48
-// CHECK-NEXT:    [[SUB_I39_I:%.*]] = add i64 [[ADD_I38_I]], [[CONV5_I37_I]]
-// CHECK-NEXT:    [[INCDEC_PTR_I40_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I]], i64 1
-// CHECK-NEXT:    br label [[CLEANUP_I44_I]]
-// CHECK:       cleanup.i44.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_1_I42_I]] = phi ptr [ [[INCDEC_PTR_I40_I]], [[IF_THEN_I41_I]] ], [ [[__TAGP_ADDR_0_I30_I]], [[WHILE_BODY_I35_I]] ]
-// CHECK-NEXT:    [[__R_1_I43_I]] = phi i64 [ [[SUB_I39_I]], [[IF_THEN_I41_I]] ], [ [[__R_0_I31_I]], [[WHILE_BODY_I35_I]] ]
-// CHECK-NEXT:    br i1 [[OR_COND_I34_I]], label [[WHILE_COND_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT:    br i1 [[OR_COND_I18_I]], label [[WHILE_COND_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]]
 // CHECK:       _ZL15__make_mantissaPKc.exit:
-// CHECK-NEXT:    [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I]] ], [ [[__R_0_I15_I]], [[WHILE_COND_I17_I]] ], [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I44_I]] ], [ [[__R_0_I31_I]], [[WHILE_COND_I33_I]] ]
+// CHECK-NEXT:    [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I44_I]] ], [ [[__R_0_I31_I]], [[WHILE_COND_I33_I]] ], [ 0, [[CLEANUP_I28_I]] ], [ [[__R_0_I15_I]], [[WHILE_COND_I17_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_0_I]]
 //
 extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
   return __make_mantissa(p);
 }
 
 // CHECK-LABEL: @test_abs(
 // CHECK-NEXT:  entry:

 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
 extern "C" __device__ float test_exp2f(float x) {
   return exp2f(x);
 }
 
-//
 // DEFAULT-LABEL: @test_exp2(
 // DEFAULT-NEXT:  entry:
 // DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]]
 // DEFAULT-NEXT:    ret double [[CALL_I]]
 //
 // FINITEONLY-LABEL: @test_exp2(
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]]

 
 // DEFAULT-LABEL: @test_jnf(
 // DEFAULT-NEXT:  entry:
 // DEFAULT-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
 // DEFAULT-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
 // DEFAULT-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
 // DEFAULT-NEXT:    ]
 // DEFAULT:       if.then.i:
-// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
 // DEFAULT-NEXT:    br label [[_ZL3JNFIF_EXIT:%.*]]
 // DEFAULT:       if.then2.i:
-// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
 // DEFAULT-NEXT:    br label [[_ZL3JNFIF_EXIT]]
 // DEFAULT:       if.end4.i:
-// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT:    br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
 // DEFAULT:       for.body.i:
-// DEFAULT-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
 // DEFAULT-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
 // DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
-// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT:    [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
+// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT:    [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
 // DEFAULT:       _ZL3jnfif.exit:
-// DEFAULT-NEXT:    [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT:    [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
 // DEFAULT-NEXT:    ret float [[RETVAL_0_I]]
 //
 // FINITEONLY-LABEL: @test_jnf(
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
 // FINITEONLY-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
 // FINITEONLY-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
 // FINITEONLY-NEXT:    ]
 // FINITEONLY:       if.then.i:
-// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    br label [[_ZL3JNFIF_EXIT:%.*]]
 // FINITEONLY:       if.then2.i:
-// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    br label [[_ZL3JNFIF_EXIT]]
 // FINITEONLY:       if.end4.i:
-// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT:    br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
 // FINITEONLY:       for.body.i:
-// FINITEONLY-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
+// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
 // FINITEONLY:       _ZL3jnfif.exit:
-// FINITEONLY-NEXT:    [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT:    [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
 // FINITEONLY-NEXT:    ret float [[RETVAL_0_I]]
 //
 extern "C" __device__ float test_jnf(int x, float y) {
   return jnf(x, y);
 }
 
 // DEFAULT-LABEL: @test_jn(
 // DEFAULT-NEXT:  entry:
 // DEFAULT-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
 // DEFAULT-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
 // DEFAULT-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
 // DEFAULT-NEXT:    ]
 // DEFAULT:       if.then.i:
-// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
 // DEFAULT-NEXT:    br label [[_ZL2JNID_EXIT:%.*]]
 // DEFAULT:       if.then2.i:
-// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
 // DEFAULT-NEXT:    br label [[_ZL2JNID_EXIT]]
 // DEFAULT:       if.end4.i:
-// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT:    br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
 // DEFAULT:       for.body.i:
-// DEFAULT-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
 // DEFAULT-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
 // DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
-// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT:    [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
+// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT:    [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
 // DEFAULT:       _ZL2jnid.exit:
-// DEFAULT-NEXT:    [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT:    [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
 // DEFAULT-NEXT:    ret double [[RETVAL_0_I]]
 //
 // FINITEONLY-LABEL: @test_jn(
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
 // FINITEONLY-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
 // FINITEONLY-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
 // FINITEONLY-NEXT:    ]
 // FINITEONLY:       if.then.i:
-// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    br label [[_ZL2JNID_EXIT:%.*]]
 // FINITEONLY:       if.then2.i:
-// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    br label [[_ZL2JNID_EXIT]]
 // FINITEONLY:       if.end4.i:
-// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT:    br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
 // FINITEONLY:       for.body.i:
-// FINITEONLY-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
+// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
 // FINITEONLY:       _ZL2jnid.exit:
-// FINITEONLY-NEXT:    [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT:    [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
 // FINITEONLY-NEXT:    ret double [[RETVAL_0_I]]
 //
 extern "C" __device__ double test_jn(int x, double y) {
   return jn(x, y);
 }
 
 // DEFAULT-LABEL: @test_ldexpf(
 // DEFAULT-NEXT:  entry:

 extern "C" __device__ double test_modf(double x, double* y) {
   return modf(x, y);
 }
 
 // CHECK-LABEL: @test_nanf(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]]
 // CHECK-NEXT:    [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT:    br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I33_I_I:%.*]]
+// CHECK-NEXT:    br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I17_I_I:%.*]]
 // CHECK:       if.then.i.i:
 // CHECK-NEXT:    [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1
 // CHECK-NEXT:    [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    switch i8 [[TMP1]], label [[WHILE_COND_I17_I_I:%.*]] [
-// CHECK-NEXT:    i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]]
-// CHECK-NEXT:    i8 88, label [[WHILE_COND_I_I_I_PREHEADER]]
+// CHECK-NEXT:    switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// CHECK-NEXT:    i8 120, label [[WHILE_COND_I33_I_I_PREHEADER:%.*]]
+// CHECK-NEXT:    i8 88, label [[WHILE_COND_I33_I_I_PREHEADER]]
 // CHECK-NEXT:    ]
-// CHECK:       while.cond.i.i.i.preheader:
-// CHECK-NEXT:    br label [[WHILE_COND_I_I_I:%.*]]
-// CHECK:       while.cond.i.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT:    [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT:    [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT:    br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]]
-// CHECK:       while.body.i.i.i:
+// CHECK:       while.cond.i33.i.i.preheader:
+// CHECK-NEXT:    br label [[WHILE_COND_I33_I_I:%.*]]
+// CHECK:       while.cond.i33.i.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I33_I_I_PREHEADER]] ]
+// CHECK-NEXT:    [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[WHILE_COND_I33_I_I_PREHEADER]] ]
+// CHECK-NEXT:    [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[CMP_NOT_I32_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I_I:%.*]]
+// CHECK:       while.body.i35.i.i:
 // CHECK-NEXT:    [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT:    [[OR_COND_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT:    br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// CHECK-NEXT:    [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// CHECK-NEXT:    br i1 [[OR_COND_I34_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
 // CHECK:       if.else.i.i.i:
 // CHECK-NEXT:    [[TMP4:%.*]] = add i8 [[TMP2]], -97
 // CHECK-NEXT:    [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
 // CHECK-NEXT:    br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label [[IF_ELSE17_I_I_I:%.*]]
 // CHECK:       if.else17.i.i.i:
 // CHECK-NEXT:    [[TMP5:%.*]] = add i8 [[TMP2]], -65
 // CHECK-NEXT:    [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT:    br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]]
+// CHECK-NEXT:    br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I44_I_I]]
 // CHECK:       if.end31.i.i.i:
-// CHECK-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT:    [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4
+// CHECK-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT:    [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I31_I_I]], 4
 // CHECK-NEXT:    [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64
 // CHECK-NEXT:    [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
 // CHECK-NEXT:    [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// CHECK-NEXT:    [[INCDEC_PTR_I42_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
+// CHECK-NEXT:    br label [[CLEANUP_I44_I_I]]
+// CHECK:       cleanup.i44.i.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I43_I_I]] = phi ptr [ [[INCDEC_PTR_I42_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT:    [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I31_I_I]], [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT:    [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT:    br i1 [[COND_I_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK:       while.cond.i.i.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
+// CHECK-NEXT:    [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// CHECK-NEXT:    [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT:    br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I_I_I:%.*]]
+// CHECK:       while.body.i.i.i:
+// CHECK-NEXT:    [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT:    [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT:    br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label [[CLEANUP_I_I_I]]
+// CHECK:       if.then.i.i.i:
+// CHECK-NEXT:    [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// CHECK-NEXT:    [[CONV5_I_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT:    [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// CHECK-NEXT:    [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
 // CHECK-NEXT:    [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
 // CHECK-NEXT:    br label [[CLEANUP_I_I_I]]
 // CHECK:       cleanup.i.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT:    [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT:    [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT:    br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// CHECK-NEXT:    [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// CHECK-NEXT:    br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP6]]
 // CHECK:       while.cond.i17.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
-// CHECK-NEXT:    [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
-// CHECK-NEXT:    [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
+// CHECK-NEXT:    [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT:    [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
 // CHECK-NEXT:    br i1 [[CMP_NOT_I16_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I19_I_I:%.*]]
 // CHECK:       while.body.i19.i.i:
-// CHECK-NEXT:    [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT:    [[OR_COND_I18_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT:    [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT:    [[OR_COND_I18_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
 // CHECK-NEXT:    br i1 [[OR_COND_I18_I_I]], label [[IF_THEN_I25_I_I:%.*]], label [[CLEANUP_I28_I_I]]
 // CHECK:       if.then.i25.i.i:
-// CHECK-NEXT:    [[MUL_I20_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3
-// CHECK-NEXT:    [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT:    [[MUL_I20_I_I:%.*]] = mul i64 [[__R_0_I15_I_I]], 10
+// CHECK-NEXT:    [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP8]] to i64
 // CHECK-NEXT:    [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
 // CHECK-NEXT:    [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
 // CHECK-NEXT:    [[INCDEC_PTR_I24_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1
 // CHECK-NEXT:    br label [[CLEANUP_I28_I_I]]
 // CHECK:       cleanup.i28.i.i:
 // CHECK-NEXT:    [[__TAGP_ADDR_1_I26_I_I]] = phi ptr [ [[INCDEC_PTR_I24_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I19_I_I]] ]
 // CHECK-NEXT:    [[__R_1_I27_I_I]] = phi i64 [ [[SUB_I23_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I19_I_I]] ]
-// CHECK-NEXT:    br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP6]]
-// CHECK:       while.cond.i33.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
-// CHECK-NEXT:    [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT:    [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT:    br i1 [[CMP_NOT_I32_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I35_I_I:%.*]]
-// CHECK:       while.body.i35.i.i:
-// CHECK-NEXT:    [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT:    [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT:    br i1 [[OR_COND_I34_I_I]], label [[IF_THEN_I41_I_I:%.*]], label [[CLEANUP_I44_I_I]]
-// CHECK:       if.then.i41.i.i:
-// CHECK-NEXT:    [[MUL_I36_I_I:%.*]] = mul i64 [[__R_0_I31_I_I]], 10
-// CHECK-NEXT:    [[CONV5_I37_I_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT:    [[ADD_I38_I_I:%.*]] = add i64 [[MUL_I36_I_I]], -48
-// CHECK-NEXT:    [[SUB_I39_I_I:%.*]] = add i64 [[ADD_I38_I_I]], [[CONV5_I37_I_I]]
-// CHECK-NEXT:    [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
-// CHECK-NEXT:    br label [[CLEANUP_I44_I_I]]
-// CHECK:       cleanup.i44.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_1_I42_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[WHILE_BODY_I35_I_I]] ]
-// CHECK-NEXT:    [[__R_1_I43_I_I]] = phi i64 [ [[SUB_I39_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_BODY_I35_I_I]] ]
-// CHECK-NEXT:    br i1 [[OR_COND_I34_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT:    br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]]
 // CHECK:       _ZL4nanfPKc.exit:
-// CHECK-NEXT:    [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ]
+// CHECK-NEXT:    [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ], [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ]
 // CHECK-NEXT:    [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
 // CHECK-NEXT:    [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
 // CHECK-NEXT:    [[BF_SET9_I:%.*]] = or i32 [[BF_VALUE_I]], 2143289344
 // CHECK-NEXT:    [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
 // CHECK-NEXT:    ret float [[TMP10]]
 //
 extern "C" __device__ float test_nanf(const char *tag) {
   return nanf(tag);
 }
 
 // CHECK-LABEL: @test_nan(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]]
 // CHECK-NEXT:    [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT:    br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I33_I_I:%.*]]
+// CHECK-NEXT:    br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I17_I_I:%.*]]
 // CHECK:       if.then.i.i:
 // CHECK-NEXT:    [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1
 // CHECK-NEXT:    [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    switch i8 [[TMP1]], label [[WHILE_COND_I17_I_I:%.*]] [
-// CHECK-NEXT:    i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]]
-// CHECK-NEXT:    i8 88, label [[WHILE_COND_I_I_I_PREHEADER]]
+// CHECK-NEXT:    switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// CHECK-NEXT:    i8 120, label [[WHILE_COND_I33_I_I_PREHEADER:%.*]]
+// CHECK-NEXT:    i8 88, label [[WHILE_COND_I33_I_I_PREHEADER]]
 // CHECK-NEXT:    ]
-// CHECK:       while.cond.i.i.i.preheader:
-// CHECK-NEXT:    br label [[WHILE_COND_I_I_I:%.*]]
-// CHECK:       while.cond.i.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT:    [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT:    [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT:    br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]]
-// CHECK:       while.body.i.i.i:
+// CHECK:       while.cond.i33.i.i.preheader:
+// CHECK-NEXT:    br label [[WHILE_COND_I33_I_I:%.*]]
+// CHECK:       while.cond.i33.i.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I33_I_I_PREHEADER]] ]
+// CHECK-NEXT:    [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[WHILE_COND_I33_I_I_PREHEADER]] ]
+// CHECK-NEXT:    [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[CMP_NOT_I32_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I_I:%.*]]
+// CHECK:       while.body.i35.i.i:
 // CHECK-NEXT:    [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT:    [[OR_COND_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT:    br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// CHECK-NEXT:    [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// CHECK-NEXT:    br i1 [[OR_COND_I34_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
 // CHECK:       if.else.i.i.i:
 // CHECK-NEXT:    [[TMP4:%.*]] = add i8 [[TMP2]], -97
 // CHECK-NEXT:    [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
 // CHECK-NEXT:    br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label [[IF_ELSE17_I_I_I:%.*]]
 // CHECK:       if.else17.i.i.i:
 // CHECK-NEXT:    [[TMP5:%.*]] = add i8 [[TMP2]], -65
 // CHECK-NEXT:    [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT:    br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]]
+// CHECK-NEXT:    br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I44_I_I]]
 // CHECK:       if.end31.i.i.i:
-// CHECK-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT:    [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4
+// CHECK-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT:    [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I31_I_I]], 4
 // CHECK-NEXT:    [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64
 // CHECK-NEXT:    [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
 // CHECK-NEXT:    [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// CHECK-NEXT:    [[INCDEC_PTR_I42_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
+// CHECK-NEXT:    br label [[CLEANUP_I44_I_I]]
+// CHECK:       cleanup.i44.i.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I43_I_I]] = phi ptr [ [[INCDEC_PTR_I42_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT:    [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I31_I_I]], [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT:    [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT:    br i1 [[COND_I_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK:       while.cond.i.i.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
+// CHECK-NEXT:    [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// CHECK-NEXT:    [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT:    br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I_I_I:%.*]]
+// CHECK:       while.body.i.i.i:
+// CHECK-NEXT:    [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT:    [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT:    br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label [[CLEANUP_I_I_I]]
+// CHECK:       if.then.i.i.i:
+// CHECK-NEXT:    [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// CHECK-NEXT:    [[CONV5_I_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT:    [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// CHECK-NEXT:    [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
 // CHECK-NEXT:    [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
 // CHECK-NEXT:    br label [[CLEANUP_I_I_I]]
 // CHECK:       cleanup.i.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT:    [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT:    [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT:    br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// CHECK-NEXT:    [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// CHECK-NEXT:    br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP6]]
 // CHECK:       while.cond.i17.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
-// CHECK-NEXT:    [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
-// CHECK-NEXT:    [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
+// CHECK-NEXT:    [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT:    [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
 // CHECK-NEXT:    br i1 [[CMP_NOT_I16_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I19_I_I:%.*]]
 // CHECK:       while.body.i19.i.i:
-// CHECK-NEXT:    [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT:    [[OR_COND_I18_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT:    [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT:    [[OR_COND_I18_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
 // CHECK-NEXT:    br i1 [[OR_COND_I18_I_I]], label [[IF_THEN_I25_I_I:%.*]], label [[CLEANUP_I28_I_I]]
 // CHECK:       if.then.i25.i.i:
-// CHECK-NEXT:    [[MUL_I20_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3
-// CHECK-NEXT:    [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT:    [[MUL_I20_I_I:%.*]] = mul i64 [[__R_0_I15_I_I]], 10
+// CHECK-NEXT:    [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP8]] to i64
 // CHECK-NEXT:    [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
 // CHECK-NEXT:    [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
 // CHECK-NEXT:    [[INCDEC_PTR_I24_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1
 // CHECK-NEXT:    br label [[CLEANUP_I28_I_I]]
 // CHECK:       cleanup.i28.i.i:
 // CHECK-NEXT:    [[__TAGP_ADDR_1_I26_I_I]] = phi ptr [ [[INCDEC_PTR_I24_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I19_I_I]] ]
 // CHECK-NEXT:    [[__R_1_I27_I_I]] = phi i64 [ [[SUB_I23_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I19_I_I]] ]
-// CHECK-NEXT:    br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP6]]
-// CHECK:       while.cond.i33.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
-// CHECK-NEXT:    [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT:    [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT:    [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT:    br i1 [[CMP_NOT_I32_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I35_I_I:%.*]]
-// CHECK:       while.body.i35.i.i:
-// CHECK-NEXT:    [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT:    [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT:    br i1 [[OR_COND_I34_I_I]], label [[IF_THEN_I41_I_I:%.*]], label [[CLEANUP_I44_I_I]]
-// CHECK:       if.then.i41.i.i:
-// CHECK-NEXT:    [[MUL_I36_I_I:%.*]] = mul i64 [[__R_0_I31_I_I]], 10
-// CHECK-NEXT:    [[CONV5_I37_I_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT:    [[ADD_I38_I_I:%.*]] = add i64 [[MUL_I36_I_I]], -48
-// CHECK-NEXT:    [[SUB_I39_I_I:%.*]] = add i64 [[ADD_I38_I_I]], [[CONV5_I37_I_I]]
-// CHECK-NEXT:    [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
-// CHECK-NEXT:    br label [[CLEANUP_I44_I_I]]
-// CHECK:       cleanup.i44.i.i:
-// CHECK-NEXT:    [[__TAGP_ADDR_1_I42_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[WHILE_BODY_I35_I_I]] ]
-// CHECK-NEXT:    [[__R_1_I43_I_I]] = phi i64 [ [[SUB_I39_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_BODY_I35_I_I]] ]
-// CHECK-NEXT:    br i1 [[OR_COND_I34_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT:    br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]]
 // CHECK:       _ZL3nanPKc.exit:
-// CHECK-NEXT:    [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ]
+// CHECK-NEXT:    [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ], [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ]
 // CHECK-NEXT:    [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
 // CHECK-NEXT:    [[BF_SET9_I:%.*]] = or i64 [[BF_VALUE_I]], 9221120237041090560
 // CHECK-NEXT:    [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
 // CHECK-NEXT:    ret double [[TMP10]]
 //
 extern "C" __device__ double test_nan(const char *tag) {
   return nan(tag);
 }

 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
 extern "C" __device__ double test_normcdfinv(double x) {
   return normcdfinv(x);
 }
 
 // DEFAULT-LABEL: @test_normf(
 // DEFAULT-NEXT:  entry:
-// DEFAULT-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT:    [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
 // DEFAULT:       while.body.i:
-// DEFAULT-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT:    [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA15]]
 // DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT:    [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT:    [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
 // DEFAULT-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
 // DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]]
 // DEFAULT:       _ZL5normfiPKf.exit:
-// DEFAULT-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
+// DEFAULT-NEXT:    [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
 // DEFAULT-NEXT:    ret float [[CALL_I]]
 //
 // FINITEONLY-LABEL: @test_normf(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT:    [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
 // FINITEONLY:       while.body.i:
-// FINITEONLY-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT:    [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA15]]
 // FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
 // FINITEONLY-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
 // FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]]
 // FINITEONLY:       _ZL5normfiPKf.exit:
-// FINITEONLY-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
+// FINITEONLY-NEXT:    [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
 extern "C" __device__ float test_normf(int x, const float *y) {
   return normf(x, y);
 }
 
 // DEFAULT-LABEL: @test_norm(
 // DEFAULT-NEXT:  entry:
-// DEFAULT-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT:    [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
 // DEFAULT:       while.body.i:
-// DEFAULT-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
+// DEFAULT-NEXT:    [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT:    [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA17]]
 // DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT:    [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT:    [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
 // DEFAULT-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
 // DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
 // DEFAULT:       _ZL4normiPKd.exit:
-// DEFAULT-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
+// DEFAULT-NEXT:    [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
 // DEFAULT-NEXT:    ret double [[CALL_I]]
 //
 // FINITEONLY-LABEL: @test_norm(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT:    [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
 // FINITEONLY:       while.body.i:
-// FINITEONLY-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
+// FINITEONLY-NEXT:    [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT:    [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA17]]
 // FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
 // FINITEONLY-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
 // FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
 // FINITEONLY:       _ZL4normiPKd.exit:
-// FINITEONLY-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
+// FINITEONLY-NEXT:    [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
 extern "C" __device__ double test_norm(int x, const double *y) {
   return norm(x, y);
 }
 
 // DEFAULT-LABEL: @test_powf(
 // DEFAULT-NEXT:  entry:

 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
 extern "C" __device__ double test_rint(double x) {
   return rint(x);
 }
 
 // DEFAULT-LABEL: @test_rnormf(
 // DEFAULT-NEXT:  entry:
-// DEFAULT-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT:    [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
 // DEFAULT:       while.body.i:
-// DEFAULT-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT:    [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA15]]
 // DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT:    [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT:    [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
 // DEFAULT-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
 // DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
 // DEFAULT:       _ZL6rnormfiPKf.exit:
-// DEFAULT-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
 // DEFAULT-NEXT:    ret float [[CALL_I]]
 //
 // FINITEONLY-LABEL: @test_rnormf(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT:    [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
 // FINITEONLY:       while.body.i:
-// FINITEONLY-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT:    [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA15]]
 // FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
 // FINITEONLY-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
 // FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
 // FINITEONLY:       _ZL6rnormfiPKf.exit:
-// FINITEONLY-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
 extern "C" __device__ float test_rnormf(int x, const float* y) {
   return rnormf(x, y);
 }
 
 // DEFAULT-LABEL: @test_rnorm(
 // DEFAULT-NEXT:  entry:
-// DEFAULT-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT:    [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
 // DEFAULT:       while.body.i:
-// DEFAULT-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
+// DEFAULT-NEXT:    [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT:    [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA17]]
 // DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT:    [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT:    [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
 // DEFAULT-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
 // DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
 // DEFAULT:       _ZL5rnormiPKd.exit:
-// DEFAULT-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
 // DEFAULT-NEXT:    ret double [[CALL_I]]
 //
 // FINITEONLY-LABEL: @test_rnorm(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT:    [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
 // FINITEONLY:       while.body.i:
-// FINITEONLY-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
+// FINITEONLY-NEXT:    [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT:    [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA17]]
 // FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
 // FINITEONLY-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
 // FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
 // FINITEONLY:       _ZL5rnormiPKd.exit:
-// FINITEONLY-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
 extern "C" __device__ double test_rnorm(int x, const double* y) {
   return rnorm(x, y);
 }
 
 // DEFAULT-LABEL: @test_rnorm3df(
 // DEFAULT-NEXT:  entry:

 
 // DEFAULT-LABEL: @test_ynf(
 // DEFAULT-NEXT:  entry:
 // DEFAULT-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
 // DEFAULT-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
 // DEFAULT-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
 // DEFAULT-NEXT:    ]
 // DEFAULT:       if.then.i:
-// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
 // DEFAULT-NEXT:    br label [[_ZL3YNFIF_EXIT:%.*]]
 // DEFAULT:       if.then2.i:
-// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
 // DEFAULT-NEXT:    br label [[_ZL3YNFIF_EXIT]]
 // DEFAULT:       if.end4.i:
-// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT:    br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
 // DEFAULT:       for.body.i:
-// DEFAULT-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
 // DEFAULT-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
 // DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
-// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT:    [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT:    [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
 // DEFAULT:       _ZL3ynfif.exit:
-// DEFAULT-NEXT:    [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT:    [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
 // DEFAULT-NEXT:    ret float [[RETVAL_0_I]]
 //
 // FINITEONLY-LABEL: @test_ynf(
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
 // FINITEONLY-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
 // FINITEONLY-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
 // FINITEONLY-NEXT:    ]
 // FINITEONLY:       if.then.i:
-// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    br label [[_ZL3YNFIF_EXIT:%.*]]
 // FINITEONLY:       if.then2.i:
-// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    br label [[_ZL3YNFIF_EXIT]]
 // FINITEONLY:       if.end4.i:
-// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT:    br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
 // FINITEONLY:       for.body.i:
-// FINITEONLY-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
 // FINITEONLY:       _ZL3ynfif.exit:
-// FINITEONLY-NEXT:    [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT:    [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
 // FINITEONLY-NEXT:    ret float [[RETVAL_0_I]]
 //
 extern "C" __device__ float test_ynf(int x, float y) {
   return ynf(x, y);
 }
 
 // DEFAULT-LABEL: @test_yn(
 // DEFAULT-NEXT:  entry:
 // DEFAULT-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
 // DEFAULT-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
 // DEFAULT-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
 // DEFAULT-NEXT:    ]
 // DEFAULT:       if.then.i:
-// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
 // DEFAULT-NEXT:    br label [[_ZL2YNID_EXIT:%.*]]
 // DEFAULT:       if.then2.i:
-// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
 // DEFAULT-NEXT:    br label [[_ZL2YNID_EXIT]]
 // DEFAULT:       if.end4.i:
-// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT:    [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT:    br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
 // DEFAULT:       for.body.i:
-// DEFAULT-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
 // DEFAULT-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
 // DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
-// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT:    [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT:    [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
 // DEFAULT:       _ZL2ynid.exit:
-// DEFAULT-NEXT:    [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT:    [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
 // DEFAULT-NEXT:    ret double [[RETVAL_0_I]]
 //
 // FINITEONLY-LABEL: @test_yn(
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
 // FINITEONLY-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
 // FINITEONLY-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
 // FINITEONLY-NEXT:    ]
 // FINITEONLY:       if.then.i:
-// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    br label [[_ZL2YNID_EXIT:%.*]]
 // FINITEONLY:       if.then2.i:
-// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    br label [[_ZL2YNID_EXIT]]
 // FINITEONLY:       if.end4.i:
-// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT:    br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
 // FINITEONLY:       for.body.i:
-// FINITEONLY-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT:    [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
 // FINITEONLY:       _ZL2ynid.exit:
-// FINITEONLY-NEXT:    [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT:    [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
 // FINITEONLY-NEXT:    ret double [[RETVAL_0_I]]
 //
 extern "C" __device__ double test_yn(int x, double y) {
   return yn(x, y);
 }
 
 // DEFAULT-LABEL: @test___cosf(
 // DEFAULT-NEXT:  entry:

clang/test/OpenMP/bug57757.cpp

 // CHECK-NEXT:    ret void
 //
 //
 // CHECK-LABEL: define {{[^@]+}}@.omp_task_entry.
 // CHECK-SAME: (i32 noundef [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP1]], i64 0, i32 2
 // CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META13:![0-9]+]])
-// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META16:![0-9]+]])
-// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4, !tbaa [[TBAA18:![0-9]+]], !alias.scope !13, !noalias !16
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4, !tbaa [[TBAA16:![0-9]+]], !alias.scope !13, !noalias !17
 // CHECK-NEXT:    switch i32 [[TMP3]], label [[DOTOMP_OUTLINED__EXIT:%.*]] [
 // CHECK-NEXT:    i32 0, label [[DOTUNTIED_JMP__I:%.*]]
 // CHECK-NEXT:    i32 1, label [[DOTUNTIED_NEXT__I:%.*]]
 // CHECK-NEXT:    ]
 // CHECK:       .untied.jmp..i:
-// CHECK-NEXT:    store i32 1, ptr [[TMP2]], align 4, !tbaa [[TBAA18]], !alias.scope !13, !noalias !16
-// CHECK-NEXT:    [[TMP4:%.*]] = tail call i32 @__kmpc_omp_task(ptr nonnull @[[GLOB1]], i32 [[TMP0]], ptr [[TMP1]]), !noalias !19
+// CHECK-NEXT:    store i32 1, ptr [[TMP2]], align 4, !tbaa [[TBAA16]], !alias.scope !13, !noalias !17
+// CHECK-NEXT:    [[TMP4:%.*]] = tail call i32 @__kmpc_omp_task(ptr nonnull @[[GLOB1]], i32 [[TMP0]], ptr [[TMP1]]), !noalias !13
 // CHECK-NEXT:    br label [[DOTOMP_OUTLINED__EXIT]]
 // CHECK:       .untied.next..i:
 // CHECK-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP1]], i64 0, i32 1
 // CHECK-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP1]], i64 0, i32 1, i32 2
 // CHECK-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP1]], i64 0, i32 1, i32 1
-// CHECK-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[TMP5]], align 8, !tbaa [[TBAA20:![0-9]+]], !alias.scope !16, !noalias !13
-// CHECK-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP7]], align 4, !tbaa [[TBAA18]], !alias.scope !16, !noalias !13
-// CHECK-NEXT:    [[TMP10:%.*]] = load float, ptr [[TMP6]], align 4, !tbaa [[TBAA21:![0-9]+]], !alias.scope !16, !noalias !13
-// CHECK-NEXT:    tail call void [[TMP8]](i32 noundef [[TMP9]], float noundef [[TMP10]]) #[[ATTR2:[0-9]+]], !noalias !19
+// CHECK-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[TMP5]], align 8, !tbaa [[TBAA19:![0-9]+]], !noalias !13
+// CHECK-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP7]], align 4, !tbaa [[TBAA16]], !noalias !13
+// CHECK-NEXT:    [[TMP10:%.*]] = load float, ptr [[TMP6]], align 4, !tbaa [[TBAA20:![0-9]+]], !noalias !13
+// CHECK-NEXT:    tail call void [[TMP8]](i32 noundef [[TMP9]], float noundef [[TMP10]]) #[[ATTR2:[0-9]+]], !noalias !13
 // CHECK-NEXT:    br label [[DOTOMP_OUTLINED__EXIT]]
 // CHECK:       .omp_outlined..exit:
 // CHECK-NEXT:    ret i32 0
 //

llvm/lib/Passes/PassBuilderPipelines.cpp

 static cl::opt<bool> EnableMemProfiler("enable-mem-prof", cl::Hidden,
                                        cl::desc("Enable memory profiler"));
 
 static cl::opt<bool> EnableModuleInliner("enable-module-inliner",
                                          cl::init(false), cl::Hidden,
                                          cl::desc("Enable module inliner"));
 
 static cl::opt<bool> PerformMandatoryInliningsFirst(
-    "mandatory-inlining-first", cl::init(true), cl::Hidden,
+    "mandatory-inlining-first", cl::init(false), cl::Hidden,
     cl::desc("Perform mandatory inlinings module-wide, before performing "
              "inlining"));
 
 static cl::opt<bool> EnableO3NonTrivialUnswitching(
     "enable-npm-O3-nontrivial-unswitch", cl::init(true), cl::Hidden,
     cl::desc("Enable non-trivial loop unswitching for -O3"));
 
 static cl::opt<bool> EnableEagerlyInvalidateAnalyses(
PassBuilder::buildModuleSimplificationPipeline(OptimizationLevel Level,
   if (PGOOpt && Phase != ThinOrFullLTOPhase::ThinLTOPostLink &&
       PGOOpt->CSAction == PGOOptions::CSIRInstr)
     MPM.addPass(PGOInstrumentationGenCreateVar(PGOOpt->CSProfileGenFile));
 
   // Synthesize function entry counts for non-PGO compilation.
   if (EnableSyntheticCounts && !PGOOpt)
     MPM.addPass(SyntheticCountsPropagation());
 
+  MPM.addPass(AlwaysInlinerPass(
+      /*InsertLifetimeIntrinsics=*/Level != OptimizationLevel::O0));
+
   if (EnableModuleInliner)
     MPM.addPass(buildModuleInlinerPipeline(Level, Phase));
   else
     MPM.addPass(buildInlinerPipeline(Level, Phase));
 
   // Remove any dead arguments exposed by cleanups, constant folding globals,
   // and argument promotion.
   MPM.addPass(DeadArgumentEliminationPass());

llvm/test/Other/new-pm-defaults.ll

 ; CHECK-O-NEXT: Running analysis: OptimizationRemarkEmitterAnalysis
 ; CHECK-O-NEXT: Running analysis: AAManager
 ; CHECK-O-NEXT: Running analysis: BasicAA
 ; CHECK-O-NEXT: Running analysis: ScopedNoAliasAA
 ; CHECK-O-NEXT: Running analysis: TypeBasedAA
 ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
 ; CHECK-EP-PEEPHOLE-NEXT: Running pass: NoOpFunctionPass
 ; CHECK-O-NEXT: Running pass: SimplifyCFGPass
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
+; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
 ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
 ; CHECK-O-NEXT: Running analysis: GlobalsAA
 ; CHECK-O-NEXT: Running analysis: CallGraphAnalysis
 ; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
 ; CHECK-O-NEXT: Invalidating analysis: AAManager
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
-; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
 ; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis
 ; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy
 ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}>
 ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
 ; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
 ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
 ; CHECK-O-NEXT: Running analysis: AAManager
 ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
 ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass on (foo)
 ; CHECK-O3-NEXT: Running pass: OpenMPOptCGSCCPass on (foo)
 ; CHECK-EP-CGSCC-LATE-NEXT: Running pass: NoOpCGSCCPass
 ; CHECK-O-NEXT: Running pass: SROAPass
 ; CHECK-O-NEXT: Running pass: EarlyCSEPass

llvm/test/Other/new-pm-print-pipeline.ll

 
 ; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(loop-vectorize<no-interleave-forced-only;no-vectorize-forced-only>,loop-vectorize<interleave-forced-only;vectorize-forced-only>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-19
 ; CHECK-19: function(loop-vectorize<no-interleave-forced-only;no-vectorize-forced-only;>,loop-vectorize<interleave-forced-only;vectorize-forced-only;>)
 
 ; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='inliner-wrapper,inliner-wrapper-no-mandatory-first' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-20
 ; CHECK-20: cgscc(inline<only-mandatory>,inline),cgscc(inline)
 
 ; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='scc-oz-module-inliner' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-21
-; CHECK-21: require<globals-aa>,function(invalidate<aa>),require<profile-summary>,cgscc(devirt<4>(inline<only-mandatory>,inline,{{.*}},instcombine{{.*}}))
+; CHECK-21: require<globals-aa>,function(invalidate<aa>),require<profile-summary>,cgscc(devirt<4>(inline,{{.*}},instcombine{{.*}}))
 
 ; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='cgscc(function<eager-inv>(no-op-function)),function<eager-inv>(no-op-function)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-22
 ; CHECK-22: cgscc(function<eager-inv>(no-op-function)),function<eager-inv>(no-op-function)
 
 ;; Test that the loop-nest-pass lnicm is printed with the other loop-passes in the pipeline.
 ; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(loop-mssa(licm,loop-rotate,loop-deletion,lnicm,loop-rotate))' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-23
 ; CHECK-23: function(loop-mssa(licm<allowspeculation>,loop-rotate,loop-deletion,lnicm<allowspeculation>,loop-rotate))
 

llvm/test/Other/new-pm-thinlto-defaults.ll

 ; CHECK-O-NEXT: Running pass: InstCombinePass
 ; CHECK-PRELINK-O-NEXT: Running analysis: OptimizationRemarkEmitterAnalysis
 ; CHECK-O-NEXT: Running analysis: AAManager
 ; CHECK-O-NEXT: Running analysis: BasicAA
 ; CHECK-O-NEXT: Running analysis: ScopedNoAliasAA
 ; CHECK-O-NEXT: Running analysis: TypeBasedAA
 ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
 ; CHECK-O-NEXT: Running pass: SimplifyCFGPass
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
+; CHECK-PRELINK-O-NEXT: Running analysis: ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
 ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
 ; CHECK-O-NEXT: Running analysis: GlobalsAA
 ; CHECK-O-NEXT: Running analysis: CallGraphAnalysis
 ; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
 ; CHECK-O-NEXT: Invalidating analysis: AAManager
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
-; CHECK-PRELINK-O-NEXT: Running analysis: ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
 ; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis
 ; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy
 ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
 ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
 ; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
 ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
 ; CHECK-O-NEXT: Running analysis: AAManager
 ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
 ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass on (foo)
 ; CHECK-O3-NEXT: Running pass: OpenMPOptCGSCCPass on (foo)
 ; CHECK-O-NEXT: Running pass: SROAPass
 ; CHECK-O-NEXT: Running pass: EarlyCSEPass
 ; CHECK-O-NEXT: Running analysis: MemorySSAAnalysis

llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll

 ; CHECK-O-NEXT: Running analysis: BlockFrequencyAnalysis on foo
 ; These next two can appear in any order since they are accessed as parameters
 ; on the same call to BlockFrequencyInfo::calculate.
 ; CHECK-OSZ-DAG: Running analysis: LoopAnalysis on foo
 ; CHECK-OSZ-DAG: Running analysis: BranchProbabilityAnalysis on foo
 ; CHECK-O123-NEXT: Running analysis: BranchProbabilityAnalysis on foo
 ; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo
 ; CHECK-O-NEXT: Running pass: SimplifyCFGPass
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
 ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
 ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
 ; CHECK-O-NEXT: Running analysis: GlobalsAA
 ; CHECK-O-NEXT: Running analysis: CallGraphAnalysis
 ; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
 ; CHECK-O-NEXT: Invalidating analysis: AAManager
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
 ; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis
 ; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy
 ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}>
 ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
 ; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
 ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
 ; CHECK-O-NEXT: Running analysis: AAManager
 ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
 ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass
 ; CHECK-O3-NEXT: Running pass: OpenMPOptCGSCCPass
 ; CHECK-O-NEXT: Running pass: SROAPass
 ; CHECK-O-NEXT: Running pass: EarlyCSEPass
 ; CHECK-O-NEXT: Running analysis: MemorySSAAnalysis

llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll

 ; These next two can appear in any order since they are accessed as parameters
 ; on the same call to BlockFrequencyInfo::calculate.
 ; CHECK-OSZ-DAG: Running analysis: LoopAnalysis on foo
 ; CHECK-OSZ-DAG: Running analysis: BranchProbabilityAnalysis on foo
 ; CHECK-O123-NEXT: Running analysis: BranchProbabilityAnalysis on foo
 ; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo
 ; CHECK-O-NEXT: Running pass: SimplifyCFGPass on foo
 
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
 ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
 ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
 ; CHECK-O-NEXT: Running analysis: GlobalsAA
 ; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
 ; CHECK-O-NEXT: Invalidating analysis: AAManager
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
 ; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis
 ; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy
 ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
 ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
 ; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
 ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
 ; CHECK-O-NEXT: Running analysis: AAManager
 ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
 ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass
 ; CHECK-O3-NEXT: Running pass: OpenMPOptCGSCCPass
 ; CHECK-O-NEXT: Running pass: SROAPass
 ; CHECK-O-NEXT: Running pass: EarlyCSEPass
 ; CHECK-O-NEXT: Running analysis: MemorySSAAnalysis

llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll

 ; CHECK-O-NEXT: Running analysis: BlockFrequencyAnalysis on foo
 ; CHECK-O-NEXT: Invalidating analysis: InnerAnalysisManagerProxy
 ; CHECK-O123SZ-NEXT: Invalidating analysis: LazyCallGraphAnalysis on
 ; CHECK-O123SZ-NEXT: Invalidating analysis: InnerAnalysisManagerProxy
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running pass: PGOIndirectCallPromotion on
 ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
 ; CHECK-O-NEXT: Running analysis: OptimizationRemarkEmitterAnalysis on foo
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
 ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
 ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
 ; CHECK-O-NEXT: Running analysis: GlobalsAA
 ; CHECK-O-NEXT: Running analysis: CallGraphAnalysis
 ; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
 ; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis
 ; CHECK-O-NEXT: Running analysis: TargetLibraryAnalysis on foo
 ; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy
 ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}>
 ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
 ; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
 ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
 ; CHECK-O-NEXT: Running analysis: AAManager
 ; CHECK-O-NEXT: Running analysis: BasicAA
 ; CHECK-O-NEXT: Running analysis: AssumptionAnalysis
 ; CHECK-O-NEXT: Running analysis: TargetIRAnalysis
 ; CHECK-O-NEXT: Running analysis: DominatorTreeAnalysis
 ; CHECK-O-NEXT: Running analysis: ScopedNoAliasAA
 ; CHECK-O-NEXT: Running analysis: TypeBasedAA

llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll

 ; CHECK-O-NEXT: Running analysis: BlockFrequencyAnalysis on foo
 ; These next two can appear in any order since they are accessed as parameters
 ; on the same call to BlockFrequencyInfo::calculate.
 ; CHECK-OSZ-DAG: Running analysis: LoopAnalysis on foo
 ; CHECK-OSZ-DAG: Running analysis: BranchProbabilityAnalysis on foo
 ; CHECK-O123-NEXT: Running analysis: BranchProbabilityAnalysis on foo
 ; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo
 ; CHECK-O-NEXT: Running pass: SimplifyCFGPass on foo
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
 ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
 ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
 ; CHECK-O-NEXT: Running analysis: GlobalsAA
 ; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
 ; CHECK-O-NEXT: Invalidating analysis: AAManager
 ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
 ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
 ; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis
 ; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy
 ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}>
 ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
 ; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
 ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
 ; CHECK-O-NEXT: Running analysis: AAManager
 ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
 ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass
 ; CHECK-O3-NEXT: Running pass: OpenMPOptCGSCCPass
 ; CHECK-O-NEXT: Running pass: SROAPass
 ; CHECK-O-NEXT: Running pass: EarlyCSEPass
 ; CHECK-O-NEXT: Running analysis: MemorySSAAnalysis

llvm/test/Transforms/Inline/always-inline-newpm.ll

+; RUN: opt --Os -pass-remarks=inline -S < %s 2>&1 | FileCheck %s
+target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
+target triple = "arm64e-apple-macosx13"
+
+; CHECK: remark: <unknown>:0:0: 'wibble' inlined into 'bar.8' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'wibble' inlined into 'pluto' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'snork' inlined into 'blam' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'wobble' inlined into 'blam' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'wobble' inlined into 'snork' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'spam' inlined into 'blam' with (cost=65, threshold=75)
+; CHECK: remark: <unknown>:0:0: 'wibble.1' inlined into 'widget' with (cost=30, threshold=75)
+; CHECK: remark: <unknown>:0:0: 'widget' inlined into 'bar.8' with (cost=30, threshold=75)
+; CHECK: remark: <unknown>:0:0: 'barney' inlined into 'wombat' with (cost=30, threshold=75)
+
+define linkonce_odr void @wombat(ptr %arg) #0 {
+bb:
+  call void @barney()
+  ret void
+}
+
+define i1 @foo() {
+bb:
+  call void @wombat(ptr null)
+  unreachable
+}
+
+define linkonce_odr void @pluto() #1 !prof !38 {
+bb:
+  call void @wibble()
+  ret void
+}
+
+; Function Attrs: alwaysinline
+define linkonce_odr void @wibble() #2 {
+bb:
+  call void @widget()
+  ret void
+}
+
+define linkonce_odr void @widget() {
+bb:
+  call void @wibble.1()
+  ret void
+}
+
+define linkonce_odr void @wibble.1() {
+bb:
+  %0 = call i32 @foo.2()
+  call void @blam()
+  ret void
+}
+
+declare i32 @foo.2()
+
+define linkonce_odr void @blam() {
+bb:
+  %tmp = call i32 @snork()
+  %tmpv1 = call ptr @wombat.3()
+  call void @eggs()
+  %tmpv2 = call ptr @wombat.3()
+  ret void
+}
+
+; Function Attrs: alwaysinline
+define linkonce_odr i32 @snork() #2 {
+bb:
+  %tmpv1 = call i32 @spam()
+  %tmpv2 = call i32 @wobble()
+  call void @widget.4(i32 %tmpv2)
+  ret i32 0
+}
+
+declare void @eggs()
+
+declare ptr @wombat.3()
+
+define linkonce_odr i32 @spam() {
+bb:
+  %tmpv1 = call i32 @wombat.6()
+  %tmpv2 = call i64 @wobble.5(i8 0)
+  %tmpv3 = call i64 @bar()
+  ret i32 0
+}
+
+; Function Attrs: alwaysinline
+define linkonce_odr i32 @wobble() #2 {
+bb:
+  %tmpv = call i64 @wobble.5(i8 0)
+  %tmpv1 = call i64 @eggs.7()
+  %tmpv2 = call i64 @wobble.5(i8 0)
+  %tmpv3 = call i64 @eggs.7()
+  %tmpv4 = lshr i64 %tmpv1, 1
+  %tmpv5 = trunc i64 %tmpv4 to i32
+  %tmpv6 = xor i32 %tmpv5, 23
+  ret i32 %tmpv6
+}
+
+declare void @widget.4(i32)
+
+declare i64 @bar()
+
+declare i64 @wobble.5(i8)
+
+declare i32 @wombat.6()
+
+declare i64 @eggs.7()
+
+define linkonce_odr void @barney() {
+bb:
+  call void @bar.8()
+  call void @pluto()
+  unreachable
+}
+
+define linkonce_odr void @bar.8() {
+bb:
+  call void @wibble()
+  ret void
+}
+
+attributes #0 = { "frame-pointer"="non-leaf" }
+attributes #1 = { "target-cpu"="apple-m1" }
+attributes #2 = { alwaysinline }
+
+!llvm.module.flags = !{!0, !1, !30, !31, !32, !36, !37}
+
+!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 13, i32 3]}
+!1 = !{i32 1, !"ProfileSummary", !2}
+!2 = !{!3, !4, !5, !6, !7, !8, !9, !10, !11, !12}
+!3 = !{!"ProfileFormat", !"InstrProf"}
+!4 = !{!"TotalCount", i64 864540306756}
+!5 = !{!"MaxCount", i64 6596759955}
+!6 = !{!"MaxInternalCount", i64 2828618424}
+!7 = !{!"MaxFunctionCount", i64 6596759955}
+!8 = !{!"NumCounts", i64 268920}
+!9 = !{!"NumFunctions", i64 106162}
+!10 = !{!"IsPartialProfile", i64 0}
+!11 = !{!"PartialProfileRatio", double 0.000000e+00}
+!12 = !{!"DetailedSummary", !13}
+!13 = !{!14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29}
+!14 = !{i32 10000, i64 5109654023, i32 2}
+!15 = !{i32 100000, i64 2480859832, i32 25}
+!16 = !{i32 200000, i64 1566552109, i32 70}
+!17 = !{i32 300000, i64 973667919, i32 140}
+!18 = !{i32 400000, i64 552159773, i32 263}
+!19 = !{i32 500000, i64 353879860, i32 463}
+!20 = !{i32 600000, i64 187122455, i32 799}
+!21 = !{i32 700000, i64 105465980, i32 1419}
+!22 = !{i32 800000, i64 49243829, i32 2620}
+!23 = !{i32 900000, i64 15198227, i32 5898}
+!24 = !{i32 950000, i64 5545670, i32 10696}
+!25 = !{i32 990000, i64 804816, i32 25738}
+!26 = !{i32 999000, i64 73999, i32 53382}
+!27 = !{i32 999900, i64 6530, i32 83503}
+!28 = !{i32 999990, i64 899, i32 110416}
+!29 = !{i32 999999, i64 120, i32 130201}
+!30 = !{i32 7, !"Dwarf Version", i32 4}
+!31 = !{i32 2, !"Debug Info Version", i32 3}
+!32 = !{i32 1, !"wchar_size", i32 4}
+!34 = !{!35}
+!35 = !{i32 0, i1 false}
+!36 = !{i32 8, !"PIC Level", i32 2}
+!37 = !{i32 7, !"frame-pointer", i32 1}
+!38 = !{!"function_entry_count", i64 15128150}

llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll

 unreachable:                                      ; preds = %cleanup
   unreachable
 }
 
 declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1
 
 attributes #0 = { nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="cortex-m55" "target-features"="+armv8.1-m.main,+dsp,+fp-armv8d16,+fp-armv8d16sp,+fp16,+fp64,+fullfp16,+hwdiv,+lob,+mve,+mve.fp,+ras,+strict-align,+thumb-mode,+vfp2,+vfp2sp,+vfp3d16,+vfp3d16sp,+vfp4d16,+vfp4d16sp,-aes,-bf16,-cdecp0,-cdecp1,-cdecp2,-cdecp3,-cdecp4,-cdecp5,-cdecp6,-cdecp7,-crc,-crypto,-d32,-dotprod,-fp-armv8,-fp-armv8sp,-fp16fml,-hwdiv-arm,-i8mm,-neon,-sb,-sha2,-vfp3,-vfp3sp,-vfp4,-vfp4sp" "unsafe-fp-math"="true" }
 attributes #1 = { argmemonly nofree nosync nounwind willreturn }
-attributes #2 = { alwaysinline nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="cortex-m55" "target-features"="+armv8.1-m.main,+dsp,+fp-armv8d16,+fp-armv8d16sp,+fp16,+fp64,+fullfp16,+hwdiv,+lob,+mve,+mve.fp,+ras,+strict-align,+thumb-mode,+vfp2,+vfp2sp,+vfp3d16,+vfp3d16sp,+vfp4d16,+vfp4d16sp,-aes,-bf16,-cdecp0,-cdecp1,-cdecp2,-cdecp3,-cdecp4,-cdecp5,-cdecp6,-cdecp7,-crc,-crypto,-d32,-dotprod,-fp-armv8,-fp-armv8sp,-fp16fml,-hwdiv-arm,-i8mm,-neon,-sb,-sha2,-vfp3,-vfp3sp,-vfp4,-vfp4sp" "unsafe-fp-math"="true" }
+attributes #2 = { nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="cortex-m55" "target-features"="+armv8.1-m.main,+dsp,+fp-armv8d16,+fp-armv8d16sp,+fp16,+fp64,+fullfp16,+hwdiv,+lob,+mve,+mve.fp,+ras,+strict-align,+thumb-mode,+vfp2,+vfp2sp,+vfp3d16,+vfp3d16sp,+vfp4d16,+vfp4d16sp,-aes,-bf16,-cdecp0,-cdecp1,-cdecp2,-cdecp3,-cdecp4,-cdecp5,-cdecp6,-cdecp7,-crc,-crypto,-d32,-dotprod,-fp-armv8,-fp-armv8sp,-fp16fml,-hwdiv-arm,-i8mm,-neon,-sb,-sha2,-vfp3,-vfp3sp,-vfp4,-vfp4sp" "unsafe-fp-math"="true" }
 attributes #3 = { nounwind }