llvm-project
111 строк · 3.2 Кб
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
13namespace Test1 {
14const int a = 1;
15
16struct B {
17static const int *const p;
18__device__ static const int *const p2;
19};
20const 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() {
27int y = a;
28const int *x = B::p;
29const 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
36namespace Test2 {
37int a = 1;
38// expected-note@-1{{host variable declared here}}
39
40struct B {
41static int *const p;
42};
43int *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() {
47int y = a;
48// expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
49const 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
57namespace Test3 {
58struct textureReference {
59int desc;
60};
61
62enum ReadMode {
63ElementType = 0,
64NormalizedFloat = 1
65};
66
67template <typename T, int dim = 1, enum ReadMode mode = ElementType>
68struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
69};
70
71struct surfaceReference {
72int desc;
73};
74
75template <typename T, int dim = 1>
76struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
77};
78
79// Partial specialization over `void`.
80template<int dim>
81struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
82};
83
84texture<float, 2, ElementType> tex;
85surface<void, 2> surf;
86
87int a = 1;
88__shared__ int b;
89__managed__ int c = 1;
90__device__ int d = 1;
91__constant__ int e = 1;
92struct 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