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
|
// REQUIRES: x86-registered-target
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -std=c++11 \
// RUN: -emit-llvm -o - %s -fsyntax-only -verify=dev,com
// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
// RUN: -emit-llvm -o - %s -fsyntax-only -verify=host,com
// Checks allowed usage of file-scope and function-scope static variables.
#include "Inputs/cuda.h"
// Checks static variables are allowed in device functions.
__device__ void f1() {
const static int b = 123;
static int a;
}
// Checks static variables are allowd in global functions.
__global__ void k1() {
const static int b = 123;
static int a;
}
// Checks static device and constant variables are allowed in device and
// host functions, and static host variables are not allowed in device
// functions.
static __device__ int x;
static __constant__ int y;
static int z; // dev-note {{host variable declared here}}
__global__ void kernel(int *a) {
a[0] = x;
a[1] = y;
a[2] = z;
// dev-error@-1 {{reference to __host__ variable 'z' in __global__ function}}
}
// Check dynamic initialization of static device variable is not allowed.
namespace TestStaticVarInLambda {
class A {
public:
A(char *);
};
class B {
public:
__device__ B(char *);
};
void fun() {
(void) [](char *c) {
static A var1(c);
static __device__ B var2(c);
// com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
(void) var1;
(void) var2;
};
}
}
int* getDeviceSymbol(int *x);
void foo() {
getDeviceSymbol(&x);
getDeviceSymbol(&y);
}
|
