112 lines · plain
1// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \2// RUN: -fsyntax-only -verify3// RUN: %clang_cc1 -triple x86_64 -x hip %s \4// RUN: -fsyntax-only -verify=host5 6// host-no-diagnostics7 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 {17 static 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 another23// 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 36namespace Test2 {37int a = 1;38// expected-note@-1{{host variable declared here}}39 40struct B {41 static 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() {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 57namespace Test3 {58struct textureReference {59 int desc;60};61 62enum ReadMode {63 ElementType = 0,64 NormalizedFloat = 165};66 67template <typename T, int dim = 1, enum ReadMode mode = ElementType>68struct __attribute__((device_builtin_texture_type)) texture : public textureReference {69};70 71struct surfaceReference {72 int 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