clang/test/SemaCUDA/static-device-var.cu

// REQUIRES: x86-registered-target
// REQUIRES: amdgpu-registered-target

// RUN: %clang_cc1 -triple nvptx -fcuda-is-device \
// RUN:    -emit-llvm -o - %s -fsyntax-only -verify=dev

// RUN: %clang_cc1 -triple x86_64-gnu-linux \
// RUN:    -emit-llvm -o - %s -fsyntax-only -verify=host

// Checks allowed usage of file-scope and function-scope static variables.

// host-no-diagnostics

#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;

__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}}
}

int* getDeviceSymbol(int *x);

void foo() {
  getDeviceSymbol(&x);
  getDeviceSymbol(&y);
}