diff --git a/IGC/DebugInfo/DwarfCompileUnit.cpp b/IGC/DebugInfo/DwarfCompileUnit.cpp index 2a31c8ca1034..a08edb4db28e 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 && 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 && 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-sliced-stack-value.cl b/IGC/ocloc_tests/DebugInfo/simd32-sliced-stack-value.cl new file mode 100644 index 000000000000..067937fbad8c --- /dev/null +++ b/IGC/ocloc_tests/DebugInfo/simd32-sliced-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-sliced-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: