[clang] Set DW_AT_language field to DW_LANG_HIP when AMD clang is used (PR #181738)

[via cfe-commits]

https://github.com/accauble created https://github.com/llvm/llvm-project/pull/181738

Before this change, HIP applications compiled with AMD clang would set the `DW_AT_language` field to `DW_LANG_C_plus_plus_14`:

```
llvm-dwarfdump hello
hello:  file format elf64-x86-64

.debug_info contents:
0x00000000: ...

0x0000000c: DW_TAG_compile_unit
              ...
  ------> DW_AT_language  (DW_LANG_C_plus_plus_14)
              DW_AT_name  ("helloworld.cpp")
...
```

This change update the language tag to `DW_LANG_HIP` for these applications:

```
llvm-dwarfdump hello
hello:  file format elf64-x86-64

.debug_info contents:
0x00000000: ...

0x0000000c: DW_TAG_compile_unit
              ...
  ------> DW_AT_language  (DW_LANG_HIP)
              DW_AT_name  ("helloworld.cpp")
...
```

Also added a test to check for the new tag (`debug-info-language-hip.hip`).

>From aa7ccc4d4ce5c943db3cb7cf0109f41b280347b5 Mon Sep 17 00:00:00 2001
From: Allyson Cauble-Chantrenne <Allyson.Cauble-Chantrenne at amd.com>
Date: Mon, 16 Feb 2026 14:06:15 -0600
Subject: [PATCH] Updated the dwarf language tag for HIP C/C++ applications to
 DW_LANG_HIP instead of DW_LANG_C_plus_plus. Added a test to check for the new
 tag (debug-info-language-hip.hip).

---
 clang/lib/CodeGen/CGDebugInfo.cpp             | 10 ++++--
 .../CodeGenHIP/debug-info-language-hip.hip    | 36 +++++++++++++++++++
 2 files changed, 44 insertions(+), 2 deletions(-)
 create mode 100644 clang/test/CodeGenHIP/debug-info-language-hip.hip

diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp
index 5e452245ee627..2cf51fb5707ff 100644
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -694,7 +694,9 @@ static llvm::dwarf::SourceLanguage GetSourceLanguage(const CodeGenModule &CGM) {
 
   llvm::dwarf::SourceLanguage LangTag;
   if (LO.CPlusPlus) {
-    if (LO.ObjC)
+    if (LO.HIP)
+      LangTag = llvm::dwarf::DW_LANG_HIP;
+    else if (LO.ObjC)
       LangTag = llvm::dwarf::DW_LANG_ObjC_plus_plus;
     else if (CGO.DebugStrictDwarf && CGO.DwarfVersion < 5)
       LangTag = llvm::dwarf::DW_LANG_C_plus_plus;
@@ -730,7 +732,9 @@ GetDISourceLanguageName(const CodeGenModule &CGM) {
   uint32_t LangVersion = 0;
   llvm::dwarf::SourceLanguageName LangTag;
   if (LO.CPlusPlus) {
-    if (LO.ObjC) {
+    if (LO.HIP) {
+      LangTag = llvm::dwarf::DW_LNAME_HIP;
+    } else if (LO.ObjC) {
       LangTag = llvm::dwarf::DW_LNAME_ObjC_plus_plus;
     } else {
       LangTag = llvm::dwarf::DW_LNAME_C_plus_plus;
@@ -1317,6 +1321,7 @@ static bool hasCXXMangling(llvm::dwarf::SourceLanguage Lang, bool IsTagDecl) {
   case llvm::dwarf::DW_LANG_C_plus_plus:
   case llvm::dwarf::DW_LANG_C_plus_plus_11:
   case llvm::dwarf::DW_LANG_C_plus_plus_14:
+  case llvm::dwarf::DW_LANG_HIP:
     return true;
   case llvm::dwarf::DW_LANG_ObjC_plus_plus:
     return IsTagDecl;
@@ -1329,6 +1334,7 @@ static bool hasCXXMangling(llvm::dwarf::SourceLanguageName Lang,
                            bool IsTagDecl) {
   switch (Lang) {
   case llvm::dwarf::DW_LNAME_C_plus_plus:
+  case llvm::dwarf::DW_LNAME_HIP:
     return true;
   case llvm::dwarf::DW_LNAME_ObjC_plus_plus:
     return IsTagDecl;
diff --git a/clang/test/CodeGenHIP/debug-info-language-hip.hip b/clang/test/CodeGenHIP/debug-info-language-hip.hip
new file mode 100644
index 0000000000000..e68a6dccd8cd9
--- /dev/null
+++ b/clang/test/CodeGenHIP/debug-info-language-hip.hip
@@ -0,0 +1,36 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -emit-llvm -triple amdgcn-amd-amdhsa -x hip \
+// RUN:   -fcuda-is-device -debug-info-kind=limited -o - %s \
+// RUN:   | FileCheck %s
+
+// Verify that HIP code gets DW_LANG_HIP as the language tag in debug info
+// instead of DW_LANG_C_plus_plus.
+
+#define __device__ __attribute__((device))
+
+// CHECK-LABEL: define dso_local void @_Z12hip_functionv(
+// CHECK-SAME: ) #[[ATTR0:[0-9]+]] !dbg [[DBG7:![0-9]+]] {
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    [[X:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
+// CHECK-NEXT:      #dbg_declare(ptr addrspace(5) [[X]], [[META12:![0-9]+]], !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef), [[META14:![0-9]+]])
+// CHECK-NEXT:    store i32 42, ptr [[X_ASCAST]], align 4, !dbg [[META14]]
+// CHECK-NEXT:    ret void, !dbg [[DBG15:![0-9]+]]
+//
+__device__ void hip_function() {
+  int x = 42;
+}
+//.
+// CHECK: [[META0:![0-9]+]] = distinct !DICompileUnit(language: DW_LANG_HIP, file: [[META1:![0-9]+]], producer: "{{.*}}clang version {{.*}}", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
+// CHECK: [[META1]] = !DIFile(filename: "{{.*}}<stdin>", directory: {{.*}})
+// CHECK: [[DBG7]] = distinct !DISubprogram(name: "hip_function", linkageName: "_Z12hip_functionv", scope: [[META8:![0-9]+]], file: [[META8]], line: 21, type: [[META9:![0-9]+]], scopeLine: 21, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: [[META0]], retainedNodes: [[META11:![0-9]+]])
+// CHECK: [[META8]] = !DIFile(filename: "{{.*}}debug-info-language-hip.hip", directory: {{.*}})
+// CHECK: [[META9]] = !DISubroutineType(types: [[META10:![0-9]+]])
+// CHECK: [[META10]] = !{null}
+// CHECK: [[META11]] = !{}
+// CHECK: [[META12]] = !DILocalVariable(name: "x", scope: [[DBG7]], file: [[META8]], line: 22, type: [[META13:![0-9]+]])
+// CHECK: [[META13]] = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+// CHECK: [[META14]] = !DILocation(line: 22, column: 7, scope: [[DBG7]])
+// CHECK: [[DBG15]] = !DILocation(line: 23, column: 1, scope: [[DBG7]])
+//.