73b22935a7
constexpr var may be initialized with address of non-const variable. In this case the initializer is not constant in device compilation. This has been handled for const vars but not for constexpr vars. This patch makes handling of const var and constexpr var consistent. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D119615 Fixes: https://github.com/llvm/llvm-project/issues/53780
106 lines
3.2 KiB
Plaintext
106 lines
3.2 KiB
Plaintext
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
|
|
// RUN: -fsyntax-only -verify
|
|
// RUN: %clang_cc1 -triple x86_64 -x hip %s \
|
|
// RUN: -fsyntax-only -verify=host
|
|
|
|
// host-no-diagnostics
|
|
|
|
#include "Inputs/cuda.h"
|
|
|
|
// Test constexpr var initialized with address of a const var.
|
|
// Both are promoted to device side.
|
|
|
|
namespace Test1 {
|
|
const int a = 1;
|
|
|
|
struct B {
|
|
static constexpr const int *p = &a;
|
|
__device__ static constexpr const int *const p2 = &a;
|
|
};
|
|
|
|
// Const variable 'a' is treated as __constant__ on device side,
|
|
// therefore its address can be used as initializer for another
|
|
// device variable.
|
|
|
|
__device__ void f() {
|
|
int y = a;
|
|
constexpr const int *x = B::p;
|
|
constexpr const int *z = B::p2;
|
|
}
|
|
}
|
|
|
|
// Test constexpr var initialized with address of a non-cost var.
|
|
// Neither is promoted to device side.
|
|
|
|
namespace Test2 {
|
|
int a = 1;
|
|
// expected-note@-1{{host variable declared here}}
|
|
|
|
struct B {
|
|
static constexpr int *const p = &a;
|
|
// expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}
|
|
};
|
|
|
|
__device__ void f() {
|
|
int y = a;
|
|
// expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
|
|
const int *const *x = &B::p;
|
|
// expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}
|
|
// ToDo: use of non-promotable constexpr variable in device compilation should be treated as
|
|
// ODR-use and diagnosed.
|
|
const int *const z = B::p;
|
|
}
|
|
}
|
|
|
|
// Test constexpr device var initialized with address of a non-const host var, __shared var,
|
|
// __managed__ var, __device__ var, __constant__ var, texture var, surface var.
|
|
|
|
namespace Test3 {
|
|
struct textureReference {
|
|
int desc;
|
|
};
|
|
|
|
enum ReadMode {
|
|
ElementType = 0,
|
|
NormalizedFloat = 1
|
|
};
|
|
|
|
template <typename T, int dim = 1, enum ReadMode mode = ElementType>
|
|
struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
|
|
};
|
|
|
|
struct surfaceReference {
|
|
int desc;
|
|
};
|
|
|
|
template <typename T, int dim = 1>
|
|
struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
|
|
};
|
|
|
|
// Partial specialization over `void`.
|
|
template<int dim>
|
|
struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
|
|
};
|
|
|
|
texture<float, 2, ElementType> tex;
|
|
surface<void, 2> surf;
|
|
|
|
int a = 1;
|
|
__shared__ int b;
|
|
__managed__ int c = 1;
|
|
__device__ int d = 1;
|
|
__constant__ int e = 1;
|
|
struct B {
|
|
__device__ static constexpr int *const p1 = &a;
|
|
// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
|
|
__device__ static constexpr int *const p2 = &b;
|
|
// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
|
|
__device__ static constexpr int *const p3 = &c;
|
|
// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
|
|
__device__ static constexpr int *const p4 = &d;
|
|
__device__ static constexpr int *const p5 = &e;
|
|
__device__ static constexpr texture<float, 2, ElementType> *const p6 = &tex;
|
|
__device__ static constexpr surface<void, 2> *const p7 = &surf;
|
|
};
|
|
}
|