constexpr-var.cu | Explore in Territory

// 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;
};
}
llvm/clang/test/SemaCUDA/constexpr-var.cu
