annotate clang/test/CodeGenCUDA/debug-info-address-class.cu @ 236:c4bab56944e8 llvm-original

LLVM 16
author kono
date Wed, 09 Nov 2022 17:45:10 +0900
parents 1d019706d866
children 1f2b6ac9f198
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
236
c4bab56944e8 LLVM 16
kono
parents: 150
diff changeset
1 // RUN: %clang_cc1 -no-opaque-pointers -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown -debug-info-kind=limited -dwarf-version=2 -debugger-tuning=gdb | FileCheck %s
150
anatofuz
parents:
diff changeset
2
anatofuz
parents:
diff changeset
3 #include "Inputs/cuda.h"
anatofuz
parents:
diff changeset
4
anatofuz
parents:
diff changeset
5 // CHECK-DAG: ![[FILEVAR0:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
anatofuz
parents:
diff changeset
6 // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR0]], expr: !DIExpression())
anatofuz
parents:
diff changeset
7 __device__ int FileVar0;
anatofuz
parents:
diff changeset
8 // CHECK-DAG: ![[FILEVAR1:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
anatofuz
parents:
diff changeset
9 // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR1]], expr: !DIExpression(DW_OP_constu, 8, DW_OP_swap, DW_OP_xderef))
anatofuz
parents:
diff changeset
10 __device__ __shared__ int FileVar1;
anatofuz
parents:
diff changeset
11 // CHECK-DAG: ![[FILEVAR2:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
anatofuz
parents:
diff changeset
12 // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR2]], expr: !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef))
anatofuz
parents:
diff changeset
13 __device__ __constant__ int FileVar2;
anatofuz
parents:
diff changeset
14
anatofuz
parents:
diff changeset
15 __device__ void kernel1(
anatofuz
parents:
diff changeset
16 // CHECK-DAG: ![[ARG:[0-9]+]] = !DILocalVariable(name: "Arg", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
anatofuz
parents:
diff changeset
17 // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[ARG]], metadata !DIExpression()), !dbg !{{[0-9]+}}
anatofuz
parents:
diff changeset
18 int Arg) {
anatofuz
parents:
diff changeset
19 // CHECK-DAG: ![[FUNCVAR0:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true)
anatofuz
parents:
diff changeset
20 // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR0]], expr: !DIExpression(DW_OP_constu, 8, DW_OP_swap, DW_OP_xderef))
anatofuz
parents:
diff changeset
21 __shared__ int FuncVar0;
anatofuz
parents:
diff changeset
22 // CHECK-DAG: ![[FUNCVAR1:[0-9]+]] = !DILocalVariable(name: "FuncVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
anatofuz
parents:
diff changeset
23 // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression()), !dbg !{{[0-9]+}}
anatofuz
parents:
diff changeset
24 int FuncVar1;
anatofuz
parents:
diff changeset
25 }