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 const 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 const int *const p;
18 __device__ static const int *const p2;
19 };
20 const int *const B::p = &a;
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 __device__ const int *const B::p2 = &a;
25
f()26 __device__ void f() {
27 int y = a;
28 const int *x = B::p;
29 const int *z = B::p2;
30 }
31 }
32
33 // Test const var initialized with address of a non-cost var.
34 // Neither is promoted to device side.
35
36 namespace Test2 {
37 int a = 1;
38 // expected-note@-1{{host variable declared here}}
39
40 struct B {
41 static int *const p;
42 };
43 int *const B::p = &a;
44 // expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}
45
f()46 __device__ void f() {
47 int y = a;
48 // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
49 const int *x = B::p;
50 // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}
51 }
52 }
53
54 // Test device var initialized with address of a non-const host var, __shared var,
55 // __managed__ var, __device__ var, __constant__ var, texture var, surface var.
56
57 namespace Test3 {
58 struct textureReference {
59 int desc;
60 };
61
62 enum ReadMode {
63 ElementType = 0,
64 NormalizedFloat = 1
65 };
66
67 template <typename T, int dim = 1, enum ReadMode mode = ElementType>
68 struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
69 };
70
71 struct surfaceReference {
72 int desc;
73 };
74
75 template <typename T, int dim = 1>
76 struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
77 };
78
79 // Partial specialization over `void`.
80 template<int dim>
81 struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
82 };
83
84 texture<float, 2, ElementType> tex;
85 surface<void, 2> surf;
86
87 int a = 1;
88 __shared__ int b;
89 __managed__ int c = 1;
90 __device__ int d = 1;
91 __constant__ int e = 1;
92 struct B {
93 __device__ static int *const p1;
94 __device__ static int *const p2;
95 __device__ static int *const p3;
96 __device__ static int *const p4;
97 __device__ static int *const p5;
98 __device__ static texture<float, 2, ElementType> *const p6;
99 __device__ static surface<void, 2> *const p7;
100 };
101 __device__ int *const B::p1 = &a;
102 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
103 __device__ int *const B::p2 = &b;
104 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
105 __device__ int *const B::p3 = &c;
106 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
107 __device__ int *const B::p4 = &d;
108 __device__ int *const B::p5 = &e;
109 __device__ texture<float, 2, ElementType> *const B::p6 = &tex;
110 __device__ surface<void, 2> *const B::p7 = &surf;
111 }
112