1 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
2 // RUN: -fsyntax-only -verify
3 // RUN: %clang_cc1 -triple x86_64 -x hip %s \
4 // RUN: -fsyntax-only -verify=host
5
6 // host-no-diagnostics
7
8 #include "Inputs/cuda.h"
9
10 // Test constexpr var initialized with address of a const var.
11 // Both are promoted to device side.
12
13 namespace Test1 {
14 const int a = 1;
15
16 struct B {
17 static constexpr const int *p = &a;
18 __device__ static constexpr const int *const p2 = &a;
19 };
20
21 // Const variable 'a' is treated as __constant__ on device side,
22 // therefore its address can be used as initializer for another
23 // device variable.
24
f()25 __device__ void f() {
26 int y = a;
27 constexpr const int *x = B::p;
28 constexpr const int *z = B::p2;
29 }
30 }
31
32 // Test constexpr var initialized with address of a non-cost var.
33 // Neither is promoted to device side.
34
35 namespace Test2 {
36 int a = 1;
37 // expected-note@-1{{host variable declared here}}
38
39 struct B {
40 static constexpr int *const p = &a;
41 // expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}
42 };
43
f()44 __device__ void f() {
45 int y = a;
46 // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
47 const int *const *x = &B::p;
48 // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}
49 // ToDo: use of non-promotable constexpr variable in device compilation should be treated as
50 // ODR-use and diagnosed.
51 const int *const z = B::p;
52 }
53 }
54
55 // Test constexpr device var initialized with address of a non-const host var, __shared var,
56 // __managed__ var, __device__ var, __constant__ var, texture var, surface var.
57
58 namespace Test3 {
59 struct textureReference {
60 int desc;
61 };
62
63 enum ReadMode {
64 ElementType = 0,
65 NormalizedFloat = 1
66 };
67
68 template <typename T, int dim = 1, enum ReadMode mode = ElementType>
69 struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
70 };
71
72 struct surfaceReference {
73 int desc;
74 };
75
76 template <typename T, int dim = 1>
77 struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
78 };
79
80 // Partial specialization over `void`.
81 template<int dim>
82 struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
83 };
84
85 texture<float, 2, ElementType> tex;
86 surface<void, 2> surf;
87
88 int a = 1;
89 __shared__ int b;
90 __managed__ int c = 1;
91 __device__ int d = 1;
92 __constant__ int e = 1;
93 struct B {
94 __device__ static constexpr int *const p1 = &a;
95 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
96 __device__ static constexpr int *const p2 = &b;
97 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
98 __device__ static constexpr int *const p3 = &c;
99 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
100 __device__ static constexpr int *const p4 = &d;
101 __device__ static constexpr int *const p5 = &e;
102 __device__ static constexpr texture<float, 2, ElementType> *const p6 = &tex;
103 __device__ static constexpr surface<void, 2> *const p7 = &surf;
104 };
105 }
106