clang/test/SemaCUDA/global-initializers.cu

// RUN: %clang_cc1 %s -triple x86_64-linux-unknown -fsyntax-only -o - -verify
// RUN: %clang_cc1 %s -fcuda-is-device -triple nvptx -fsyntax-only -o - -verify

#include "Inputs/cuda.h"

// Check that we get an error if we try to call a __device__ function from a
// module initializer.

struct S {
  // expected-note@-1 {{candidate constructor (the implicit copy constructor) not viable: requires 1 argument, but 0 were provided}}
  // expected-note@-2 {{candidate constructor (the implicit move constructor) not viable: requires 1 argument, but 0 were provided}}
  __device__ S() {}
  // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}
};

S s;
// expected-error@-1 {{no matching constructor for initialization of 'S'}}

struct T {
  __host__ __device__ T() {}
};
T t;  // No error, this is OK.

struct U {
  // expected-note@-1 {{candidate constructor (the implicit copy constructor) not viable: no known conversion from 'int' to 'const U' for 1st argument}}
  // expected-note@-2 {{candidate constructor (the implicit move constructor) not viable: no known conversion from 'int' to 'U' for 1st argument}}
  __host__ U() {}
  // expected-note@-1 {{candidate constructor not viable: requires 0 arguments, but 1 was provided}}
  __device__ U(int) {}
  // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}
};
U u(42);
// expected-error@-1 {{no matching constructor for initialization of 'U'}}

__device__ int device_fn() { return 42; }
// expected-note@-1 {{candidate function not viable: call to __device__ function from __host__ function}}
int n = device_fn();
// expected-error@-1 {{no matching function for call to 'device_fn'}}

// Check host/device-based overloding resolution in global variable initializer.
double pow(double, double);

__device__ double pow(double, int);

double X = pow(1.0, 1);
__device__ double Y = pow(2.0, 2); // expected-error{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}

constexpr double cpow(double, double) { return 1.0; }

constexpr __device__ double cpow(double, int) { return 2.0; }

const double CX = cpow(1.0, 1);
const __device__ double CY = cpow(2.0, 2);

struct A {
  double pow(double, double);

  __device__ double pow(double, int);

  constexpr double cpow(double, double) const { return 1.0; }

  constexpr __device__ double cpow(double, int) const { return 1.0; }

};

A a;
double AX = a.pow(1.0, 1);
__device__ double AY = a.pow(2.0, 2); // expected-error{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}

const A ca;
const double CAX = ca.cpow(1.0, 1);
const __device__ double CAY = ca.cpow(2.0, 2);