[DebugInfo] Only emit stack value at the end of SIMD32 exprs

Artem Gindinson · igcbot · commit d08d9ce9cc91 · 2024-11-29T10:58:25.000+01:00
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.
diff --git 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<DbgDecoder::LiveIntervalsVISA> *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<DbgDecoder::LiveIntervalsVISA> *vars, bool firstHalf) {
+    DbgVariable &var, IGC::DIEBlock *Block, const VISAVariableLocation &loc,
+    const std::vector<DbgDecoder::LiveIntervalsVISA> *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<int32_t>(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<DbgDecoder::LiveIntervalsVISA> *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
@@ -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<DbgDecoder::LiveIntervalsVISA> *,
                          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<DbgDecoder::LiveIntervalsVISA> *);
-  bool buildValidVar(const DbgVariable &, IGC::DIEBlock *,
+  bool buildValidVar(DbgVariable &, IGC::DIEBlock *,
                      const VISAVariableLocation &,
-                     const std::vector<DbgDecoder::LiveIntervalsVISA> *, bool);
+                     const std::vector<DbgDecoder::LiveIntervalsVISA> *,
+                     DbgRegisterType);
 
   // Variables, used in buildGeneral-algorithm:
   bool emitLocation = false;
diff --git 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
@@ -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
@@ -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: <End of list>
+// 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: <End of list>
