[DebugInfo] Reduce SalvageDebugInfo() functions

[chrisjackson]

Summary

Having both salvageDebugInfo() and salvageDebugInfoOrMarkUndef() has led to situations when the incorrect function was called. Eliminate the potential for error by combining the two functions. First suggested in D76930.

Now all SalvageDebugInfo() calls will mark undef if the salvage attempt fails. This did not require changes to call sites, except the function name, but some tests resulted in extra debug info (undef) being generated where previously there was none.

[Orlando]

Removing a class of bug is always nice.

I have a question for the debug-info reviewers here. Does anyone know what these tests that had to be updated are for? Maybe I'm missing something, but they seem brittle and it's not obvious to me what they are checking for.

llvm/lib/Transforms/Utils/Local.cpp

1631

Is there any value in returning a result at all (i.e. is the return value used anywhere?).

IIRC the only reason the result of salvageDebugInfo was checked was to see if we need to replaceDbgUsesWithUndef (which is now done inside salvageDebugInfo.

[aprantl]

llvm/include/llvm/Transforms/Utils/Local.h

359

Should we rename this to updateDebugInfoForDeletedInsn() or some variant of that to make it more clear that it isn't optional to call this method?

[vsk]

Thanks! This lgtm pending outstanding reviewer comments.

In D78369#1989058, @Orlando wrote:

I have a question for the debug-info reviewers here. Does anyone know what these tests that had to be updated are for? Maybe I'm missing something, but they seem brittle and it's not obvious to me what they are checking for.

Looks like end-to-end tests for DWARF production under NVPTX. It's probably checking at a more granular level than it needs to. Paging @ABataev, would it be all right to rewrite the test so it pattern-matches against llvm-dwarfdump output?

llvm/include/llvm/Transforms/Utils/Local.h

359

I don't think the current name is unclear (I'm not opposed to changing it, just don't see a strong need to right now).

However, what /does/ seem missing is some way to enforce that this method is called. I like @alok's approach in D70419: wdyt of just making Instruction::eraseFromParent() call Instruction::salvageDebugInfo()? (This would be for a follow-up, my preference would be to consolidate APIs here first.)

[vsk]

llvm/include/llvm/Transforms/Utils/Local.h

359

Ah, one issue with this is that the common deletion idiom is: "I.RAUW(replacement); I.eraseFromParent();" -- so if the salageDebugInfo operation occurs in eraseFromParent, in many cases it's too late, because the debug uses are gone.

[ABataev]

In D78369#1989363, @vsk wrote:

Thanks! This lgtm pending outstanding reviewer comments.

In D78369#1989058, @Orlando wrote:

I have a question for the debug-info reviewers here. Does anyone know what these tests that had to be updated are for? Maybe I'm missing something, but they seem brittle and it's not obvious to me what they are checking for.

Looks like end-to-end tests for DWARF production under NVPTX. It's probably checking at a more granular level than it needs to. Paging @ABataev, would it be all right to rewrite the test so it pattern-matches against llvm-dwarfdump output?

No, llvm-dwarfdump works with object files and we cannot produce it for NVPTX target, it emits only .PTX assembler.

[jmorse]

LGTM with comments addressed,

llvm/include/llvm/Transforms/Utils/Local.h

359

I've no opinion on changing the name; for the "RAUW; Erase" deletion idiom I don't think that's a problem, it just means those locations were saved by the RAUW instead of through salvaging. A salvage call in eraseFromParent would be highly useful in catching the more complicated circumstances, for example I believe that during inlining large numbers of instructions are pruned for being unused, without ever being RAUW'd.

[chrisjackson]

llvm/include/llvm/Transforms/Utils/Local.h

359

Yes this is something I'd like to address too.

359

I'm not sure that a method name is a good way of reflecting its necessity.

[chrisjackson]

Removed unused return value from SalvageDebugInfo().

[dblaikie]

llvm/include/llvm/Transforms/Utils/Local.h

359

I'll +1 @aprantl's comment - it sounds like at least from the comment the previous version of this function was "best effort" - if you didn't call it you just got worse debug info. But the new behavior seems like it might break (result in incorrect debug info) if you don't call this function? Is that the case? (if so, then a new name seems justified)

[Orlando]

llvm/include/llvm/Transforms/Utils/Local.h

359

it sounds like at least from the comment the previous version of this
function was "best effort" - if you didn't call it you just got worse debug
info.

This is only true if replaceDbgUsesWithUndef was being called where
appropriate. With D80264 we do actually get this for free in a lot of cases,
but there are still times where we want to undef a dbg.value without deleting
the instruction. E.g. When sinking an instruction with a corresponding
dbg.value, we want to leave behind a salvage-or-undef, and sink a clone
of the original dbg.value to the new instruction position.

[dblaikie]

llvm/include/llvm/Transforms/Utils/Local.h

359

it sounds like at least from the comment the previous version of this

function was "best effort" - if you didn't call it you just got worse debug
info.

This is only true if replaceDbgUsesWithUndef was being called where

appropriate.

Sorry, not sure I understand - you mean previously you had to call one of either replaceDbgUsesWithUndef (not anymore, since D80264, now that's implicit) or salvageDebugInfo - and now you have to always call salvageDebugInfo ?

(maybe I shuold just bow out of this review - I haven't gotten very involved in LLVM's debug info variable location tracking, so I don't have a lot of context for this - just wanted to express some agreement with @aprantl's perspective there)

[Orlando]

llvm/include/llvm/Transforms/Utils/Local.h

359

Ah sorry for being unclear, I hope this helps. I don't have a strong opinion on the name change.

AFAIK it is correct to call replaceDbgUsesWithUndef if salvageDebugInfo failed (returns false), and not doing so can lead to wrong debug-info by allowing previous dbg.value definitions to run on for too long. This patch changes salvageDebugInfo to call replaceDbgUsesWithUndef itself if it fails to salvage. Because replaceDbgUsesWithUndef is now called automatically when we delete an instruction (D80264), not calling salvageDebugInfo just reduces coverage, but doesn't produce incorrect debug-info when deleting instructions.

But the new behavior seems like it might break (result in incorrect debug info) if you don't call this function? Is that the case? (if so, then a new name seems justified)

So I think the answer to this is no, now that we have D80264.

[Orlando]

LGTM. There are two other text only LGTMs on this patch and AFAICT the only blocker is the (inline) discussion on whether salvageDebugInfo needs renaming, but there doesn't seem to be a consensus. Does anyone have a strong opinion on this subject?

I feel like everyone agrees that the patch is good, but I don't want to press the button and to have inadvertently shut someone down.

[vsk]

This looks good to me. We should consolidate the salvageDebugInfo* APIs. I think the conversation about renaming the API is a bit moot post-D80264 (afaict the motivation for the rename was that /not/ calling the API could result in invalid debug info -- that should be fixed now).

llvm/include/llvm/Transforms/Utils/Local.h

356

Maybe 's/mark it/change its debug uses to/'?

[chrisjackson]

This revision was closed with commit c6c65164. I'm uncertain why the commit's 'differential revision' tag failed to trigger the autoclose.

llvm/include/llvm/Transforms/Utils/Local.h

 void replaceDbgValueForAlloca(AllocaInst *AI, Value *NewAllocaAddress,
                               DIBuilder &Builder, int Offset = 0);
 
 /// Finds alloca where the value comes from.
 AllocaInst *findAllocaForValue(Value *V,
                                DenseMap<Value *, AllocaInst *> &AllocaForValue);
 
 /// Assuming the instruction \p I is going to be deleted, attempt to salvage
-/// debug users of \p I by writing the effect of \p I in a DIExpression.
+/// debug users of \p I by writing the effect of \p I in a DIExpression. If it
+/// cannot be salvaged mark it undef.
 /// Returns true if any debug users were updated.
 bool salvageDebugInfo(Instruction &I);
 
-/// Salvage all debug users of the instruction \p I or mark it as undef if it
-/// cannot be salvaged.
-void salvageDebugInfoOrMarkUndef(Instruction &I);
 
 /// Implementation of salvageDebugInfo, applying only to instructions in
 /// \p Insns, rather than all debug users of \p I.
 bool salvageDebugInfoForDbgValues(Instruction &I,
                                   ArrayRef<DbgVariableIntrinsic *> Insns);
 
 /// Given an instruction \p I and DIExpression \p DIExpr operating on it, write
 /// the effects of \p I into the returned DIExpression, or return nullptr if

llvm/lib/Transforms/InstCombine/InstCombineInternal.h

public:
   /// Combiner aware instruction erasure.
   ///
   /// When dealing with an instruction that has side effects or produces a void
   /// value, we can't rely on DCE to delete the instruction. Instead, visit
   /// methods should return the value returned by this function.
   Instruction *eraseInstFromFunction(Instruction &I) {
     LLVM_DEBUG(dbgs() << "IC: ERASE " << I << '\n');
     assert(I.use_empty() && "Cannot erase instruction that is used!");
-    salvageDebugInfoOrMarkUndef(I);
+    salvageDebugInfo(I);
 
     // Make sure that we reprocess all operands now that we reduced their
     // use counts.
     for (Use &Operand : I.operands())
       if (auto *Inst = dyn_cast<Instruction>(Operand))
         Worklist.add(Inst);
 
     Worklist.remove(&I);

llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp

bool InstCombiner::SimplifyDemandedBits(Instruction *I, unsigned OpNo,
                                         const APInt &DemandedMask,
                                         KnownBits &Known,
                                         unsigned Depth) {
   Use &U = I->getOperandUse(OpNo);
   Value *NewVal = SimplifyDemandedUseBits(U.get(), DemandedMask, Known,
                                           Depth, I);
   if (!NewVal) return false;
   if (Instruction* OpInst = dyn_cast<Instruction>(U))
-    salvageDebugInfoOrMarkUndef(*OpInst);
+    salvageDebugInfo(*OpInst);
     
   replaceUse(U, NewVal);
   return true;
 }
 
 
 /// This function attempts to replace V with a simpler value based on the
 /// demanded bits. When this function is called, it is known that only the bits

llvm/lib/Transforms/InstCombine/InstructionCombining.cpp

static bool prepareICWorklistFromFunction(Function &F, const DataLayout &DL,
   // some N^2 behavior in pathological cases.
   ICWorklist.reserve(InstrsForInstCombineWorklist.size());
   for (Instruction *Inst : reverse(InstrsForInstCombineWorklist)) {
     // DCE instruction if trivially dead. As we iterate in reverse program
     // order here, we will clean up whole chains of dead instructions.
     if (isInstructionTriviallyDead(Inst, TLI)) {
       ++NumDeadInst;
       LLVM_DEBUG(dbgs() << "IC: DCE: " << *Inst << '\n');
-      salvageDebugInfoOrMarkUndef(*Inst);
+      salvageDebugInfo(*Inst);
       Inst->eraseFromParent();
       MadeIRChange = true;
       continue;
     }
 
     ICWorklist.push(Inst);
   }
 

llvm/lib/Transforms/Scalar/BDCE.cpp

if (I.mayHaveSideEffects() && I.use_empty())
       continue;
 
     // Remove instructions that are dead, either because they were not reached
     // during analysis or have no demanded bits.
     if (DB.isInstructionDead(&I) ||
         (I.getType()->isIntOrIntVectorTy() &&
          DB.getDemandedBits(&I).isNullValue() &&
          wouldInstructionBeTriviallyDead(&I))) {
-      salvageDebugInfoOrMarkUndef(I);
+      salvageDebugInfo(I);
       Worklist.push_back(&I);
       I.dropAllReferences();
       Changed = true;
       continue;
     }
 
     for (Use &U : I.operands()) {
       // DemandedBits only detects dead integer uses.

llvm/lib/Transforms/Scalar/DeadStoreElimination.cpp

do {
     Instruction *DeadInst = NowDeadInsts.pop_back_val();
     // Mark the DeadInst as dead in the list of throwable instructions.
     auto It = ThrowableInst.find(DeadInst);
     if (It != ThrowableInst.end())
       ThrowableInst[It->first] = false;
     ++NumFastOther;
 
     // Try to preserve debug information attached to the dead instruction.
-    salvageDebugInfoOrMarkUndef(*DeadInst);
+    salvageDebugInfo(*DeadInst);
     salvageKnowledge(DeadInst);
 
     // This instruction is dead, zap it, in stages.  Start by removing it from
     // MemDep, which needs to know the operands and needs it to be in the
     // function.
     MD.removeInstruction(DeadInst);
 
     for (unsigned op = 0, e = DeadInst->getNumOperands(); op != e; ++op) {

llvm/lib/Transforms/Scalar/EarlyCSE.cpp

for (Instruction &Inst : make_early_inc_range(BB->getInstList())) {
     if (isInstructionTriviallyDead(&Inst, &TLI)) {
       LLVM_DEBUG(dbgs() << "EarlyCSE DCE: " << Inst << '\n');
       if (!DebugCounter::shouldExecute(CSECounter)) {
         LLVM_DEBUG(dbgs() << "Skipping due to debug counter\n");
         continue;
       }
 
       salvageKnowledge(&Inst, &AC);
-      salvageDebugInfoOrMarkUndef(Inst);
+      salvageDebugInfo(Inst);
       removeMSSA(Inst);
       Inst.eraseFromParent();
       Changed = true;
       ++NumSimplify;
       continue;
     }
 
     // Skip assume intrinsics, they don't really have side effects (although

llvm/lib/Transforms/Scalar/LICM.cpp

for (BasicBlock::iterator II = BB->end(); II != BB->begin();) {
       bool FreeInLoop = false;
       if (isNotUsedOrFreeInLoop(I, CurLoop, SafetyInfo, TTI, FreeInLoop) &&
           canSinkOrHoistInst(I, AA, DT, CurLoop, CurAST, MSSAU, true, &Flags,
                              ORE) &&
           !I.mayHaveSideEffects()) {
         if (sink(I, LI, DT, CurLoop, SafetyInfo, MSSAU, ORE)) {
           if (!FreeInLoop) {
             ++II;
-            salvageDebugInfoOrMarkUndef(I);
+            salvageDebugInfo(I);
             eraseInstruction(I, *SafetyInfo, CurAST, MSSAU);
           }
           Changed = true;
         }
       }
     }
   }
   if (MSSAU && VerifyMemorySSA)

llvm/lib/Transforms/Scalar/Reassociate.cpp

 // Insts so they will be removed as well.
 void ReassociatePass::RecursivelyEraseDeadInsts(Instruction *I,
                                                 OrderedSet &Insts) {
   assert(isInstructionTriviallyDead(I) && "Trivially dead instructions only!");
   SmallVector<Value *, 4> Ops(I->op_begin(), I->op_end());
   ValueRankMap.erase(I);
   Insts.remove(I);
   RedoInsts.remove(I);
-  llvm::salvageDebugInfoOrMarkUndef(*I);
+  llvm::salvageDebugInfo(*I);
   I->eraseFromParent();
   for (auto Op : Ops)
     if (Instruction *OpInst = dyn_cast<Instruction>(Op))
       if (OpInst->use_empty())
         Insts.insert(OpInst);
 }
 
 /// Zap the given instruction, adding interesting operands to the work list.
 void ReassociatePass::EraseInst(Instruction *I) {
   assert(isInstructionTriviallyDead(I) && "Trivially dead instructions only!");
   LLVM_DEBUG(dbgs() << "Erasing dead inst: "; I->dump());
 
   SmallVector<Value*, 8> Ops(I->op_begin(), I->op_end());
   // Erase the dead instruction.
   ValueRankMap.erase(I);
   RedoInsts.remove(I);
-  llvm::salvageDebugInfoOrMarkUndef(*I);
+  llvm::salvageDebugInfo(*I);
   I->eraseFromParent();
   // Optimize its operands.
   SmallPtrSet<Instruction *, 8> Visited; // Detect self-referential nodes.
   for (unsigned i = 0, e = Ops.size(); i != e; ++i)
     if (Instruction *Op = dyn_cast<Instruction>(Ops[i])) {
       // If this is a node in an expression tree, climb to the expression root
       // and add that since that's where optimization actually happens.
       unsigned Opcode = Op->getOpcode();

llvm/lib/Transforms/Utils/Local.cpp

if (auto *MDV = MetadataAsValue::getIfExists(AI->getContext(), L))
       }
 }
 
 /// Wrap \p V in a ValueAsMetadata instance.
 static MetadataAsValue *wrapValueInMetadata(LLVMContext &C, Value *V) {
   return MetadataAsValue::get(C, ValueAsMetadata::get(V));
 }
 
-bool llvm::salvageDebugInfo(Instruction &I) {
+static bool attemptToSalvageDebugInfo(Instruction &I) {
   SmallVector<DbgVariableIntrinsic *, 1> DbgUsers;
   findDbgUsers(DbgUsers, &I);
   if (DbgUsers.empty())
     return false;
 
   return salvageDebugInfoForDbgValues(I, DbgUsers);
 }
 
-void llvm::salvageDebugInfoOrMarkUndef(Instruction &I) {
-  if (!salvageDebugInfo(I))
+bool llvm::salvageDebugInfo(Instruction &I) {
+  bool salvaged = attemptToSalvageDebugInfo(I);
+  if (!salvaged)
     replaceDbgUsesWithUndef(&I);
+  return salvaged;
 }
 
 bool llvm::salvageDebugInfoForDbgValues(
     Instruction &I, ArrayRef<DbgVariableIntrinsic *> DbgUsers) {
   auto &Ctx = I.getContext();
   auto wrapMD = [&](Value *V) { return wrapValueInMetadata(Ctx, V); };
 
   for (auto *DII : DbgUsers) {
for (auto *DII : Users) {
     DII->setOperand(0, wrapValueInMetadata(Ctx, &To));
     DII->setOperand(2, MetadataAsValue::get(Ctx, *DVR));
     LLVM_DEBUG(dbgs() << "REWRITE:  " << *DII << '\n');
     Changed = true;
   }
 
   if (!UndefOrSalvage.empty()) {
     // Try to salvage the remaining debug users.
-    salvageDebugInfoOrMarkUndef(From);
+    salvageDebugInfo(From);
     Changed = true;
   }
 
   return Changed;
 }
 
 /// Check if a bitcast between a value of type \p FromTy to type \p ToTy would
 /// losslessly preserve the bits and semantics of the value. This predicate is

llvm/test/DebugInfo/NVPTX/debug-addr-class.ll

 ; CHECK-NEXT: .b8 58                                  // DW_AT_decl_file
 ; CHECK-NEXT: .b8 11                                  // DW_FORM_data1
 ; CHECK-NEXT: .b8 59                                  // DW_AT_decl_line
 ; CHECK-NEXT: .b8 11                                  // DW_FORM_data1
 ; CHECK-NEXT: .b8 73                                  // DW_AT_type
 ; CHECK-NEXT: .b8 19                                  // DW_FORM_ref4
 ; CHECK-NEXT: .b8 0                                   // EOM(1)
 ; CHECK-NEXT: .b8 0                                   // EOM(2)
+; CHECK-NEXT: .b8 6                                   // Abbreviation Code
+; CHECK-NEXT: .b8 15                                  // DW_TAG_pointer_type
+; CHECK-NEXT: .b8 0                                   // DW_CHILDREN_no
+; CHECK-NEXT: .b8 73                                  // DW_AT_type
+; CHECK-NEXT: .b8 19                                  // DW_FORM_ref4
+; CHECK-NEXT: .b8 0                                   // EOM(1)
+; CHECK-NEXT: .b8 0                                   // EOM(2)
 ; CHECK-NEXT: .b8 0                                   // EOM(3)
 ; CHECK-NEXT: }
 ; CHECK-NEXT: .section .debug_info
 ; CHECK-NEXT: {
-; CHECK-NEXT: .b32 217                                // Length of Unit
+; CHECK-NEXT: .b32 240                                // Length of Unit
 ; CHECK-NEXT: .b8 2                                   // DWARF version number
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b32 .debug_abbrev                      // Offset Into Abbrev. Section
 ; CHECK-NEXT: .b8 8                                   // Address Size (in bytes)
-; CHECK-NEXT: .b8 1                                   // Abbrev [1] 0xb:0xd2 DW_TAG_compile_unit
+; CHECK-NEXT: .b8 1                                   // Abbrev [1] 0xb:0xe9 DW_TAG_compile_unit
 ; CHECK-NEXT: .b8 99                                  // DW_AT_producer
 ; CHECK-NEXT: .b8 108
 ; CHECK-NEXT: .b8 97
 ; CHECK-NEXT: .b8 110
 ; CHECK-NEXT: .b8 103
 ; CHECK-NEXT: .b8 32
 ; CHECK-NEXT: .b8 118
 ; CHECK-NEXT: .b8 101

 ; CHECK-NEXT: .b32 127                                // DW_AT_type
 ; CHECK-NEXT: .b8 1                                   // DW_AT_external
 ; CHECK-NEXT: .b8 1                                   // DW_AT_decl_file
 ; CHECK-NEXT: .b8 4                                   // DW_AT_decl_line
 ; CHECK-NEXT: .b8 8                                   // DW_AT_address_class
 ; CHECK-NEXT: .b8 9                                   // DW_AT_location
 ; CHECK-NEXT: .b8 3
 ; CHECK-NEXT: .b64 SHARED
-; CHECK-NEXT: .b8 4                                   // Abbrev [4] 0xa0:0x33 DW_TAG_subprogram
+; CHECK-NEXT: .b8 4                                   // Abbrev [4] 0xa0:0x45 DW_TAG_subprogram
 ; CHECK-NEXT: .b64 Lfunc_begin0                       // DW_AT_low_pc
 ; CHECK-NEXT: .b64 Lfunc_end0                         // DW_AT_high_pc
 ; CHECK-NEXT: .b8 1                                   // DW_AT_frame_base
 ; CHECK-NEXT: .b8 156
 ; CHECK-NEXT: .b8 116                                 // DW_AT_MIPS_linkage_name
 ; CHECK-NEXT: .b8 101
 ; CHECK-NEXT: .b8 115
 ; CHECK-NEXT: .b8 116
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b8 116                                 // DW_AT_name
 ; CHECK-NEXT: .b8 101
 ; CHECK-NEXT: .b8 115
 ; CHECK-NEXT: .b8 116
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b8 1                                   // DW_AT_decl_file
 ; CHECK-NEXT: .b8 6                                   // DW_AT_decl_line
 ; CHECK-NEXT: .b8 1                                   // DW_AT_external
 ; CHECK-NEXT: .b8 5                                   // Abbrev [5] 0xc0:0x9 DW_TAG_formal_parameter
 ; CHECK-NEXT: .b8 97                                  // DW_AT_name
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b8 1                                   // DW_AT_decl_file
 ; CHECK-NEXT: .b8 6                                   // DW_AT_decl_line
-; CHECK-NEXT: .b32 211                                // DW_AT_type
+; CHECK-NEXT: .b32 234                                // DW_AT_type
 ; CHECK-NEXT: .b8 5                                   // Abbrev [5] 0xc9:0x9 DW_TAG_formal_parameter
+; CHECK-NEXT: .b8 120                                 // DW_AT_name
+; CHECK-NEXT: .b8 0
+; CHECK-NEXT: .b8 1                                   // DW_AT_decl_file
+; CHECK-NEXT: .b8 6                                   // DW_AT_decl_line
+; CHECK-NEXT: .b32 229                                // DW_AT_type
+; CHECK-NEXT: .b8 5                                   // Abbrev [5] 0xd2:0x9 DW_TAG_formal_parameter
+; CHECK-NEXT: .b8 121                                 // DW_AT_name
+; CHECK-NEXT: .b8 0
+; CHECK-NEXT: .b8 1                                   // DW_AT_decl_file
+; CHECK-NEXT: .b8 6                                   // DW_AT_decl_line
+; CHECK-NEXT: .b32 229                                // DW_AT_type
+; CHECK-NEXT: .b8 5                                   // Abbrev [5] 0xdb:0x9 DW_TAG_formal_parameter
 ; CHECK-NEXT: .b8 105                                 // DW_AT_name
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b8 1                                   // DW_AT_decl_file
 ; CHECK-NEXT: .b8 6                                   // DW_AT_decl_line
 ; CHECK-NEXT: .b32 127                                // DW_AT_type
 ; CHECK-NEXT: .b8 0                                   // End Of Children Mark
-; CHECK-NEXT: .b8 3                                   // Abbrev [3] 0xd3:0x9 DW_TAG_base_type
+; CHECK-NEXT: .b8 6                                   // Abbrev [6] 0xe5:0x5 DW_TAG_pointer_type
+; CHECK-NEXT: .b32 234                                // DW_AT_type
+; CHECK-NEXT: .b8 3                                   // Abbrev [3] 0xea:0x9 DW_TAG_base_type
 ; CHECK-NEXT: .b8 102                                 // DW_AT_name
 ; CHECK-NEXT: .b8 108
 ; CHECK-NEXT: .b8 111
 ; CHECK-NEXT: .b8 97
 ; CHECK-NEXT: .b8 116
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b8 4                                   // DW_AT_encoding
 ; CHECK-NEXT: .b8 4                                   // DW_AT_byte_size
 ; CHECK-NEXT: .b8 0                                   // End Of Children Mark
-; CHECK-NEXT: }
-; CHECK-NEXT: .section .debug_loc { }
-; CHECK-NOT: debug_
-
+; CHECK-NEXT:   }
+; CHECK-NEXT:   .section  .debug_loc  { }

llvm/test/DebugInfo/NVPTX/debug-info.ll

 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
 
-; // Bitcode int this test case is reduced version of compiled code below:
+; // Bitcode in this test case is a reduced version of compiled code below:
 ;__device__ inline void res(float x, float y, float *res) { *res = x + y; }
 ;
 ;__global__ void saxpy(int n, float a, float *x, float *y) {
 ;  int i = blockIdx.x * blockDim.x + threadIdx.x;
 ;  if (i < n)
 ;    res(a * x[i], y[i], &y[i]);
 ;}
 

 ; CHECK-NEXT: .b8 49                               // DW_AT_abstract_origin
 ; CHECK-NEXT: .b8 19                               // DW_FORM_ref4
 ; CHECK-NEXT: .b8 0                                // EOM(1)
 ; CHECK-NEXT: .b8 0                                // EOM(2)
 ; CHECK-NEXT: .b8 0                                // EOM(3)
 ; CHECK-NEXT: }
 ; CHECK-NEXT: .section .debug_info
 ; CHECK-NEXT: {
-; CHECK-NEXT: .b32 10034                           // Length of Unit
+; CHECK-NEXT: .b32 10039                           // Length of Unit
 ; CHECK-NEXT: .b8 2                                // DWARF version number
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b32 .debug_abbrev                   // Offset Into Abbrev. Section
 ; CHECK-NEXT: .b8 8                                // Address Size (in bytes)
-; CHECK-NEXT: .b8 1                                // Abbrev [1] 0xb:0x272b DW_TAG_compile_unit
+; CHECK-NEXT: .b8 1                                // Abbrev [1] 0xb:0x2730 DW_TAG_compile_unit
 ; CHECK-NEXT: .b8 0                                // DW_AT_producer
 ; CHECK-NEXT: .b8 4                                // DW_AT_language
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b8 100                              // DW_AT_name
 ; CHECK-NEXT: .b8 101
 ; CHECK-NEXT: .b8 98
 ; CHECK-NEXT: .b8 117
 ; CHECK-NEXT: .b8 103

 ; CHECK-NEXT: .b8 114                              // DW_AT_name
 ; CHECK-NEXT: .b8 101
 ; CHECK-NEXT: .b8 115
 ; CHECK-NEXT: .b8 0
 ; CHECK-NEXT: .b8 12                               // DW_AT_decl_file
 ; CHECK-NEXT: .b8 3                                // DW_AT_decl_line
 ; CHECK-NEXT: .b32 3345                            // DW_AT_type
 ; CHECK-NEXT: .b8 0                                // End Of Children Mark
-; CHECK-NEXT: .b8 40                               // Abbrev [40] 0x2671:0xc4 DW_TAG_subprogram
+; CHECK-NEXT: .b8 40                               // Abbrev [40] 0x2671:0xc9 DW_TAG_subprogram
 ; CHECK-NEXT: .b64 Lfunc_begin0                    // DW_AT_low_pc
 ; CHECK-NEXT: .b64 Lfunc_end0                      // DW_AT_high_pc
 ; CHECK-NEXT: .b8 1                                // DW_AT_frame_base
 ; CHECK-NEXT: .b8 156
 ; CHECK-NEXT: .b8 95                               // DW_AT_MIPS_linkage_name
 ; CHECK-NEXT: .b8 90
 ; CHECK-NEXT: .b8 53
 ; CHECK-NEXT: .b8 115

 ; CHECK-NEXT: .b8 24                               // DW_AT_call_column
 ; CHECK-NEXT: .b8 42                               // Abbrev [42] 0x26f9:0x18 DW_TAG_inlined_subroutine
 ; CHECK-NEXT: .b32 9785                            // DW_AT_abstract_origin
 ; CHECK-NEXT: .b64 Ltmp2                           // DW_AT_low_pc
 ; CHECK-NEXT: .b64 Ltmp3                           // DW_AT_high_pc
 ; CHECK-NEXT: .b8 12                               // DW_AT_call_file
 ; CHECK-NEXT: .b8 6                                // DW_AT_call_line
 ; CHECK-NEXT: .b8 37                               // DW_AT_call_column
-; CHECK-NEXT: .b8 43                               // Abbrev [43] 0x2711:0x23 DW_TAG_inlined_subroutine
+; CHECK-NEXT: .b8 43                                  // Abbrev [43] 0x2711:0x28 DW_TAG_inlined_subroutine
 ; CHECK-NEXT: .b32 9791                            // DW_AT_abstract_origin
 ; CHECK-NEXT: .b64 Ltmp9                           // DW_AT_low_pc
 ; CHECK-NEXT: .b64 Ltmp10                          // DW_AT_high_pc
 ; CHECK-NEXT: .b8 12                               // DW_AT_call_file
 ; CHECK-NEXT: .b8 8                                // DW_AT_call_line
 ; CHECK-NEXT: .b8 5                                // DW_AT_call_column
 ; CHECK-NEXT: .b8 44                               // Abbrev [44] 0x2729:0x5 DW_TAG_formal_parameter
 ; CHECK-NEXT: .b32 9811                            // DW_AT_abstract_origin
 ; CHECK-NEXT: .b8 44                               // Abbrev [44] 0x272e:0x5 DW_TAG_formal_parameter
 ; CHECK-NEXT: .b32 9820                            // DW_AT_abstract_origin
+; CHECK-NEXT: .b8 44                               // Abbrev [44] 0x2733:0x5 DW_TAG_formal_parameter
+; CHECK-NEXT: .b32 9829                            // DW_AT_abstract_origin
 ; CHECK-NEXT: .b8 0                                // End Of Children Mark
 ; CHECK-NEXT: .b8 0                                // End Of Children Mark
 ; CHECK-NEXT: .b8 0                                // End Of Children Mark
 ; CHECK-NEXT: }
 ; CHECK-NEXT: .section .debug_loc { }
 ; CHECK-NOT: debug_
 
 ; Function Attrs: nounwind readnone