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
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
|
//RUN: %clang_cc1 %s -pedantic -ast-dump -verify | FileCheck %s
//expected-no-diagnostics
//CHECK: |-VarDecl {{.*}} foo 'const __global int'
constexpr int foo = 0;
//CHECK: |-VarDecl {{.*}} foo1 'T' cinit
//CHECK: `-VarTemplateSpecializationDecl {{.*}} used foo1 '__global long':'__global long' cinit
template <typename T>
T foo1 = 0;
class c {
public:
//CHECK: `-VarDecl {{.*}} foo2 'const __global int'
static constexpr int foo2 = 0;
};
struct c1 {};
// We only deduce addr space in type alias in pointer types.
//CHECK: TypeAliasDecl {{.*}} alias_c1 'c1'
using alias_c1 = c1;
//CHECK: TypeAliasDecl {{.*}} alias_c1_ptr '__generic c1 *'
using alias_c1_ptr = c1 *;
struct c2 {
alias_c1 y;
alias_c1_ptr ptr = &y;
};
// Addr spaces for pointee of dependent types are not deduced
// during parsing but during template instantiation instead.
template <class T>
struct x1 {
//CHECK: -CXXMethodDecl {{.*}} operator= 'x1<T> &(const x1<T> &){{( __attribute__.*)?}} __generic'
//CHECK: -CXXMethodDecl {{.*}} operator= '__generic x1<int> &(const __generic x1<int> &__private){{( __attribute__.*)?}} __generic'
x1<T>& operator=(const x1<T>& xx) {
y = xx.y;
return *this;
}
int y;
};
template <class T>
struct x2 {
//CHECK: -CXXMethodDecl {{.*}} foo 'void (x1<T> *){{( __attribute__.*)?}} __generic'
//CHECK: -CXXMethodDecl {{.*}} foo 'void (__generic x1<int> *__private){{( __attribute__.*)?}} __generic'
void foo(x1<T>* xx) {
m[0] = *xx;
}
//CHECK: -FieldDecl {{.*}} m 'x1<int> [2]'
x1<T> m[2];
};
void bar(__global x1<int> *xx, __global x2<int> *bar) {
bar->foo(xx);
}
template <typename T>
class x3 : public T {
public:
//CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic'
x3(const x3 &t);
};
//CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic'
template <typename T>
x3<T>::x3(const x3<T> &t) {}
template <class T>
T xxx(T *in1, T in2) {
// This pointer can't be deduced to generic because addr space
// will be taken from the template argument.
//CHECK: `-VarDecl {{.*}} 'T *' cinit
//CHECK: `-VarDecl {{.*}} i '__private int *__private' cinit
T *i = in1;
T ii;
__private T *ptr = ⅈ
ptr = &in2;
return *i;
}
__kernel void test() {
int foo[10];
xxx<__private int>(&foo[0], foo[0]);
// FIXME: Template param deduction fails here because
// temporaries are not in the __private address space.
// It is probably reasonable to put them in __private
// considering that stack and function params are
// implicitly in __private.
// However, if temporaries are left in default addr
// space we should at least pretty print the __private
// addr space. Otherwise diagnostic apprears to be
// confusing.
//xxx(&foo[0], foo[0]);
}
// Addr space for pointer/reference to an array
//CHECK: FunctionDecl {{.*}} t1 'void (const float (__generic &__private)[2])'
void t1(const float (&fYZ)[2]);
//CHECK: FunctionDecl {{.*}} t2 'void (const float (__generic *__private)[2])'
void t2(const float (*fYZ)[2]);
//CHECK: FunctionDecl {{.*}} t3 'void (float (((__generic *__private)))[2])'
void t3(float(((*fYZ)))[2]);
//CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *__private)))[2])'
void t4(float(((**fYZ)))[2]);
//CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *__private))[2])'
void t5(float (*(*fYZ))[2]);
__kernel void k() {
__local float x[2];
float(*p)[2];
t1(x);
t2(&x);
t3(&x);
t4(&p);
t5(&p);
long f1 = foo1<long>;
}
|
