106 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 constexpr 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 constexpr const int *p = &a;18 __device__ static constexpr const int *const p2 = &a;19};20 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 25__device__ void f() {26 int y = a;27 constexpr const int *x = B::p;28 constexpr const int *z = B::p2;29}30}31 32// Test constexpr var initialized with address of a non-cost var.33// Neither is promoted to device side.34 35namespace Test2 {36int a = 1;37// expected-note@-1{{host variable declared here}}38 39struct B {40 static constexpr int *const p = &a;41 // expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}42};43 44__device__ void f() {45 int y = a;46 // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}47 const int *const *x = &B::p;48 // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}49 // ToDo: use of non-promotable constexpr variable in device compilation should be treated as50 // ODR-use and diagnosed.51 const int *const z = B::p;52}53}54 55// Test constexpr device var initialized with address of a non-const host var, __shared var,56// __managed__ var, __device__ var, __constant__ var, texture var, surface var.57 58namespace Test3 {59struct textureReference {60 int desc;61};62 63enum ReadMode {64 ElementType = 0,65 NormalizedFloat = 166};67 68template <typename T, int dim = 1, enum ReadMode mode = ElementType>69struct __attribute__((device_builtin_texture_type)) texture : public textureReference {70};71 72struct surfaceReference {73 int desc;74};75 76template <typename T, int dim = 1>77struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {78};79 80// Partial specialization over `void`.81template<int dim>82struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {83};84 85texture<float, 2, ElementType> tex;86surface<void, 2> surf;87 88int a = 1;89__shared__ int b;90__managed__ int c = 1;91__device__ int d = 1;92__constant__ int e = 1;93struct B {94 __device__ static constexpr int *const p1 = &a;95 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}96 __device__ static constexpr int *const p2 = &b;97 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}98 __device__ static constexpr int *const p3 = &c;99 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}100 __device__ static constexpr int *const p4 = &d;101 __device__ static constexpr int *const p5 = &e;102 __device__ static constexpr texture<float, 2, ElementType> *const p6 = &tex;103 __device__ static constexpr surface<void, 2> *const p7 = &surf;104};105}106