llvm-project

1
// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify
2

3
#include "Inputs/cuda.h"
4

5
struct A {
6
  int a;
7
  __device__ A() { a = 1; }
8
  __device__ ~A() { a = 2; }
9
};
10

11
// This can be a global var since ctor/dtors of data members are not called.
12
union B {
13
  A a;
14
  __device__ B() {}
15
  __device__ ~B() {}
16
};
17

18
// This cannot be a global var since it has a dynamic ctor.
19
union C {
20
  A a;
21
  __device__ C() { a.a = 3; }
22
  __device__ ~C() {}
23
};
24

25
// This cannot be a global var since it has a dynamic dtor.
26
union D {
27
  A a;
28
  __device__ D() { }
29
  __device__ ~D() { a.a = 4; }
30
};
31

32
__device__ B b;
33
__device__ C c;
34
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
35
__device__ D d;
36
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
37

38
__device__ void foo() {
39
  __shared__ B b;
40
  __shared__ C c;
41
  // expected-error@-1 {{initialization is not supported for __shared__ variables}}
42
  __shared__ D d;
43
  // expected-error@-1 {{initialization is not supported for __shared__ variables}}
44
}
45