From d08d9ce9cc91162a67326e46dfe464ed72ec475e Mon Sep 17 00:00:00 2001 From: Artem Gindinson Date: Fri, 22 Nov 2024 16:05:29 +0000 Subject: [PATCH] [DebugInfo] Only emit stack value at the end of SIMD32 exprs Since our emitter splits SIMD32 programs into SIMD16 subprograms, we should adapt the debug info by only emitting `DW_OP_stack_value` at the end of the source variable's DI Block, i.e. after the merge point that follows the upper SIMD16 register. The approach is to note the information about the split in an `IGC::DbgVariable` instance when generating the upper register variable. Potentially, we could also consider checking for the presence of `DW_OP_skip` when evaluating the DI expression, however it would seem less future-proof in case of future use cases for skips. The change is accompanied by minor in-place refactoring where appropriate. --- IGC/DebugInfo/DwarfCompileUnit.cpp | 44 ++++++--- IGC/DebugInfo/DwarfCompileUnit.hpp | 9 +- IGC/DebugInfo/DwarfDebug.cpp | 19 ++-- IGC/DebugInfo/DwarfDebug.hpp | 29 +++++- .../DebugInfo/simd32-split-stack-value.cl | 92 +++++++++++++++++++ 5 files changed, 163 insertions(+), 30 deletions(-) create mode 100644 IGC/ocloc_tests/DebugInfo/simd32-split-stack-value.cl diff --git a/IGC/DebugInfo/DwarfCompileUnit.cpp b/IGC/DebugInfo/DwarfCompileUnit.cpp index 2a31c8ca1034..cfab6f5a0f21 100644 --- a/IGC/DebugInfo/DwarfCompileUnit.cpp +++ b/IGC/DebugInfo/DwarfCompileUnit.cpp @@ -2257,7 +2257,7 @@ IGC::DIE *CompileUnit::constructVariableDIE(DbgVariable &DV, // Check if variable is described by a DBG_VALUE instruction. const Instruction *pDbgInst = DV.getDbgInst(); - if (!pDbgInst || !DV.isLocationInlined) { + if (!pDbgInst || !DV.currentLocationIsInlined()) { DV.setDIE(VariableDie); LLVM_DEBUG(dbgs() << " done. No dbg.inst assotiated\n"); return VariableDie; @@ -2603,8 +2603,7 @@ bool CompileUnit::buildFpBasedLoc(const DbgVariable &var, IGC::DIEBlock *Block, } bool CompileUnit::buildSlicedLoc( - const DbgVariable &var, IGC::DIEBlock *Block, - const VISAVariableLocation &loc, + DbgVariable &var, IGC::DIEBlock *Block, const VISAVariableLocation &loc, const std::vector *vars) { LLVM_DEBUG(dbgs() << " sliced variable, pushing lane \n"); // DW_OP_push_simd_lane @@ -2628,7 +2627,7 @@ bool CompileUnit::buildSlicedLoc( unsigned int offsetNotTaken = Block->ComputeSizeOnTheFly(Asm); // Emit first register - if (!buildValidVar(var, Block, loc, vars, true)) + if (!buildValidVar(var, Block, loc, vars, DbgRegisterType::FirstHalf)) return false; // Emit second half register @@ -2645,16 +2644,16 @@ bool CompileUnit::buildSlicedLoc( // register in buildValidVar(), which always processes the 1st register only. VISAVariableLocation second_loc(loc); second_loc.SetRegister(loc.GetSecondReg()); - if (!buildValidVar(var, Block, second_loc, vars, false)) + if (!buildValidVar(var, Block, second_loc, vars, DbgRegisterType::SecondHalf)) return false; return true; } bool CompileUnit::buildValidVar( - const DbgVariable &var, IGC::DIEBlock *Block, - const VISAVariableLocation &loc, - const std::vector *vars, bool firstHalf) { + DbgVariable &var, IGC::DIEBlock *Block, const VISAVariableLocation &loc, + const std::vector *vars, + DbgRegisterType regType) { const DbgDecoder::VarInfo *VarInfo = nullptr; const auto *VISAMod = loc.GetVISAModule(); @@ -2671,11 +2670,26 @@ bool CompileUnit::buildValidVar( LLVM_DEBUG(dbgs() << " warning: could not get vISA Variable info\n"); } - if (VarInfo || (vars && vars->size() >= (firstHalf ? 1u : 2u))) { - const auto &lrToUse = - vars ? vars->at(firstHalf ? 0 : 1) : VarInfo->lrs.front(); + const bool isSecondHalf = regType == DbgRegisterType::SecondHalf; + const unsigned NumVarsExpected = isSecondHalf ? 2 : 1; + // TODO: If neither condition is fulfilled, should we do an early + // 'return false' as in "invalid variable"? In that case, we could improve + // the logic in the following way: + // + // DbgDecoder::LiveIntervalsVISA *lrToUse = nullptr; + // if (vars->size() >= NumVarsExpected) + // lrToUse = vars->at(LRIndex); + // else if (VarInfo) + // lrToUse = VarInfo->lrs.front(); + // if (!lrToUse) + // return false; + // /* remaining code from the if block */ + if (VarInfo || vars->size() >= NumVarsExpected) { + const unsigned LRIndex = isSecondHalf ? 1 : 0; + const auto &lrToUse = vars ? vars->at(LRIndex) : VarInfo->lrs.front(); LLVM_DEBUG(dbgs() << " emitting variable location at LR: <"; lrToUse.print(dbgs()); dbgs() << ">\n"); + var.setLocationRegisterType(regType); emitLocation = true; if (lrToUse.isGRF()) { if (loc.IsVectorized() == false) { @@ -2702,7 +2716,7 @@ bool CompileUnit::buildValidVar( SimdOffset < MaxUI16); if (loc.IsRegister()) addSimdLane(Block, var, loc, &lrToUse, (uint16_t)(SimdOffset), - false, !firstHalf); + false, isSecondHalf); } } } else if (lrToUse.isSpill()) { @@ -2738,7 +2752,7 @@ bool CompileUnit::buildValidVar( static_cast(VectorOffset)); addBE_FP(Block); addSimdLane(Block, var, loc, &lrToUse, 0, false, - !firstHalf); // Emit SIMD lane for spill (unpacked) + isSecondHalf); // Emit SIMD lane for spill (unpacked) } } } else { @@ -2751,7 +2765,7 @@ bool CompileUnit::buildValidVar( } IGC::DIEBlock *CompileUnit::buildGeneral( - const DbgVariable &var, const VISAVariableLocation &loc, + DbgVariable &var, const VISAVariableLocation &loc, const std::vector *vars, IGC::DIE *VariableDie) { IGC::DIEBlock *Block = new (DIEValueAllocator) IGC::DIEBlock(); @@ -2788,7 +2802,7 @@ IGC::DIEBlock *CompileUnit::buildGeneral( if (loc.HasLocationSecondReg()) { buildSlicedLoc(var, Block, loc, vars); } else { - buildValidVar(var, Block, loc, vars, true); + buildValidVar(var, Block, loc, vars, DbgRegisterType::Regular); } } diff --git a/IGC/DebugInfo/DwarfCompileUnit.hpp b/IGC/DebugInfo/DwarfCompileUnit.hpp index 8195a61a3b62..03e678355324 100644 --- a/IGC/DebugInfo/DwarfCompileUnit.hpp +++ b/IGC/DebugInfo/DwarfCompileUnit.hpp @@ -524,7 +524,7 @@ class CompileUnit { // buildSLM - Build expression for location described as offset in SLM memory. DIEBlock *buildSLM(const DbgVariable &, const VISAVariableLocation &, IGC::DIE *); - DIEBlock *buildGeneral(const DbgVariable &, const VISAVariableLocation &, + DIEBlock *buildGeneral(DbgVariable &, const VISAVariableLocation &, const std::vector *, IGC::DIE *); @@ -533,12 +533,13 @@ class CompileUnit { const VISAVariableLocation &); bool buildFpBasedLoc(const DbgVariable &, IGC::DIEBlock *, const VISAVariableLocation &); - bool buildSlicedLoc(const DbgVariable &, IGC::DIEBlock *, + bool buildSlicedLoc(DbgVariable &, IGC::DIEBlock *, const VISAVariableLocation &, const std::vector *); - bool buildValidVar(const DbgVariable &, IGC::DIEBlock *, + bool buildValidVar(DbgVariable &, IGC::DIEBlock *, const VISAVariableLocation &, - const std::vector *, bool); + const std::vector *, + DbgRegisterType); // Variables, used in buildGeneral-algorithm: bool emitLocation = false; diff --git a/IGC/DebugInfo/DwarfDebug.cpp b/IGC/DebugInfo/DwarfDebug.cpp index 0f7f97fbcc9a..96de1c976add 100644 --- a/IGC/DebugInfo/DwarfDebug.cpp +++ b/IGC/DebugInfo/DwarfDebug.cpp @@ -211,14 +211,19 @@ void DbgVariable::emitExpression(CompileUnit *CU, IGC::DIEBlock *Block) const { } I->appendToVector(Elements); } - bool isStackValueNeeded = false; - if (currentLocationIsSimpleIndirectValue()) { - // drop OP_deref and don't emit DW_OP_stack_value. + const bool isSimpleIndirect = currentLocationIsSimpleIndirectValue(); + if (isSimpleIndirect) + // drop OP_deref Elements.erase(Elements.begin()); - } else if (!currentLocationIsMemoryAddress() && - !currentLocationIsImplicit() && !currentLocationIsVector()) { - isStackValueNeeded = true; + bool shouldResetStackValue = currentLocationIsImplicit(); + if (shouldResetStackValue && !Elements.empty() && + *Elements.rbegin() == dwarf::DW_OP_stack_value) { + Elements.pop_back(); } + const bool isFirstHalf = this->RegType == DbgRegisterType::FirstHalf; + bool isStackValueNeeded = !isSimpleIndirect && + !currentLocationIsMemoryAddress() && + !currentLocationIsVector() && !isFirstHalf; for (auto elem : Elements) { auto BF = DIEInteger::BestForm(false, elem); @@ -1648,7 +1653,7 @@ void DwarfDebug::collectVariableInfo( (pInst->getMetadata("StorageOffset") || Loc.HasSurface() || Loc.IsSLM()))) { RegVar->setDbgInst(pInst); - RegVar->isLocationInlined = true; + RegVar->setLocationInlined(true); break; } } diff --git a/IGC/DebugInfo/DwarfDebug.hpp b/IGC/DebugInfo/DwarfDebug.hpp index 079e4fc90028..63db6335353f 100644 --- a/IGC/DebugInfo/DwarfDebug.hpp +++ b/IGC/DebugInfo/DwarfDebug.hpp @@ -123,6 +123,15 @@ class DotDebugLocEntry { void setSymbol(llvm::MCSymbol *S) { Symbol = S; } }; +//===----------------------------------------------------------------------===// +/// \brief This enum is used to describe whether a register represents one of +/// the SIMD32 register halves. +enum class DbgRegisterType : uint8_t { + Regular = 0, // Represents all SIMD channels for a source variable, no slice + FirstHalf = 1, // SIMD32 sliced - lower channels + SecondHalf = 2 // SIMD32 sliced - upper channels +}; + //===----------------------------------------------------------------------===// /// \brief This class is used to track local variable information. class DbgVariable { @@ -143,6 +152,12 @@ class DbgVariable { // DBG_VALUE instruction of the variable const llvm::DbgVariableIntrinsic *m_pDbgInst = nullptr; + // isLocationInlined is true when we expect location to be inlined in + // DW_AT_location. + bool isLocationInlined = false; + + DbgRegisterType RegType = DbgRegisterType::Regular; + public: // AbsVar may be NULL. DbgVariable(const llvm::DILocalVariable *V, @@ -196,6 +211,16 @@ class DbgVariable { bool currentLocationIsSimpleIndirectValue() const; bool currentLocationIsVector() const; + bool currentLocationIsInlined() const { return isLocationInlined; } + void setLocationInlined(bool isInlined = true) { + isLocationInlined = isInlined; + } + + DbgRegisterType getLocationRegisterType() const { return RegType; } + void setLocationRegisterType(DbgRegisterType RegType) { + this->RegType = RegType; + } + void emitExpression(CompileUnit *CU, IGC::DIEBlock *Block) const; // Translate tag to proper Dwarf tag. @@ -223,10 +248,6 @@ class DbgVariable { return false; } - // isLocationInlined is true when we expect location to be inlined in - // DW_AT_location. - bool isLocationInlined = false; - bool isBlockByrefVariable() const; llvm::DIType *getType() const; diff --git a/IGC/ocloc_tests/DebugInfo/simd32-split-stack-value.cl b/IGC/ocloc_tests/DebugInfo/simd32-split-stack-value.cl new file mode 100644 index 000000000000..e7927c5b5181 --- /dev/null +++ b/IGC/ocloc_tests/DebugInfo/simd32-split-stack-value.cl @@ -0,0 +1,92 @@ +//========================== begin_copyright_notice ============================ +// +// Copyright (C) 2024 Intel Corporation +// +// SPDX-License-Identifier: MIT +// +//=========================== end_copyright_notice ============================= + +// UNSUPPORTED: sys32 +// REQUIRES: oneapi-readelf, dg2-supported + +// RUN: ocloc compile -file %s -options "-g -igc_opts 'PrintToConsole=1 PrintAfter=EmitPass'" -device dg2 2>&1 | FileCheck %s --check-prefix=CHECK-LLVM +// RUN: ocloc compile -file %s -options "-g -igc_opts 'ElfDumpEnable=1 DumpUseShorterName=0 DebugDumpNamePrefix=%t_dg2_'" -device dg2 +// RUN: oneapi-readelf --debug-dump %t_dg2_OCL_simd32_entry_0001.elf | FileCheck %s --check-prefix=CHECK-DWARF + +// CHECK-LLVM-LABEL: define spir_kernel void @test +// CHECK-LLVM-SAME: ({{.*}} %in,{{.*}} %out,{{.*}} i16 %localIdX{{.*}}) +__attribute__((intel_reqd_sub_group_size(32))) +kernel void test(global int* in, global int* out) { + // COM: The routine instructions for local ID extraction are largely skipped in the checks below + // CHECK-LLVM: %[[LOCAL_ID_X:.+]] = zext i16 %localIdX to i32, !dbg !{{[0-9]+}} + // CHECK-LLVM: %[[LOCAL_ID_TMP_0:.+]] = add i32 %{{.+}}, %localIdX4, !dbg !{{[0-9]+}} + // CHECK-LLVM: %[[LOCAL_ID_TMP_1:.+]] = add i32 %[[LOCAL_ID_TMP_0]], %{{.*}}, !dbg !{{[0-9]+}} + // COM: 'gid' is implicitly marked as a stack value before the emitter + // CHECK-LLVM: call void @llvm.dbg.value(metadata i32 %[[LOCAL_ID_TMP_1]], metadata ![[GID_DI_VAR_MD:[0-9]+]] + // CHECK-LLVM-SAME: metadata !DIExpression(DW_OP_LLVM_convert, 32, DW_ATE_unsigned, DW_OP_LLVM_convert, 64, DW_ATE_unsigned, DW_OP_stack_value)) + size_t gid = get_global_id(0); + // CHECK-LLVM: %[[IN_LOAD:.+]] = call i32 @llvm.genx.GenISA.ldraw.indexed.i32{{.*}}({{.*}}), !dbg !{{.*}} + // CHECK-LLVM: %[[MUL:.+]] = mul nsw i32 %[[IN_LOAD]], 42, !dbg !{{.*}} + // COM: 'mul' is to be marked as a stack value during the emitter phase + // CHECK-LLVM: call void @llvm.dbg.value(metadata i32 %[[MUL]], metadata ![[MUL_DI_VAR_MD:[0-9]+]], metadata !DIExpression()) + int mul = in[gid] * 42; + out[gid] = mul; +} +// CHECK-LLVM-DAG: !{{[0-9]+}} = !{!"sub_group_size", i32 32} +// +// CHECK-LLVM-DAG: ![[GID_DI_VAR_MD]] = !DILocalVariable(name: "gid", {{.+}}, type: ![[SIZE_T_DI_TY_MD:[0-9]+]]) +// CHECK-LLVM-DAG: ![[SIZE_T_DI_TY_MD]] = !DIDerivedType(tag: DW_TAG_typedef, name: "size_t", file: !{{[0-9]+}}, baseType: !{{[0-9]+}}) +// CHECK-LLVM-DAG: ![[MUL_DI_VAR_MD]] = !DILocalVariable(name: "mul", {{.+}}, type: ![[INT_DI_TY_MD:[0-9]+]]) +// CHECK-LLVM-DAG: ![[INT_DI_TY_MD]] = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) + +// CHECK-DWARF: Contents of the .debug_info section: +// COM: Briefly check global layout +// CHECK-DWARF: Abbrev Number: 1 (DW_TAG_compile_unit) +// CHECK-DWARF: DW_AT_name{{ *}}: simd32-split-stack-value.cl +// CHECK-DWARF: DW_AT_INTEL_simd_width{{ *}}: 32 +// CHECK-DWARF: Abbrev Number: 2 (DW_TAG_subprogram) +// CHECK-DWARF: DW_AT_name{{ *}}: test +// CHECK-DWARF: DW_AT_INTEL_simd_width{{ *}}: 32 +// CHECK-DWARF: Abbrev Number: 3 (DW_TAG_formal_parameter) +// CHECK-DWARF: DW_AT_name{{ *}}: in +// CHECK-DWARF: Abbrev Number: 3 (DW_TAG_formal_parameter) +// CHECK-DWARF: DW_AT_name{{ *}}: out +// COM: Relevant variable checks/type captures +// CHECK-DWARF: Abbrev Number: 4 (DW_TAG_variable) +// CHECK-DWARF: DW_AT_name{{ *}}: gid +// CHECK-DWARF: DW_AT_type{{ *}}: <0x[[SIZE_T_TY:[0-9a-f]+]]> +// CHECK-DWARF: DW_AT_location{{ *}}: [[GID_LOC:0]] (location list) +// CHECK-DWARF: Abbrev Number: 4 (DW_TAG_variable) +// CHECK-DWARF: DW_AT_name{{ *}}: mul +// CHECK-DWARF: DW_AT_type{{ *}}: <0x[[INT_TY:[0-9a-f]+]]> +// CHECK-DWARF: DW_AT_location{{ *}}: 0x[[MUL_LOC:[0-9a-f]+]] (location list) +// COM: Type checks +// CHECK-DWARF: <[[INT_TY]]>: Abbrev Number: 6 (DW_TAG_base_type) +// CHECK-DWARF-NEXT: DW_AT_name{{ *}}: int +// CHECK-DWARF-NEXT: DW_AT_encoding{{ *}}: 5{{ *}} (signed) +// CHECK-DWARF: <[[SIZE_T_TY]]>: Abbrev Number: 7 (DW_TAG_typedef) +// CHECK-DWARF-NEXT: DW_AT_type +// CHECK-DWARF-NEXT: DW_AT_name{{ *}}: size_t +// +// CHECK-DWARF: Contents of the .debug_loc section: +// COM: Check SIMD 32 location expressions. We only expect DW_OP_stack_value at the end of each +// expression, never before a DW_OP_skip. +// CHECK-DWARF-NOT: DW_OP_stack_value; DW_OP_skip +// COM: 'gid' source variable (implicit stack value) +// CHECK-DWARF: {{0+}}[[GID_LOC]] {{[0-9a-f]+}} {{[0-9a-f]+}} +// CHECK-DWARF-SAME: (DW_OP_INTEL_push_simd_lane; DW_OP_lit16; DW_OP_ge; DW_OP_bra: [[GID_BR:[0-9]+]]; +// CHECK-DWARF-SAME: DW_OP_INTEL_push_simd_lane; DW_OP_lit3; DW_OP_shr; DW_OP_plus_uconst: {{[0-9]+}}; +// CHECK-DWARF-SAME: DW_OP_INTEL_push_simd_lane; [[GID_MAIN_EXPR:DW_OP_lit7; DW_OP_and; DW_OP_const1u: 32; DW_OP_mul; DW_OP_INTEL_regval_bits: 32; DW_OP_const4u: 4294967295; DW_OP_and]]; +// CHECK-DWARF-SAME: DW_OP_skip: [[GID_BR]]; +// CHECK-DWARF-SAME: DW_OP_INTEL_push_simd_lane; DW_OP_lit16; DW_OP_minus; DW_OP_lit3; DW_OP_shr; DW_OP_plus_uconst: {{[0-9]+}}; +// CHECK-DWARF-SAME: DW_OP_INTEL_push_simd_lane; [[GID_MAIN_EXPR]]; DW_OP_stack_value) +// CHECK-DWARF-NEXT: +// COM: 'mul' source variable (explicitly marked as stack value) +// CHECK-DWARF: {{0+}}[[MUL_LOC]] {{[0-9a-f]+}} {{[0-9a-f]+}} +// CHECK-DWARF-SAME: (DW_OP_INTEL_push_simd_lane; DW_OP_lit16; DW_OP_ge; DW_OP_bra: [[MUL_BR:[0-9]+]]; +// CHECK-DWARF-SAME: DW_OP_INTEL_push_simd_lane; DW_OP_lit3; DW_OP_shr; DW_OP_plus_uconst: {{[0-9]+}}; +// CHECK-DWARF-SAME: DW_OP_INTEL_push_simd_lane; [[MUL_MAIN_EXPR:DW_OP_lit7; DW_OP_and; DW_OP_const1u: 32; DW_OP_mul; DW_OP_INTEL_regval_bits: 32]]; +// CHECK-DWARF-SAME: DW_OP_skip: [[MUL_BR]]; +// CHECK-DWARF-SAME: DW_OP_INTEL_push_simd_lane; DW_OP_lit16; DW_OP_minus; DW_OP_lit3; DW_OP_shr; DW_OP_plus_uconst: {{[0-9]+}}; +// CHECK-DWARF-SAME: DW_OP_INTEL_push_simd_lane; [[MUL_MAIN_EXPR]]; DW_OP_stack_value) +// CHECK-DWARF-NEXT: