Author: accauble Date: 2026-02-26T09:47:50-08:00 New Revision: f55416b6308a4a89f71d7610e993a20404d11841
URL: https://github.com/llvm/llvm-project/commit/f55416b6308a4a89f71d7610e993a20404d11841 DIFF: https://github.com/llvm/llvm-project/commit/f55416b6308a4a89f71d7610e993a20404d11841.diff LOG: [DebugInfo][HIP] Set DW_AT_language field to DW_LANG_HIP when AMD clang is used (#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`). The primary motivation for this change is that ROCgdb has to use a workaround to detect that a program is a HIP application. This removes the guessing and allows proper detection of the HIP language. ROCgdb [already has support for DW_LANG_HIP](https://github.com/ROCm/ROCgdb/commit/316e8393f2ad). Added: clang/test/CodeGenHIP/debug-info-language-hip.hip Modified: clang/lib/CodeGen/CGDebugInfo.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index 401d8bc16d290..b6084a69382ba 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; @@ -1322,6 +1326,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; @@ -1334,6 +1339,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]]) +//. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
