[llvm-branch-commits] [llvm] [SPIRV] Emit NonSemantic DebugDeclare, DebugExpression and DebugOperation (PR #222011)
Manuel Carrasco via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Wed Sep 9 01:16:51 PDT 2026
https://github.com/mgcarrasco updated https://github.com/llvm/llvm-project/pull/222011
>From 1f115a0e2e7aa83f040fb762db8018259879d770 Mon Sep 17 00:00:00 2001
From: Manuel Carrasco <Manuel.Carrasco at amd.com>
Date: Tue, 8 Sep 2026 09:14:20 -0500
Subject: [PATCH 1/2] [SPIRV] Emit NonSemantic DebugDeclare, DebugExpression
and DebugOperation.
Add support for
[DebugDeclare](https://github.khronos.org/SPIRV-Registry/nonsemantic/NonSemantic.Shader.DebugInfo.html#DebugDeclare),
[DebugExpression](https://github.khronos.org/SPIRV-Registry/nonsemantic/NonSemantic.Shader.DebugInfo.html#DebugExpression)
and
[DebugOperation](https://github.khronos.org/SPIRV-Registry/nonsemantic/NonSemantic.Shader.DebugInfo.html#DebugOperation).
Changes:
- Emit DebugDeclare from an indirect DBG_VALUE whose location register is
defined by OpVariable, which is what the spec requires of the Variable
operand. #dbg_declare does not survive as its own opcode in MIR, IRTranslator
lowers it to an indirect DBG_VALUE.
- Skip every other shape instead of emitting an invalid instruction: an access
chain into a field, an OpFunctionParameter for a byval argument, a constant
address, a register left without a def after dead storage was erased, a
variadic #dbg_value, an expression using an unmapped operation, and a variable
whose DebugLocalVariable was not emitted.
- Collect expressions from MIR, so DebugExpression and DebugOperation are
emitted at module scope and shared by every declare using the same
DIExpression.
- Map the nine operations the instruction set defines, and reject an
expression whole when an argument does not fit in 32 bits (as per spec).
- Cache DebugLocalVariable result ids, which emitDebugLocalVariable already
returned but nothing stored.
- Add tests for each emitted and skipped case, for the interaction with
DebugLine and DebugScope, and for the expression operations.
spirv-val and the spec disagree in two places, in opposite directions:
- DebugDeclare's Variable operand: spirv-val accepts an OpFunctionParameter,
while the spec names only OpVariable. This patch follows the spec, so the byval
case is skipped even though spirv-val would take it, see
debug-declare-function-parameter.ll.
- DebugGlobalVariable's Variable operand: spirv-val rejects a DebugExpression
there although the spec permits one, KhronosGroup/SPIRV-Tools#6469, see debug-global-variable-constant-value.ll.
---
.../SPIRV/SPIRVNonSemanticDebugHandler.cpp | 192 +++++++++++++++++-
.../SPIRV/SPIRVNonSemanticDebugHandler.h | 53 ++++-
.../debug-info/debug-declare-access-chain.ll | 36 ++++
.../debug-declare-dbg-value-variadic.ll | 36 ++++
.../debug-info/debug-declare-dead-alloca.ll | 32 +++
.../debug-declare-expression-unsupported.ll | 31 +++
.../debug-declare-expression-xderef.ll | 60 ++++++
.../debug-declare-function-parameter.ll | 39 ++++
.../debug-info/debug-declare-int-storage.ll | 34 ++++
.../debug-info/debug-declare-line-scope.ll | 98 +++++++++
.../debug-declare-module-scope-variable.ll | 45 ++++
.../debug-info/debug-declare-null-storage.ll | 30 +++
.../debug-declare-poison-storage.ll | 28 +++
.../debug-info/debug-declare-skip-type.ll | 33 +++
.../CodeGen/SPIRV/debug-info/debug-declare.ll | 55 +++++
.../debug-info/debug-expression-bit-piece.ll | 40 ++++
.../debug-info/debug-expression-operations.ll | 54 +++++
.../debug-expression-out-of-range.ll | 76 +++++++
.../debug-global-variable-constant-value.ll | 76 +++++++
.../debug-global-variable-init-expr.ll | 40 +++-
.../debug-local-variable-dbg-value.ll | 7 +-
21 files changed, 1074 insertions(+), 21 deletions(-)
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-access-chain.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dbg-value-variadic.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dead-alloca.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-unsupported.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-xderef.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-function-parameter.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-int-storage.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-line-scope.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-null-storage.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-poison-storage.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-skip-type.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-expression-bit-piece.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-expression-operations.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-expression-out-of-range.ll
create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-constant-value.ll
diff --git a/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.cpp b/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.cpp
index af553d75f6c04..8088555404562 100644
--- a/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.cpp
@@ -18,6 +18,7 @@
#include "llvm/CodeGen/AsmPrinter.h"
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/CodeGen/MachineInstr.h"
+#include "llvm/CodeGen/MachineModuleInfo.h"
#include "llvm/IR/DebugInfo.h"
#include "llvm/IR/DebugInfoMetadata.h"
#include "llvm/IR/DebugProgramInstruction.h"
@@ -28,6 +29,7 @@
#include "llvm/MC/MCInst.h"
#include "llvm/MC/MCStreamer.h"
#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/MathExtras.h"
#include "llvm/Support/Path.h"
#include <cassert>
@@ -320,6 +322,8 @@ void SPIRVNonSemanticDebugHandler::beginModule(Module *M) {
UniqueDebugLocations.clear();
GlobalVariableDebugInfoMap.clear();
LocalVariables.clear();
+ DebugLocalVariableRegs.clear();
+ DebugExpressionRegs.clear();
LexicalBlocks.clear();
DebugScopeRegs.clear();
DebugInlinedAtRegs.clear();
@@ -842,10 +846,82 @@ std::optional<MCRegister> SPIRVNonSemanticDebugHandler::mapDISignatureTypeToReg(
return lookupOptReg(DebugScopeRegs, Ty);
}
-// Unimplemented no-op; see emitDebugExpression declaration.
+// NonSemantic.Shader.DebugInfo.100 debug operation encodings
+// (section 4.5, "Debug Operations").
+namespace NonSemanticDebugOp {
+enum : uint32_t {
+ Deref = 0,
+ Plus = 1,
+ Minus = 2,
+ PlusUconst = 3,
+ BitPiece = 4,
+ Swap = 5,
+ Xderef = 6,
+ StackValue = 7,
+ Constu = 8,
+ Fragment = 9
+};
+} // namespace NonSemanticDebugOp
+
+static std::optional<uint32_t> mapDwarfOpToNonSemanticOp(uint64_t DwarfOp) {
+ switch (DwarfOp) {
+ case dwarf::DW_OP_deref:
+ return NonSemanticDebugOp::Deref;
+ case dwarf::DW_OP_plus:
+ return NonSemanticDebugOp::Plus;
+ case dwarf::DW_OP_minus:
+ return NonSemanticDebugOp::Minus;
+ case dwarf::DW_OP_plus_uconst:
+ return NonSemanticDebugOp::PlusUconst;
+ case dwarf::DW_OP_bit_piece:
+ return NonSemanticDebugOp::BitPiece;
+ case dwarf::DW_OP_swap:
+ return NonSemanticDebugOp::Swap;
+ case dwarf::DW_OP_xderef:
+ return NonSemanticDebugOp::Xderef;
+ case dwarf::DW_OP_stack_value:
+ return NonSemanticDebugOp::StackValue;
+ case dwarf::DW_OP_constu:
+ return NonSemanticDebugOp::Constu;
+ case dwarf::DW_OP_LLVM_fragment:
+ return NonSemanticDebugOp::Fragment;
+ default:
+ return std::nullopt;
+ }
+}
+
std::optional<MCRegister> SPIRVNonSemanticDebugHandler::emitDebugExpression(
- const DIExpression *, MCRegister, MCRegister, SPIRV::ModuleAnalysisInfo &) {
- return std::nullopt;
+ const DIExpression *Expr, MCRegister VoidTypeReg, MCRegister I32TypeReg,
+ MCRegister ExtInstSetReg, SPIRV::ModuleAnalysisInfo &MAI) {
+ assert(Expr && "Expr must not be null in emitDebugExpression");
+
+ // Check the whole expression before emitting anything, and bail out if
+ // unsupported. Verify that the operation is supported, and that each argument
+ // is a 32-bit constant (as per spec).
+ for (const DIExpression::ExprOperand &Op : Expr->expr_ops()) {
+ if (!mapDwarfOpToNonSemanticOp(Op.getOp()))
+ return std::nullopt;
+ for (unsigned I = 0, E = Op.getNumArgs(); I != E; ++I)
+ if (!isUInt<32>(Op.getArg(I)))
+ return std::nullopt;
+ }
+
+ SmallVector<MCRegister, 4> OperationRegs;
+ for (const DIExpression::ExprOperand &Op : Expr->expr_ops()) {
+ SmallVector<MCRegister, 3> Operands{emitOpConstantI32(
+ *mapDwarfOpToNonSemanticOp(Op.getOp()), I32TypeReg, MAI)};
+ for (unsigned I = 0, E = Op.getNumArgs(); I != E; ++I)
+ // Operands are truncated to 32 bits but we already checked that they are
+ // in range.
+ Operands.push_back(emitOpConstantI32(static_cast<uint32_t>(Op.getArg(I)),
+ I32TypeReg, MAI));
+ OperationRegs.push_back(
+ emitExtInst(SPIRV::NonSemanticExtInst::DebugOperation, VoidTypeReg,
+ ExtInstSetReg, Operands, MAI));
+ }
+
+ return emitExtInst(SPIRV::NonSemanticExtInst::DebugExpression, VoidTypeReg,
+ ExtInstSetReg, OperationRegs, MAI);
}
std::optional<MCRegister> SPIRVNonSemanticDebugHandler::emitDebugGlobalVariable(
@@ -894,15 +970,18 @@ std::optional<MCRegister> SPIRVNonSemanticDebugHandler::emitDebugGlobalVariable(
MCRegister ColReg = emitOpConstantI32(0, I32TypeReg, MAI);
// Variable: @g OpVariable id when !dbg matches; else a DebugExpression for
- // the GVE init value when no @g exists; else DebugInfoNone.
+ // the GVE init value when no @g exists and the expression is non-empty; else
+ // DebugInfoNone. As per spec, the DebugExpression must contains the constant
+ // value of the variable that was optimized out. An empty expression contains
+ // no value, so we emit DebugInfoNone instead.
MCRegister VariableReg = CachedDebugInfoNoneReg;
if (const GlobalVariable *LLVMGV = Info.LLVMGV) {
MCRegister GVReg = MAI.getGlobalObjReg(LLVMGV);
if (GVReg.isValid())
VariableReg = GVReg;
- } else if (Info.Expr) {
- if (auto ExprReg =
- emitDebugExpression(Info.Expr, VoidTypeReg, ExtInstSetReg, MAI))
+ } else if (Info.Expr && Info.Expr->getNumElements() != 0) {
+ if (auto ExprReg = emitDebugExpression(Info.Expr, VoidTypeReg, I32TypeReg,
+ ExtInstSetReg, MAI))
VariableReg = *ExprReg;
}
@@ -1300,6 +1379,72 @@ void SPIRVNonSemanticDebugHandler::beginInstruction(const MachineInstr *MI) {
emitDebugScopeForInstruction(*Target);
emitDebugLineForInstruction(*Target);
+
+ emitDebugDeclare(MI);
+}
+
+// The register that holds the variable's address in \p MI, or std::nullopt
+// when \p MI is not a declare this backend can describe.
+//
+// The spec requires DebugDeclare's Variable operand to be "the <id> of an
+// OpVariable instruction that defines the local variable". MIR has no
+// DBG_DECLARE, so what this looks for is an indirect DBG_VALUE whose location
+// register an OpVariable defines.
+static std::optional<Register>
+getDebugDeclareStorageReg(const MachineInstr &MI) {
+ // #dbg_declare is an indirect DBG_VALUE in MIR; #dbg_value is normally a
+ // direct one except for the variadic case.
+ if (!MI.isIndirectDebugValue())
+ return std::nullopt;
+
+ // A variadic #dbg_value becomes DBG_VALUE $noreg, 0, ... which is indirect
+ // too, and $noreg is not virtual.
+ Register LocReg = MI.getDebugOperand(0).getReg();
+ if (!LocReg.isVirtual())
+ return std::nullopt;
+
+ // DebugDeclare can only encode the address of an OpVariable.
+ // Other legitimate #dbg_declare cannot be encoded.
+ // Examples: an access chain for a field, an OpFunctionParameter for a byval
+ // argument, or a module-scope constant for a null or fixed address.
+
+ // LocReg may also have no def at all: erasing dead storage leaves the
+ // DBG_VALUE pointing at an undefined register. MachineVerifier permits that
+ // because LiveDebugVariables normally clears it, but this pipeline has no
+ // register allocation, so LiveDebugVariables never runs.
+ const MachineInstr *Def = MI.getMF()->getRegInfo().getUniqueVRegDef(LocReg);
+ if (!Def || Def->getOpcode() != SPIRV::OpVariable)
+ return std::nullopt;
+
+ return LocReg;
+}
+
+void SPIRVNonSemanticDebugHandler::emitDebugDeclare(const MachineInstr *MI) {
+ assert(DebugFunctionDefinitionEmitted &&
+ "DebugFunctionDefinition must be emitted");
+ assert(CurrentMAI && "CurrentMAI must be set");
+
+ std::optional<Register> LocReg = getDebugDeclareStorageReg(*MI);
+ if (!LocReg)
+ return;
+
+ auto VarRegOpt = lookupOptReg(DebugLocalVariableRegs, MI->getDebugVariable());
+ if (!VarRegOpt)
+ return;
+
+ auto ExprRegOpt = lookupOptReg(DebugExpressionRegs, MI->getDebugExpression());
+ if (!ExprRegOpt)
+ return;
+
+ SPIRV::ModuleAnalysisInfo &MAI = *CurrentMAI;
+ MCRegister StorageReg = MAI.getRegisterAlias(MI->getMF(), *LocReg);
+ if (!StorageReg.isValid())
+ return;
+
+ MCRegister VoidTypeReg = getOrEmitOpTypeVoidReg(MAI);
+ MCRegister ExtInstSetReg = MAI.getExtInstSetReg(NSSet);
+ emitExtInst(SPIRV::NonSemanticExtInst::DebugDeclare, VoidTypeReg,
+ ExtInstSetReg, {*VarRegOpt, StorageReg, *ExprRegOpt}, MAI);
}
static bool isMergeInstruction(unsigned Opcode) {
@@ -1519,6 +1664,22 @@ void SPIRVNonSemanticDebugHandler::notifyEntryLabelEmitted(
tryEmitDebugFunctionDefinition(*CurrentMAI);
}
+void SPIRVNonSemanticDebugHandler::collectDebugExpressions(
+ SetVector<const DIExpression *> &Out) const {
+ MachineModuleInfo *ModuleInfo = Asm->MMI;
+ assert(ModuleInfo && "MachineModuleInfo must be set during module output");
+
+ for (const Function &F : *ModuleInfo->getModule()) {
+ const MachineFunction *MF = ModuleInfo->getMachineFunction(F);
+ if (!MF)
+ continue;
+ for (const MachineBasicBlock &MBB : *MF)
+ for (const MachineInstr &MI : MBB)
+ if (MI.isDebugValueLike())
+ Out.insert(MI.getDebugExpression());
+ }
+}
+
void SPIRVNonSemanticDebugHandler::emitNonSemanticGlobalDebugInfo(
SPIRV::ModuleAnalysisInfo &MAI) {
if (GlobalDIEmitted)
@@ -1729,9 +1890,22 @@ void SPIRVNonSemanticDebugHandler::emitNonSemanticGlobalDebugInfo(
}
// Emit DebugLocalVariable after DebugFunction and their lexical blocks so the
- // Parent operand can resolve.
+ // Parent operand can resolve. Record the ids for DebugDeclare.
for (const DILocalVariable *LV : LocalVariables)
- emitDebugLocalVariable(LV, VoidTypeReg, I32TypeReg, ExtInstSetReg, MAI);
+ if (auto LVReg = emitDebugLocalVariable(LV, VoidTypeReg, I32TypeReg,
+ ExtInstSetReg, MAI))
+ DebugLocalVariableRegs[LV] = *LVReg;
+
+ // Opcodes like DebugDeclare are part of the function body, but
+ // DebugExpression cannot. For such opcodes, we collect the expressions
+ // directly from the MIR to avoid inconsistencies with those in the LLVM IR
+ // module.
+ SetVector<const DIExpression *> Expressions;
+ collectDebugExpressions(Expressions);
+ for (const DIExpression *Expr : Expressions)
+ if (auto ExprReg = emitDebugExpression(Expr, VoidTypeReg, I32TypeReg,
+ ExtInstSetReg, MAI))
+ DebugExpressionRegs[Expr] = *ExprReg;
// Emit DebugGlobalVariable for each collected DIGlobalVariable.
for (const auto &[GV, Info] : GlobalVariableDebugInfoMap)
diff --git a/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.h b/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.h
index 2650402d7c796..fbfb7bfce1666 100644
--- a/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.h
+++ b/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.h
@@ -108,6 +108,16 @@ class SPIRVNonSemanticDebugHandler : public DebugHandlerBase {
// records, dbg intrinsics, and subprogram retained nodes.
SetVector<const DILocalVariable *> LocalVariables;
+ // DebugLocalVariable result id per variable that module-scope emission
+ // actually emitted. DebugDeclare needs it for its Local Variable operand; a
+ // variable missing here (skipped type or scope) gets no declare.
+ DenseMap<const DILocalVariable *, MCRegister> DebugLocalVariableRegs;
+
+ // DebugExpression result id per DIExpression that could be lowered. An
+ // expression missing here uses operations with no NonSemantic counterpart,
+ // so declares referencing it are skipped rather than described wrongly.
+ DenseMap<const DIExpression *, MCRegister> DebugExpressionRegs;
+
// Distinct DILexicalBlock and DINamespace scopes, parent-before-child
// order, collected in beginModule() for DebugLexicalBlock emission.
SetVector<const DIScope *> LexicalBlocks;
@@ -394,14 +404,51 @@ class SPIRVNonSemanticDebugHandler : public DebugHandlerBase {
MCRegister VoidTypeReg, MCRegister I32TypeReg, MCRegister ExtInstSetReg,
SPIRV::ModuleAnalysisInfo &MAI);
- /// Emit \c DebugExpression for \p Expr. Unimplemented: defined as a no-op
- /// (\returns \c std::nullopt, emits nothing) so \c emitDebugGlobalVariable
- /// can complete Variable-operand resolution for the opcodes we support today.
+ /// Collect the \c DIExpression of every debug value in the module
+ /// (\c DBG_VALUE, \c DBG_VALUE_LIST, \c DBG_INSTR_REF), in MIR order.
+ ///
+ /// Reads MIR rather than IR because only MIR shows which debug values
+ /// survived codegen and in what form, and because an expression synthesized
+ /// during lowering never appears in the IR at all. Must be called from
+ /// module-scope emission, which is where the resulting \c DebugExpression
+ /// instructions have to be emitted; every \c MachineFunction is still
+ /// reachable at that point through \c MachineModuleInfo.
+ ///
+ /// Deliberately independent of what the consumers can currently emit, so
+ /// that adding an instruction that needs an expression (\c DebugValue) needs
+ /// no change here. The cost is a \c DebugExpression that nothing references
+ /// yet, for a debug value no instruction is emitted for.
+ void collectDebugExpressions(SetVector<const DIExpression *> &Out) const;
+
+ /// Emit one \c DebugOperation per element of \p Expr followed by the
+ /// \c DebugExpression that lists them. An empty \p Expr yields a
+ /// \c DebugExpression with no operands, which is what a plain
+ /// \c !DIExpression() means.
+ ///
+ /// Must be called from module-scope emission only: \c DebugExpression and
+ /// \c DebugOperation are not in the spec's list of instructions allowed
+ /// inside a function, and forward references were removed in Rev 2.
+ ///
+ /// \returns The result id register on success. Returns \c std::nullopt and
+ /// emits nothing if any element has no NonSemantic counterpart, or carries an
+ /// argument too large for the 32-bit \c OpConstant operands this set
+ /// requires.
std::optional<MCRegister> emitDebugExpression(const DIExpression *Expr,
MCRegister VoidTypeReg,
+ MCRegister I32TypeReg,
MCRegister ExtInstSetReg,
SPIRV::ModuleAnalysisInfo &MAI);
+ /// Emit \c DebugDeclare for \p MI when it is an indirect \c DBG_VALUE whose
+ /// location register is defined by \c OpVariable, which is the shape
+ /// \c IRTranslator gives a \c #dbg_declare on storage the backend kept.
+ ///
+ /// Emits nothing when \p MI is not such a declare, when the variable has no
+ /// \c DebugLocalVariable, when the expression was not lowered, or when the
+ /// storage is anything other than an \c OpVariable (an access chain, a
+ /// constant, a function parameter, or a dead alloca with no def at all).
+ void emitDebugDeclare(const MachineInstr *MI);
+
/// Emit \c DebugTypeVector for the vector composite type \p VT.
///
/// \returns The result id register on success. Returns \c std::nullopt and
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-access-chain.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-access-chain.ll
new file mode 100644
index 0000000000000..9c051a6ac9622
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-access-chain.ll
@@ -0,0 +1,36 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A declare whose address is a field of an aggregate. The location register is
+; defined by OpInBoundsPtrAccessChain, not OpVariable, so the declare is
+; dropped.
+; clang seems to emit dbg_declare on the aggregate's own alloca but not in the GEP.
+
+; CHECK: OpExtInst {{.*}} DebugLocalVariable
+
+target triple = "spirv64-unknown-unknown"
+
+%struct.S = type { i32, i32 }
+
+define spir_func void @f() !dbg !5 {
+entry:
+ %s = alloca %struct.S, align 4
+ %b = getelementptr inbounds %struct.S, ptr %s, i32 0, i32 1
+ #dbg_declare(ptr %b, !9, !DIExpression(), !10)
+ store i32 1, ptr %b, align 4, !dbg !10
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-access-chain.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "b", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dbg-value-variadic.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dbg-value-variadic.ll
new file mode 100644
index 0000000000000..ee22e024a2397
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dbg-value-variadic.ll
@@ -0,0 +1,36 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare --implicit-check-not=DebugExpression
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A variadic #dbg_value. IRTranslator cannot lower the DIArgList.
+; It emits DBG_VALUE $noreg, 0, which is an indirect DBG_VALUE just like a
+; declare is.
+; The test checks that it must not become a DebugDeclare.
+
+; Clang seem to emit no DIArgList. The optimizer does,
+; when it deletes a dead binary operation and rewrites the variable as an
+; expression over the two operands.
+
+; CHECK: OpExtInst {{.*}} DebugLocalVariable
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func i32 @sum(i32 %a, i32 %b) !dbg !5 {
+entry:
+ %add = add nsw i32 %a, %b, !dbg !11
+ #dbg_value(!DIArgList(i32 %a, i32 %b), !9, !DIExpression(DW_OP_LLVM_arg, 0, DW_OP_LLVM_arg, 1, DW_OP_plus, DW_OP_stack_value), !11)
+ ret i32 %add, !dbg !11
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-dbg-value-variadic.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{!7, !7, !7}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "sum", linkageName: "sum", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "total", scope: !5, file: !1, line: 2, type: !7)
+!11 = !DILocation(line: 2, column: 7, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dead-alloca.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dead-alloca.ll
new file mode 100644
index 0000000000000..0c7cde2bcaeb5
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dead-alloca.ll
@@ -0,0 +1,32 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; An alloca whose only use is the declare, which is what clang emits at -O0 for
+; an unused local. Nothing keeps the alloca alive, so no OpVariable is left to
+; point at and the declare is dropped. The DebugLocalVariable still describes
+; the variable.
+
+; CHECK: OpExtInst {{.*}} DebugLocalVariable
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ %unused = alloca i32, align 4
+ #dbg_declare(ptr %unused, !9, !DIExpression(), !10)
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-dead-alloca.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "unused", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-unsupported.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-unsupported.ll
new file mode 100644
index 0000000000000..a816590c2604b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-unsupported.ll
@@ -0,0 +1,31 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare --implicit-check-not=DebugExpression --implicit-check-not=DebugOperation
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A declare whose expression contains an operation with no counterpart in
+; NonSemantic.Shader.DebugInfo.100. No DebugExpression is built for it.
+
+; CHECK: OpExtInst {{.*}} DebugLocalVariable
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ %x = alloca i32, align 4
+ #dbg_declare(ptr %x, !9, !DIExpression(DW_OP_LLVM_convert, 32, DW_ATE_signed), !10)
+ store i32 1, ptr %x, align 4, !dbg !10
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-expression-unsupported.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-xderef.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-xderef.ll
new file mode 100644
index 0000000000000..d00075a58bae3
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-xderef.ll
@@ -0,0 +1,60 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; The address-space expression clang puts on every declare for a SPIR-V target.
+
+; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
+; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
+; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0
+; CHECK-DAG: [[X:%[0-9]+]] = OpString "x"
+; CHECK-DAG: [[Y:%[0-9]+]] = OpString "y"
+; The trailing anchors keep e.g. [[C8]] from binding to "OpConstant %3 80".
+; CHECK-DAG: [[C0:%[0-9]+]] = OpConstant [[I32T]] 0{{ *$}}
+; CHECK-DAG: [[C5:%[0-9]+]] = OpConstant [[I32T]] 5{{ *$}}
+; CHECK-DAG: [[C6:%[0-9]+]] = OpConstant [[I32T]] 6{{ *$}}
+; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}}
+; CHECK-DAG: [[XVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[X]]
+; CHECK-DAG: [[YVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[Y]]
+
+; CHECK-DAG: [[CONSTU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C0]]{{ *$}}
+; CHECK-DAG: [[SWAP:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C5]]{{ *$}}
+; CHECK-DAG: [[XDEREF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C6]]{{ *$}}
+; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[CONSTU]] [[SWAP]] [[XDEREF]]{{ *$}}
+
+; CHECK: [[XADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function
+; CHECK: [[YADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[XVAR]] [[XADDR]] [[EXPR]]
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[YVAR]] [[YADDR]] [[EXPR]]
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func i32 @sum(i32 noundef %x) !dbg !5 {
+entry:
+ %x.addr = alloca i32, align 4
+ %y = alloca i32, align 4
+ store i32 %x, ptr %x.addr, align 4
+ #dbg_declare(ptr %x.addr, !9, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !11)
+ #dbg_declare(ptr %y, !10, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !12)
+ %0 = load i32, ptr %x.addr, align 4, !dbg !12
+ store i32 %0, ptr %y, align 4, !dbg !12
+ %1 = load i32, ptr %y, align 4, !dbg !12
+ ret i32 %1, !dbg !12
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-expression-xderef.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+
+!4 = !DISubroutineType(types: !6)
+!6 = !{!7, !7}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+
+!5 = distinct !DISubprogram(name: "sum", linkageName: "sum", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "x", arg: 1, scope: !5, file: !1, line: 1, type: !7)
+!10 = !DILocalVariable(name: "y", scope: !5, file: !1, line: 2, type: !7)
+!11 = !DILocation(line: 1, column: 13, scope: !5)
+!12 = !DILocation(line: 2, column: 7, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-function-parameter.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-function-parameter.ll
new file mode 100644
index 0000000000000..1de2ea6f1947b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-function-parameter.ll
@@ -0,0 +1,39 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A declare whose address is an incoming argument rather than an alloca, which
+; is the shape clang produces for a byval parameter on this triple. The
+; location register is defined by OpFunctionParameter, so the declare is
+; dropped as per spec.
+
+; At the time of writing, spirv-val accepts it, since its rule allows OpVariable or
+; OpFunctionParameter, but the spec restricts the Variable operand to
+; OpVariable.
+;
+; spirv64-amd-amdhsa never reaches this shape: its ABI passes the struct byref,
+; so the declare lands on the callee's own copy, an alloca.
+
+; CHECK: OpExtInst {{.*}} DebugLocalVariable
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f(ptr %p) !dbg !5 {
+entry:
+ #dbg_declare(ptr %p, !9, !DIExpression(), !10)
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-function-parameter.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null, !8}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!8 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !7, size: 64, dwarfAddressSpace: 4)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "p", arg: 1, scope: !5, file: !1, line: 1, type: !8)
+!10 = !DILocation(line: 1, column: 20, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-int-storage.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-int-storage.ll
new file mode 100644
index 0000000000000..5dd2010ccfb5b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-int-storage.ll
@@ -0,0 +1,34 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A declare on a fixed numeric address, which the IR verifier accepts. The
+; address is a constant rather than an OpVariable, and nothing defines the
+; location register inside the function, so the declare is dropped.
+;
+; Clang seems not to emit declares on constants. The verifier permits an
+; integer location only because deleting an inttoptr rewrites the declare onto
+; the integer itself.
+
+; CHECK: OpExtInst {{.*}} DebugLocalVariable
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ #dbg_declare(ptr inttoptr (i64 4096 to ptr), !9, !DIExpression(), !10)
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-int-storage.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "mapped", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-line-scope.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-line-scope.ll
new file mode 100644
index 0000000000000..eb03ddee8869f
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-line-scope.ll
@@ -0,0 +1,98 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A declare has a location of its own, so it takes part in DebugLine and
+; DebugScope tracking like a real instruction. The body interleaves declares
+; and instructions whose locations disagree on purpose:
+;
+; #dbg_declare(ptr %x, !9, !DIExpression(), !20) ; !20 is line 20
+; store i32 1, ptr %x, align 4, !dbg !21 ; !21 is line 5
+;
+
+; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
+; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
+; CHECK-DAG: [[I32:%[0-9]+]] = OpTypeInt 32 0
+; CHECK-DAG: [[PATH:%[0-9]+]] = OpString "{{[/\\]}}src{{[/\\]}}debug-declare-line-scope.c"
+; CHECK-DAG: [[X:%[0-9]+]] = OpString "x"
+; CHECK-DAG: [[Y:%[0-9]+]] = OpString "y"
+; CHECK-DAG: [[DS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugSource [[PATH]]
+; CHECK-DAG: [[DF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugFunction {{.*}}
+; CHECK-DAG: [[LB:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLexicalBlock [[DS]] {{.*}} [[DF]]
+; CHECK-DAG: [[XVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[X]]
+; CHECK-DAG: [[YVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[Y]]
+; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression{{ *$}}
+; CHECK-DAG: [[V3:%[0-9]+]] = OpConstant [[I32]] 3{{$}}
+; CHECK-DAG: [[V4:%[0-9]+]] = OpConstant [[I32]] 4{{$}}
+; CHECK-DAG: [[V5:%[0-9]+]] = OpConstant [[I32]] 5{{$}}
+; CHECK-DAG: [[V6:%[0-9]+]] = OpConstant [[I32]] 6{{$}}
+; CHECK-DAG: [[V7:%[0-9]+]] = OpConstant [[I32]] 7{{$}}
+; CHECK-DAG: [[V8:%[0-9]+]] = OpConstant [[I32]] 8{{$}}
+; CHECK-DAG: [[V9:%[0-9]+]] = OpConstant [[I32]] 9{{$}}
+; CHECK-DAG: [[V10:%[0-9]+]] = OpConstant [[I32]] 10{{$}}
+; CHECK-DAG: [[V20:%[0-9]+]] = OpConstant [[I32]] 20{{$}}
+; CHECK-DAG: [[V31:%[0-9]+]] = OpConstant [[I32]] 31{{$}}
+
+; CHECK: [[FN:%[0-9]+]] = OpFunction
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[XADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function
+; CHECK-NEXT: [[YADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugFunctionDefinition [[DF]] [[FN]]
+
+; The declare's own location, line 20, not the line 5 of the store below it.
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugScope [[DF]]
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugLine [[DS]] [[V20]] [[V20]] [[V7]] [[V8]]
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[XVAR]] [[XADDR]] [[EXPR]]
+; CHECK-NEXT: ;DEBUG_VALUE:
+
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugLine [[DS]] [[V5]] [[V5]] [[V3]] [[V4]]
+; CHECK-NEXT: OpStore [[XADDR]]
+
+; A declare in a lexical block moves the scope, on the declare alone.
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugScope [[LB]]
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugLine [[DS]] [[V31]] [[V31]] [[V9]] [[V10]]
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[YVAR]] [[YADDR]] [[EXPR]]
+; CHECK-NEXT: ;DEBUG_VALUE:
+
+; The unlocated store drops both, rather than keeping the declare's line.
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugNoScope
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugNoLine
+; CHECK-NEXT: OpStore [[YADDR]]
+
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugScope [[DF]]
+; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugLine [[DS]] [[V6]] [[V6]] [[V3]] [[V4]]
+; CHECK-NEXT: OpReturn
+; CHECK-NEXT: OpFunctionEnd
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ %x = alloca i32, align 4
+ %y = alloca i32, align 4
+ #dbg_declare(ptr %x, !9, !DIExpression(), !20) ; function scope, line 20
+ store i32 1, ptr %x, align 4, !dbg !21 ; function scope, line 5
+ #dbg_declare(ptr %y, !10, !DIExpression(), !22) ; lexical block, line 31
+ store i32 2, ptr %y, align 4 ; no debug location
+ ret void, !dbg !23 ; function scope, line 6
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-line-scope.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!8 = distinct !DILexicalBlock(scope: !5, file: !1, line: 30, column: 3)
+!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 20, type: !7)
+!10 = !DILocalVariable(name: "y", scope: !8, file: !1, line: 31, type: !7)
+!20 = !DILocation(line: 20, column: 7, scope: !5)
+!21 = !DILocation(line: 5, column: 3, scope: !5)
+!22 = !DILocation(line: 31, column: 9, scope: !8)
+!23 = !DILocation(line: 6, column: 3, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll
new file mode 100644
index 0000000000000..710f3ef0221f5
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll
@@ -0,0 +1,45 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A #dbg_declare whose address is a module-scope global rather than an alloca.
+; The G_GLOBAL_VALUE is selected to the Workgroup OpVariable, so the storage
+; operand resolves and the declare is emitted from inside the function.
+;
+; This is a synthetic case, clang doesn't seem to emit this.
+; clang describes an OpenCL __local or a HIP __shared__ variable with a
+; DIGlobalVariable, which takes the DebugGlobalVariable path instead.
+
+; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
+; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
+; CHECK-DAG: [[NAME:%[0-9]+]] = OpString "shared"
+; CHECK-DAG: [[STORAGE:%[0-9]+]] = OpVariable {{%[0-9]+}} Workgroup
+; CHECK-DAG: [[VAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[NAME]]
+; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression{{ *$}}
+
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugFunctionDefinition
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[VAR]] [[STORAGE]] [[EXPR]]
+
+target triple = "spirv64-unknown-unknown"
+
+ at shared = internal addrspace(3) global i32 undef, align 4
+
+define spir_func void @k() !dbg !5 {
+entry:
+ #dbg_declare(ptr addrspace(3) @shared, !9, !DIExpression(), !10)
+ store i32 1, ptr addrspace(3) @shared, align 4, !dbg !10
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-module-scope-variable.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "k", linkageName: "k", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "shared", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 3, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-null-storage.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-null-storage.ll
new file mode 100644
index 0000000000000..615b35e98408a
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-null-storage.ll
@@ -0,0 +1,30 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A declare on a null address. The pointer is a module-scope OpConstantNull, so
+; nothing defines the location register inside the function and the declare is
+; dropped. The register still has a SPIR-V id.
+
+; CHECK: OpExtInst {{.*}} DebugLocalVariable
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ #dbg_declare(ptr null, !9, !DIExpression(), !10)
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-null-storage.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-poison-storage.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-poison-storage.ll
new file mode 100644
index 0000000000000..c5961ec21a697
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-poison-storage.ll
@@ -0,0 +1,28 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare --implicit-check-not=DebugExpression
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A declare on poison, which is how a dropped variable location is treated.
+
+; CHECK: OpExtInst {{.*}} DebugLocalVariable
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ #dbg_declare(ptr poison, !9, !DIExpression(), !10)
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-poison-storage.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "dropped", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-skip-type.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-skip-type.ll
new file mode 100644
index 0000000000000..bee90184131ef
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-skip-type.ll
@@ -0,0 +1,33 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugLocalVariable --implicit-check-not=DebugDeclare
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A declare on live storage, so the OpVariable is there, but the variable's
+; type has no DWARF address space and its DebugTypePointer is skipped.
+
+; CHECK: OpVariable {{%[0-9]+}} Function
+; CHECK: OpExtInst {{.*}} DebugFunctionDefinition
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ %p = alloca ptr, align 8
+ #dbg_declare(ptr %p, !9, !DIExpression(), !10)
+ store ptr null, ptr %p, align 8, !dbg !10
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare-skip-type.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!8 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !7, size: 64)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "p", scope: !5, file: !1, line: 2, type: !8)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare.ll
new file mode 100644
index 0000000000000..fd8fa157c2466
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare.ll
@@ -0,0 +1,55 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; A #dbg_declare on an alloca that survives to MIR: the parameter copy and a
+; local. Both use an empty DIExpression, so a single DebugExpression with no
+; operations is shared by both declares.
+
+; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
+; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
+; CHECK-DAG: [[VALUE:%[0-9]+]] = OpString "value"
+; CHECK-DAG: [[RESULT:%[0-9]+]] = OpString "result"
+; CHECK-DAG: [[VALUEVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[VALUE]]
+; CHECK-DAG: [[RESULTVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[RESULT]]
+; An empty DIExpression lowers to a DebugExpression with no operands.
+; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression{{ *$}}
+
+; CHECK: [[VALUEADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function
+; CHECK: [[RESULTADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugFunctionDefinition
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[VALUEVAR]] [[VALUEADDR]] [[EXPR]]
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[RESULTVAR]] [[RESULTADDR]] [[EXPR]]
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func i32 @add_one(i32 %value) !dbg !5 {
+entry:
+ %value.addr = alloca i32, align 4
+ %result = alloca i32, align 4
+ store i32 %value, ptr %value.addr, align 4
+ #dbg_declare(ptr %value.addr, !9, !DIExpression(), !11)
+ %0 = load i32, ptr %value.addr, align 4, !dbg !11
+ %add = add nsw i32 %0, %0, !dbg !11
+ store i32 %add, ptr %result, align 4, !dbg !11
+ #dbg_declare(ptr %result, !10, !DIExpression(), !12)
+ %1 = load i32, ptr %result, align 4, !dbg !12
+ ret i32 %1, !dbg !12
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-declare.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+
+!4 = !DISubroutineType(types: !6)
+!6 = !{!7, !7}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+
+!5 = distinct !DISubprogram(name: "add_one", linkageName: "add_one", scope: !1, file: !1, line: 7, type: !4, scopeLine: 7, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "value", arg: 1, scope: !5, file: !1, line: 7, type: !7)
+!10 = !DILocalVariable(name: "result", scope: !5, file: !1, line: 11, type: !7)
+!11 = !DILocation(line: 7, column: 20, scope: !5)
+!12 = !DILocation(line: 11, column: 30, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-bit-piece.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-bit-piece.ll
new file mode 100644
index 0000000000000..cfda50da35b0c
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-bit-piece.ll
@@ -0,0 +1,40 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o /dev/null 2>&1 | FileCheck %s --check-prefix=VERIFY
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=Debug
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; DW_OP_bit_piece is the DWARF counterpart of NonSemantic.Shader.DebugInfo.100
+; BitPiece (encoding 4). The IR verifier currently rejects it.
+; If rejected, the debug information is dropped, and llc continues the execution.
+
+; Future implementation must be careful with the following:
+; DW_OP_bit_piece is (size, offset); NSDI BitPiece is (offset, size).
+
+; VERIFY: invalid expression
+; VERIFY: !DIExpression(157, 32, 8)
+; VERIFY: warning: ignoring invalid debug info
+
+; CHECK: OpFunction
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ %x = alloca i32, align 4
+ #dbg_declare(ptr %x, !9, !DIExpression(DW_OP_bit_piece, 32, 8), !10)
+ store i32 1, ptr %x, align 4, !dbg !10
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-expression-bit-piece.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-operations.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-operations.ll
new file mode 100644
index 0000000000000..84e5c766afc41
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-operations.ll
@@ -0,0 +1,54 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; One expression using the operations no other test reaches: Deref (0),
+; Plus (1), Minus (2), PlusUconst (3), StackValue (7) and Fragment (9).
+
+; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
+; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
+; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0
+; CHECK-DAG: [[C0:%[0-9]+]] = OpConstant [[I32T]] 0{{ *$}}
+; CHECK-DAG: [[C1:%[0-9]+]] = OpConstant [[I32T]] 1{{ *$}}
+; CHECK-DAG: [[C2:%[0-9]+]] = OpConstant [[I32T]] 2{{ *$}}
+; CHECK-DAG: [[C3:%[0-9]+]] = OpConstant [[I32T]] 3{{ *$}}
+; CHECK-DAG: [[C4:%[0-9]+]] = OpConstant [[I32T]] 4{{ *$}}
+; CHECK-DAG: [[C7:%[0-9]+]] = OpConstant [[I32T]] 7{{ *$}}
+; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}}
+; CHECK-DAG: [[C9:%[0-9]+]] = OpConstant [[I32T]] 9{{ *$}}
+
+; CHECK-DAG: [[DEREF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C0]]{{ *$}}
+; CHECK-DAG: [[PLUSU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C3]] [[C4]]{{ *$}}
+; CHECK-DAG: [[CONSTU8:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C8]]{{ *$}}
+; CHECK-DAG: [[MINUS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C2]]{{ *$}}
+; CHECK-DAG: [[CONSTU2:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C2]]{{ *$}}
+; CHECK-DAG: [[PLUS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C1]]{{ *$}}
+; CHECK-DAG: [[STACK:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C7]]{{ *$}}
+; CHECK-DAG: [[FRAG:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C9]] [[C0]] [[C8]]{{ *$}}
+; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[DEREF]] [[PLUSU]] [[CONSTU8]] [[MINUS]] [[CONSTU2]] [[PLUS]] [[STACK]] [[FRAG]]{{ *$}}
+
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare {{%[0-9]+}} {{%[0-9]+}} [[EXPR]]
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !5 {
+entry:
+ %x = alloca i32, align 4
+ #dbg_declare(ptr %x, !9, !DIExpression(DW_OP_deref, DW_OP_plus_uconst, 4, DW_OP_constu, 8, DW_OP_minus, DW_OP_constu, 2, DW_OP_plus, DW_OP_stack_value, DW_OP_LLVM_fragment, 0, 8), !10)
+ ; 99 keeps this constant clear of the operation encodings captured above.
+ store i32 99, ptr %x, align 4, !dbg !10
+ ret void, !dbg !10
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "debug-expression-operations.c", directory: "/src")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !DISubroutineType(types: !6)
+!6 = !{null}
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 2, type: !7)
+!10 = !DILocation(line: 3, column: 1, scope: !5)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-out-of-range.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-out-of-range.ll
new file mode 100644
index 0000000000000..108b1f7338084
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-out-of-range.ll
@@ -0,0 +1,76 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --check-prefix=ONE --implicit-check-not=DebugOperation --implicit-check-not=DebugExpression --implicit-check-not=DebugDeclare
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; NonSemantic sets take no literals, so every DIExpression argument becomes a
+; 32-bit OpConstant. An argument that does not fit is dropped.
+
+; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
+; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
+; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0
+; CHECK-DAG: [[OK:%[0-9]+]] = OpString "ok"
+; CHECK-DAG: [[BIG:%[0-9]+]] = OpString "big"
+; CHECK-DAG: [[GBIG:%[0-9]+]] = OpString "gbig"
+; The trailing anchors keep e.g. [[C4]] from binding to "OpConstant %3 40".
+; CHECK-DAG: [[C0:%[0-9]+]] = OpConstant [[I32T]] 0{{ *$}}
+; CHECK-DAG: [[C4:%[0-9]+]] = OpConstant [[I32T]] 4{{ *$}}
+; CHECK-DAG: [[C5:%[0-9]+]] = OpConstant [[I32T]] 5{{ *$}}
+; CHECK-DAG: [[C6:%[0-9]+]] = OpConstant [[I32T]] 6{{ *$}}
+; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}}
+; CHECK-DAG: [[NONE:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugInfoNone
+
+; Both variables are collected and emitted; only the expressions differ.
+; CHECK-DAG: [[OKVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[OK]]
+; CHECK-DAG: [[BIGVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[BIG]]
+
+; CHECK-DAG: [[CONSTU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C4]]{{ *$}}
+; CHECK-DAG: [[SWAP:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C5]]{{ *$}}
+; CHECK-DAG: [[XDEREF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C6]]{{ *$}}
+; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[CONSTU]] [[SWAP]] [[XDEREF]]{{ *$}}
+
+; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[GBIG]] {{.*}} [[GBIG]] [[NONE]]
+
+; Both allocas still become OpVariable, in declaration order; only the declare
+; for the first one can be emitted.
+; CHECK: [[OKADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function
+; CHECK: [[BIGADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[OKVAR]] [[OKADDR]] [[EXPR]]
+
+; ONE: OpExtInst {{.*}} DebugOperation
+; ONE: OpExtInst {{.*}} DebugOperation
+; ONE: OpExtInst {{.*}} DebugOperation
+; ONE: OpExtInst {{.*}} DebugExpression
+; ONE: OpExtInst {{.*}} DebugDeclare
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func i32 @f(i32 noundef %x) !dbg !10 {
+entry:
+ %ok = alloca i32, align 4
+ %big = alloca i32, align 4
+ store i32 %x, ptr %ok, align 4
+ #dbg_declare(ptr %ok, !11, !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef), !13)
+ #dbg_declare(ptr %big, !12, !DIExpression(DW_OP_constu, 4294967296, DW_OP_swap, DW_OP_xderef), !13)
+ %0 = load i32, ptr %ok, align 4, !dbg !13
+ store i32 %0, ptr %big, align 4, !dbg !13
+ %1 = load i32, ptr %big, align 4, !dbg !13
+ ret i32 %1, !dbg !13
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!1, !2}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !3, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, globals: !4, splitDebugInlining: false, nameTableKind: None)
+!1 = !{i32 7, !"Dwarf Version", i32 5}
+!2 = !{i32 2, !"Debug Info Version", i32 3}
+!3 = !DIFile(filename: "t.c", directory: "/src")
+!4 = !{!5}
+!5 = !DIGlobalVariableExpression(var: !6, expr: !DIExpression(DW_OP_constu, 4294967296, DW_OP_swap, DW_OP_xderef))
+!6 = distinct !DIGlobalVariable(name: "gbig", linkageName: "gbig", scope: !0, file: !3, line: 1, type: !7, isLocal: false, isDefinition: true)
+!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!8 = !DISubroutineType(cc: DW_CC_LLVM_SpirFunction, types: !9)
+!9 = !{!7, !7}
+!10 = distinct !DISubprogram(name: "f", scope: !3, file: !3, line: 3, type: !8, scopeLine: 3, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!11 = !DILocalVariable(name: "ok", scope: !10, file: !3, line: 4, type: !7)
+!12 = !DILocalVariable(name: "big", scope: !10, file: !3, line: 5, type: !7)
+!13 = !DILocation(line: 4, column: 7, scope: !10)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-constant-value.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-constant-value.ll
new file mode 100644
index 0000000000000..b5f73c39bedb4
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-constant-value.ll
@@ -0,0 +1,76 @@
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --check-prefix=ONE --implicit-check-not=DebugOperation --implicit-check-not=DebugExpression
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | not spirv-val 2>&1 | FileCheck %s --check-prefix=VAL %}
+
+; This HIP code compiles to the metadata below with no optimisations as it is the case for amdgcnspirv:
+;
+; constexpr unsigned long long Hash = 0xff51afd7ed558ccdULL;
+; constexpr double Scale = 1.5;
+; constexpr unsigned Small = 42u;
+;
+
+; DebugOperation requires int32 operands and DI for Hash and Scale cannot be encoded.
+
+; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
+; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
+; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0
+; CHECK-DAG: [[HASH:%[0-9]+]] = OpString "Hash"
+; CHECK-DAG: [[SCALE:%[0-9]+]] = OpString "Scale"
+; CHECK-DAG: [[SMALL:%[0-9]+]] = OpString "Small"
+; The trailing anchors keep e.g. [[C7]] from binding to "OpConstant %3 72".
+; CHECK-DAG: [[C7:%[0-9]+]] = OpConstant [[I32T]] 7{{ *$}}
+; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}}
+; CHECK-DAG: [[C42:%[0-9]+]] = OpConstant [[I32T]] 42{{ *$}}
+; CHECK-DAG: [[NONE:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugInfoNone
+
+; CHECK-DAG: [[CONSTU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C42]]{{ *$}}
+; CHECK-DAG: [[SV:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C7]]{{ *$}}
+; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[CONSTU]] [[SV]]{{ *$}}
+; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[SMALL]] {{.*}} [[SMALL]] [[EXPR]]
+
+; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[HASH]] {{.*}} [[HASH]] [[NONE]]
+; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[SCALE]] {{.*}} [[SCALE]] [[NONE]]
+
+; Small's two operations and one expression are the only ones emitted.
+; ONE: OpExtInst {{.*}} DebugOperation
+; ONE: OpExtInst {{.*}} DebugOperation
+; ONE: OpExtInst {{.*}} DebugExpression
+
+; spirv-val is broken and is rejecting an expression in a DebugGlobalVariable (KhronosGroup/SPIRV-Tools#6469).
+; VAL: DebugGlobalVariable: expected operand Variable must be a result id of
+
+target triple = "spirv64-unknown-unknown"
+
+define spir_func void @f() !dbg !20 {
+entry:
+ ret void, !dbg !21
+}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!1, !2}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_HIP, file: !3, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, globals: !4, splitDebugInlining: false, nameTableKind: None)
+!1 = !{i32 7, !"Dwarf Version", i32 5}
+!2 = !{i32 2, !"Debug Info Version", i32 3}
+!3 = !DIFile(filename: "fe.hip", directory: "/src")
+!4 = !{!5, !8, !11}
+
+; constexpr unsigned long long Hash = 0xff51afd7ed558ccdULL;
+!5 = !DIGlobalVariableExpression(var: !6, expr: !DIExpression(DW_OP_constu, 18397679294719823053, DW_OP_stack_value))
+!6 = distinct !DIGlobalVariable(name: "Hash", linkageName: "Hash", scope: !0, file: !3, line: 2, type: !7, isLocal: true, isDefinition: true)
+!7 = !DIBasicType(name: "unsigned long long", size: 64, encoding: DW_ATE_unsigned)
+
+; constexpr double Scale = 1.5; (0x3FF8000000000000)
+!8 = !DIGlobalVariableExpression(var: !9, expr: !DIExpression(DW_OP_constu, 4609434218613702656, DW_OP_stack_value))
+!9 = distinct !DIGlobalVariable(name: "Scale", linkageName: "Scale", scope: !0, file: !3, line: 3, type: !10, isLocal: true, isDefinition: true)
+!10 = !DIBasicType(name: "double", size: 64, encoding: DW_ATE_float)
+
+; constexpr unsigned Small = 42u;
+!11 = !DIGlobalVariableExpression(var: !12, expr: !DIExpression(DW_OP_constu, 42, DW_OP_stack_value))
+!12 = distinct !DIGlobalVariable(name: "Small", linkageName: "Small", scope: !0, file: !3, line: 4, type: !13, isLocal: true, isDefinition: true)
+!13 = !DIBasicType(name: "unsigned int", size: 32, encoding: DW_ATE_unsigned)
+
+!18 = !DISubroutineType(cc: DW_CC_LLVM_SpirFunction, types: !19)
+!19 = !{null}
+!20 = distinct !DISubprogram(name: "f", scope: !3, file: !3, line: 6, type: !18, scopeLine: 6, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!21 = !DILocation(line: 7, column: 1, scope: !20)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-init-expr.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-init-expr.ll
index fb83a6350ce03..a62fe822555d6 100644
--- a/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-init-expr.ll
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-init-expr.ll
@@ -1,22 +1,48 @@
; RUN: llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | not spirv-val 2>&1 | FileCheck %s --check-prefix=VAL %}
; A DIGlobalVariable with no backing llvm::GlobalVariable but whose
; DIGlobalVariableExpression carries a non-empty DIExpression (a constant
-; initializer). Since DebugExpression emission is not implemented, the Variable
-; operand falls back to DebugInfoNone. Flags encode IsLocal|IsDefinition (12).
+; initializer). Both operations map, so the Variable operand is the resulting
+; DebugExpression rather than DebugInfoNone: Constu (8) carrying 42, then
+; StackValue (7) to say the operand stack holds the value itself.
+; Flags encode IsLocal|IsDefinition (12).
+;
+; spirv-val rejects that, so the second RUN line expects it to fail. The
+; extension permits it: "If the variable is optimized out, this operand can be
+; the <id> of a DebugExpression instruction that contains the constant value of
+; the variable that was optimized out." The validator instead shares one rule
+; with OpenCL.DebugInfo.100, whose wording stops at DebugInfoNone, and checks
+; the operand against a fixed list of OpVariable and constant opcodes.
+;
+; This is KhronosGroup/SPIRV-Tools#6469, open, where the maintainers agree the
+; validator is at fault twice over: it should accept a DebugExpression for
+; NonSemantic.Shader.DebugInfo.100, and it should not be accepting the OpConstant
+; variants that only the OpenCL wording allows. Every version tried rejects it,
+; from 2022.2 to 2026.2, and SPIRV-LLVM-Translator emits the same thing, so this
+; is not a stale-binary problem. When the fix lands this RUN line will start
+; failing, which is the signal to restore a plain spirv-val invocation.
; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0
; CHECK-DAG: [[NAME:%[0-9]+]] = OpString "constg"
; CHECK-DAG: [[STR_INT:%[0-9]+]] = OpString "int"
-; CHECK-DAG: [[C42:%[0-9]+]] = OpConstant [[I32T]] 42
-; CHECK-DAG: [[C12:%[0-9]+]] = OpConstant [[I32T]] 12
-; CHECK-DAG: [[NONE:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugInfoNone
+; The trailing anchors keep e.g. [[C7]] from binding to "OpConstant %3 72".
+; CHECK-DAG: [[C7:%[0-9]+]] = OpConstant [[I32T]] 7{{ *$}}
+; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}}
+; CHECK-DAG: [[C12:%[0-9]+]] = OpConstant [[I32T]] 12{{ *$}}
+; CHECK-DAG: [[C42:%[0-9]+]] = OpConstant [[I32T]] 42{{ *$}}
; CHECK-DAG: [[DS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugSource
; CHECK-DAG: [[DTI:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugTypeBasic [[STR_INT]]
-; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[NAME]] [[DTI]] [[DS]] [[C42]] {{%[0-9]+}} {{%[0-9]+}} [[NAME]] [[NONE]] [[C12]]
+; CHECK-DAG: [[CONSTU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C42]]{{ *$}}
+; CHECK-DAG: [[SV:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C7]]{{ *$}}
+; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[CONSTU]] [[SV]]{{ *$}}
+; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[NAME]] [[DTI]] [[DS]] [[C42]] {{%[0-9]+}} {{%[0-9]+}} [[NAME]] [[EXPR]] [[C12]]
+
+; The operand list the validator will accept has varied across releases, so
+; match only up to it.
+; VAL: DebugGlobalVariable: expected operand Variable must be a result id of
target triple = "spirv64-unknown-unknown"
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-local-variable-dbg-value.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-local-variable-dbg-value.ll
index 83483ac8d5805..91699d80cbe54 100644
--- a/llvm/test/CodeGen/SPIRV/debug-info/debug-local-variable-dbg-value.ll
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-local-variable-dbg-value.ll
@@ -1,7 +1,10 @@
-; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s
+; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare
; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
; Collect DILocalVariable from a #dbg_value record. No retainedNodes.
+;
+; The record is a plain #dbg_value, so it lowers to a direct DBG_VALUE and gets
+; no DebugDeclare.
; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100"
; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid
@@ -16,7 +19,7 @@
; CHECK-DAG: [[DS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugSource [[PATH]]
; CHECK-DAG: [[INT:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugTypeBasic [[INTNAME]] {{.*}} [[C0]]
; CHECK-DAG: [[DF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugFunction {{.*}}
-; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[XNAME]] [[INT]] [[DS]] [[C8]] [[C0]] [[DF]] [[C0]] [[C1]]
+; CHECK: OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[XNAME]] [[INT]] [[DS]] [[C8]] [[C0]] [[DF]] [[C0]] [[C1]]
target triple = "spirv64-unknown-unknown"
>From bab4f38c78425872c58c993ead7d1abd4e771af4 Mon Sep 17 00:00:00 2001
From: Manuel Carrasco <Manuel.Carrasco at amd.com>
Date: Tue, 8 Sep 2026 09:43:42 -0500
Subject: [PATCH 2/2] Fix linter.
---
.../SPIRV/debug-info/debug-declare-module-scope-variable.ll | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll
index 710f3ef0221f5..a8c33b59d29e0 100644
--- a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll
@@ -21,7 +21,7 @@
target triple = "spirv64-unknown-unknown"
- at shared = internal addrspace(3) global i32 undef, align 4
+ at shared = internal addrspace(3) global i32 poison, align 4
define spir_func void @k() !dbg !5 {
entry:
More information about the llvm-branch-commits
mailing list