1 // REQUIRES: x86-registered-target
2 // REQUIRES: amdgpu-registered-target
3
4 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device -std=c++11 \
5 // RUN: -o - %s -fsyntax-only -verify=dev,com
6
7 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
8 // RUN: -o - %s -fsyntax-only -verify=host,com
9
10 // Checks allowed usage of file-scope and function-scope static variables.
11
12 #include "Inputs/cuda.h"
13
14 // Checks static variables are allowed in device functions.
15
f1()16 __device__ void f1() {
17 const static int b = 123;
18 static int a;
19 }
20
21 // Checks static variables are allowd in global functions.
22
k1()23 __global__ void k1() {
24 const static int b = 123;
25 static int a;
26 }
27
28 // Checks static device and constant variables are allowed in device and
29 // host functions, and static host variables are not allowed in device
30 // functions.
31
32 static __device__ int x;
33 static __constant__ int y;
34 static int z; // dev-note {{host variable declared here}}
35
kernel(int * a)36 __global__ void kernel(int *a) {
37 a[0] = x;
38 a[1] = y;
39 a[2] = z;
40 // dev-error@-1 {{reference to __host__ variable 'z' in __global__ function}}
41 }
42
43 // Check dynamic initialization of static device variable is not allowed.
44
45 namespace TestStaticVarInLambda {
46 class A {
47 public:
48 A(char *);
49 };
50 class B {
51 public:
52 __device__ B(char *);
53 };
fun()54 void fun() {
55 (void) [](char *c) {
56 static A var1(c);
57 static __device__ B var2(c);
58 // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
59 (void) var1;
60 (void) var2;
61 };
62 }
63 }
64
65 int* getDeviceSymbol(int *x);
66
foo()67 void foo() {
68 getDeviceSymbol(&x);
69 getDeviceSymbol(&y);
70 }
71