Index: llvm/include/llvm/Transforms/Utils/Local.h =================================================================== --- llvm/include/llvm/Transforms/Utils/Local.h +++ llvm/include/llvm/Transforms/Utils/Local.h @@ -352,13 +352,10 @@ DenseMap &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. -/// Returns true if any debug users were updated. -bool salvageDebugInfo(Instruction &I); +/// debug users of \p I by writing the effect of \p I in a DIExpression. If it +/// cannot be salvaged mark it undef. +void 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. Index: llvm/lib/Transforms/InstCombine/InstCombineInternal.h =================================================================== --- llvm/lib/Transforms/InstCombine/InstCombineInternal.h +++ llvm/lib/Transforms/InstCombine/InstCombineInternal.h @@ -721,7 +721,7 @@ 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. Index: llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp =================================================================== --- llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp +++ llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp @@ -88,7 +88,7 @@ Depth, I); if (!NewVal) return false; if (Instruction* OpInst = dyn_cast(U)) - salvageDebugInfoOrMarkUndef(*OpInst); + salvageDebugInfo(*OpInst); replaceUse(U, NewVal); return true; Index: llvm/lib/Transforms/InstCombine/InstructionCombining.cpp =================================================================== --- llvm/lib/Transforms/InstCombine/InstructionCombining.cpp +++ llvm/lib/Transforms/InstCombine/InstructionCombining.cpp @@ -3626,7 +3626,7 @@ if (isInstructionTriviallyDead(Inst, TLI)) { ++NumDeadInst; LLVM_DEBUG(dbgs() << "IC: DCE: " << *Inst << '\n'); - salvageDebugInfoOrMarkUndef(*Inst); + salvageDebugInfo(*Inst); Inst->eraseFromParent(); MadeIRChange = true; continue; Index: llvm/lib/Transforms/Scalar/BDCE.cpp =================================================================== --- llvm/lib/Transforms/Scalar/BDCE.cpp +++ llvm/lib/Transforms/Scalar/BDCE.cpp @@ -102,7 +102,7 @@ (I.getType()->isIntOrIntVectorTy() && DB.getDemandedBits(&I).isNullValue() && wouldInstructionBeTriviallyDead(&I))) { - salvageDebugInfoOrMarkUndef(I); + salvageDebugInfo(I); Worklist.push_back(&I); I.dropAllReferences(); Changed = true; Index: llvm/lib/Transforms/Scalar/DeadStoreElimination.cpp =================================================================== --- llvm/lib/Transforms/Scalar/DeadStoreElimination.cpp +++ llvm/lib/Transforms/Scalar/DeadStoreElimination.cpp @@ -145,7 +145,7 @@ ++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 Index: llvm/lib/Transforms/Scalar/EarlyCSE.cpp =================================================================== --- llvm/lib/Transforms/Scalar/EarlyCSE.cpp +++ llvm/lib/Transforms/Scalar/EarlyCSE.cpp @@ -949,7 +949,7 @@ } salvageKnowledge(&Inst, &AC); - salvageDebugInfoOrMarkUndef(Inst); + salvageDebugInfo(Inst); removeMSSA(Inst); Inst.eraseFromParent(); Changed = true; Index: llvm/lib/Transforms/Scalar/LICM.cpp =================================================================== --- llvm/lib/Transforms/Scalar/LICM.cpp +++ llvm/lib/Transforms/Scalar/LICM.cpp @@ -504,7 +504,7 @@ if (sink(I, LI, DT, CurLoop, SafetyInfo, MSSAU, ORE)) { if (!FreeInLoop) { ++II; - salvageDebugInfoOrMarkUndef(I); + salvageDebugInfo(I); eraseInstruction(I, *SafetyInfo, CurAST, MSSAU); } Changed = true; Index: llvm/lib/Transforms/Scalar/Reassociate.cpp =================================================================== --- llvm/lib/Transforms/Scalar/Reassociate.cpp +++ llvm/lib/Transforms/Scalar/Reassociate.cpp @@ -1900,7 +1900,7 @@ 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(Op)) @@ -1917,7 +1917,7 @@ // Erase the dead instruction. ValueRankMap.erase(I); RedoInsts.remove(I); - llvm::salvageDebugInfoOrMarkUndef(*I); + llvm::salvageDebugInfo(*I); I->eraseFromParent(); // Optimize its operands. SmallPtrSet Visited; // Detect self-referential nodes. Index: llvm/lib/Transforms/Utils/Local.cpp =================================================================== --- llvm/lib/Transforms/Utils/Local.cpp +++ llvm/lib/Transforms/Utils/Local.cpp @@ -1628,7 +1628,7 @@ return MetadataAsValue::get(C, ValueAsMetadata::get(V)); } -bool llvm::salvageDebugInfo(Instruction &I) { +static bool attemptToSalvageDebugInfo(Instruction &I) { SmallVector DbgUsers; findDbgUsers(DbgUsers, &I); if (DbgUsers.empty()) @@ -1637,8 +1637,8 @@ return salvageDebugInfoForDbgValues(I, DbgUsers); } -void llvm::salvageDebugInfoOrMarkUndef(Instruction &I) { - if (!salvageDebugInfo(I)) +void llvm::salvageDebugInfo(Instruction &I) { + if (!attemptToSalvageDebugInfo(I)) replaceDbgUsesWithUndef(&I); } @@ -1822,7 +1822,7 @@ if (!UndefOrSalvage.empty()) { // Try to salvage the remaining debug users. - salvageDebugInfoOrMarkUndef(From); + salvageDebugInfo(From); Changed = true; } Index: llvm/test/DebugInfo/NVPTX/debug-addr-class.ll =================================================================== --- llvm/test/DebugInfo/NVPTX/debug-addr-class.ll +++ llvm/test/DebugInfo/NVPTX/debug-addr-class.ll @@ -168,16 +168,23 @@ ; 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 @@ -289,7 +296,7 @@ ; 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 @@ -312,15 +319,29 @@ ; 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 @@ -330,7 +351,5 @@ ; 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 { } Index: llvm/test/DebugInfo/NVPTX/debug-info.ll =================================================================== --- llvm/test/DebugInfo/NVPTX/debug-info.ll +++ llvm/test/DebugInfo/NVPTX/debug-info.ll @@ -1,6 +1,6 @@ ; 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) { @@ -702,12 +702,12 @@ ; 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 @@ -8306,7 +8306,7 @@ ; 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 @@ -8386,7 +8386,7 @@ ; 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 @@ -8397,6 +8397,8 @@ ; 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