[llvm] [NVPTX] Do not emit `.target debug` if only line tables are requested (PR #140146)

[Sergei Lebedev via llvm-commits]

https://github.com/superbobry created https://github.com/llvm/llvm-project/pull/140146

According to [*], only the deprecated @@DWARF directives require setting `.target debug`.

[*]: https://docs.nvidia.com/cuda/parallel-thread-execution/#debugging-directives

>From 1a30286770e9b0264aeb71f46eec15b64b690f3b Mon Sep 17 00:00:00 2001
From: Sergei Lebedev <slebedev at google.com>
Date: Thu, 15 May 2025 21:52:41 +0000
Subject: [PATCH] [NVPTX] Do not emit `.target debug` if only line tables are
 requested

According to [*], only the deprecated @@DWARF directives require setting
`.target debug`.

[*]: https://docs.nvidia.com/cuda/parallel-thread-execution/#debugging-directives
---
 llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp   | 2 +-
 llvm/test/DebugInfo/NVPTX/debug-file-loc.ll | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 0e5207cf9b04c..5b5d72e310ce2 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -745,8 +745,8 @@ void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
     switch(CU->getEmissionKind()) {
     case DICompileUnit::NoDebug:
     case DICompileUnit::DebugDirectivesOnly:
-      break;
     case DICompileUnit::LineTablesOnly:
+      break;
     case DICompileUnit::FullDebug:
       HasFullDebugInfo = true;
       break;
diff --git a/llvm/test/DebugInfo/NVPTX/debug-file-loc.ll b/llvm/test/DebugInfo/NVPTX/debug-file-loc.ll
index 9d1d3506a3f19..258f8e79ad63a 100644
--- a/llvm/test/DebugInfo/NVPTX/debug-file-loc.ll
+++ b/llvm/test/DebugInfo/NVPTX/debug-file-loc.ll
@@ -9,7 +9,7 @@
 ;__device__ void bar() {}
 ;}
 
-; CHECK: .target sm_{{[0-9]+}}, debug
+; CHECK: .target sm_{{[0-9]+}}
 
 ; CHECK: .visible .func foo()
 ; CHECK: .loc [[FOO:[0-9]+]] 1 31