llvm-project

1
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
2
// RUN:   -fsyntax-only -verify
3
// RUN: %clang_cc1 -triple x86_64 -x hip %s \
4
// RUN:   -fsyntax-only -verify=host
5

6
// host-no-diagnostics
7

8
#include "Inputs/cuda.h"
9

10
// Test const var initialized with address of a const var.
11
// Both are promoted to device side.
12

13
namespace Test1 {
14
const int a = 1;
15

16
struct B {
17
    static const int *const p;
18
    __device__ static const int *const p2;
19
};
20
const int *const B::p = &a;
21
// Const variable 'a' is treated as __constant__ on device side,
22
// therefore its address can be used as initializer for another
23
// device variable.
24
__device__ const int *const B::p2 = &a;
25

26
__device__ void f() {
27
  int y = a;
28
  const int *x = B::p;
29
  const int *z = B::p2;
30
}
31
}
32

33
// Test const var initialized with address of a non-cost var.
34
// Neither is promoted to device side.
35

36
namespace Test2 {
37
int a = 1;
38
// expected-note@-1{{host variable declared here}}
39

40
struct B {
41
    static int *const p;
42
};
43
int *const B::p = &a;
44
// expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}
45

46
__device__ void f() {
47
  int y = a;
48
  // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
49
  const int *x = B::p;
50
  // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}
51
}
52
}
53

54
// Test device var initialized with address of a non-const host var, __shared var,
55
// __managed__ var, __device__ var, __constant__ var, texture var, surface var.
56

57
namespace Test3 {
58
struct textureReference {
59
  int desc;
60
};
61

62
enum ReadMode {
63
  ElementType = 0,
64
  NormalizedFloat = 1
65
};
66

67
template <typename T, int dim = 1, enum ReadMode mode = ElementType>
68
struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
69
};
70

71
struct surfaceReference {
72
  int desc;
73
};
74

75
template <typename T, int dim = 1>
76
struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
77
};
78

79
// Partial specialization over `void`.
80
template<int dim>
81
struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
82
};
83

84
texture<float, 2, ElementType> tex;
85
surface<void, 2> surf;
86

87
int a = 1;
88
__shared__ int b;
89
__managed__ int c = 1;
90
__device__ int d = 1;
91
__constant__ int e = 1;
92
struct B {
93
    __device__ static int *const p1;
94
    __device__ static int *const p2;
95
    __device__ static int *const p3;
96
    __device__ static int *const p4;
97
    __device__ static int *const p5;
98
    __device__ static texture<float, 2, ElementType> *const p6;
99
    __device__ static surface<void, 2> *const p7;
100
};
101
__device__ int *const B::p1 = &a;
102
// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
103
__device__ int *const B::p2 = &b;
104
// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
105
__device__ int *const B::p3 = &c;
106
// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
107
__device__ int *const B::p4 = &d;
108
__device__ int *const B::p5 = &e;
109
__device__ texture<float, 2, ElementType> *const B::p6 = &tex;
110
__device__ surface<void, 2> *const B::p7 = &surf;
111
}
112