| // 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]]) |
| //. |
