1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
|
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST %s
// Negative tests.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV-NEG %s
#include "Inputs/cuda.h"
template <class T>
class A {
static int h_member;
__device__ static int d_member;
__constant__ static int c_member;
__managed__ static int m_member;
const static int const_member = 0;
};
template <class T>
int A<T>::h_member;
template <class T>
__device__ int A<T>::d_member;
template <class T>
__constant__ int A<T>::c_member;
template <class T>
__managed__ int A<T>::m_member;
template <class T>
const int A<T>::const_member;
template class A<int>;
//DEV-DAG: @_ZN1AIiE8d_memberE = internal addrspace(1) global i32 0, comdat, align 4
//DEV-DAG: @_ZN1AIiE8c_memberE = internal addrspace(4) global i32 0, comdat, align 4
//DEV-DAG: @_ZN1AIiE8m_memberE = internal addrspace(1) externally_initialized global ptr addrspace(1) null
//DEV-DAG: @_ZN1AIiE12const_memberE = internal addrspace(4) constant i32 0, comdat, align 4
//DEV-NEG-NOT: @_ZN1AIiE8h_memberE
//HOST-DAG: @_ZN1AIiE8h_memberE = weak_odr global i32 0, comdat, align 4
//HOST-DAG: @_ZN1AIiE8d_memberE = internal global i32 undef, comdat, align 4
//HOST-DAG: @_ZN1AIiE8c_memberE = internal global i32 undef, comdat, align 4
//HOST-DAG: @_ZN1AIiE8m_memberE = internal externally_initialized global ptr null
//HOST-DAG: @_ZN1AIiE12const_memberE = weak_odr constant i32 0, comdat, align 4
|
