71 lines · plain
1// REQUIRES: x86-registered-target2// REQUIRES: amdgpu-registered-target3 4// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -std=c++11 \5// RUN: -o - %s -fsyntax-only -verify=dev,com6 7// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \8// RUN: -o - %s -fsyntax-only -verify=host,com9 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 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 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 and29// host functions, and static host variables are not allowed in device30// functions.31 32static __device__ int x;33static __constant__ int y;34static int z; // dev-note {{host variable declared here}}35 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 45namespace TestStaticVarInLambda {46class A {47public:48 A(char *);49};50class B {51public:52 __device__ B(char *);53};54void 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 65int* getDeviceSymbol(int *x);66 67void foo() {68 getDeviceSymbol(&x);69 getDeviceSymbol(&y);70}71