1 // RUN: %clang_cc1 %s -triple x86_64-linux-unknown -fsyntax-only -o - -verify
2 // RUN: %clang_cc1 %s -fcuda-is-device -triple nvptx -fsyntax-only -o - -verify
3
4 #include "Inputs/cuda.h"
5
6 // Check that we get an error if we try to call a __device__ function from a
7 // module initializer.
8
9 struct S {
10 // expected-note@-1 {{candidate constructor (the implicit copy constructor) not viable: requires 1 argument, but 0 were provided}}
11 // expected-note@-2 {{candidate constructor (the implicit move constructor) not viable: requires 1 argument, but 0 were provided}}
SS12 __device__ S() {}
13 // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}
14 };
15
16 S s;
17 // expected-error@-1 {{no matching constructor for initialization of 'S'}}
18
19 struct T {
TT20 __host__ __device__ T() {}
21 };
22 T t; // No error, this is OK.
23
24 struct U {
25 // expected-note@-1 {{candidate constructor (the implicit copy constructor) not viable: no known conversion from 'int' to 'const U' for 1st argument}}
26 // expected-note@-2 {{candidate constructor (the implicit move constructor) not viable: no known conversion from 'int' to 'U' for 1st argument}}
UU27 __host__ U() {}
28 // expected-note@-1 {{candidate constructor not viable: requires 0 arguments, but 1 was provided}}
UU29 __device__ U(int) {}
30 // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}
31 };
32 U u(42);
33 // expected-error@-1 {{no matching constructor for initialization of 'U'}}
34
device_fn()35 __device__ int device_fn() { return 42; }
36 // expected-note@-1 {{candidate function not viable: call to __device__ function from __host__ function}}
37 int n = device_fn();
38 // expected-error@-1 {{no matching function for call to 'device_fn'}}
39
40 // Check host/device-based overloding resolution in global variable initializer.
41 double pow(double, double);
42
43 __device__ double pow(double, int);
44
45 double X = pow(1.0, 1);
46 __device__ double Y = pow(2.0, 2); // expected-error{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
47
cpow(double,double)48 constexpr double cpow(double, double) { return 1.0; }
49
cpow(double,int)50 constexpr __device__ double cpow(double, int) { return 2.0; }
51
52 const double CX = cpow(1.0, 1);
53 const __device__ double CY = cpow(2.0, 2);
54
55 struct A {
56 double pow(double, double);
57
58 __device__ double pow(double, int);
59
cpowA60 constexpr double cpow(double, double) const { return 1.0; }
61
cpowA62 constexpr __device__ double cpow(double, int) const { return 1.0; }
63
64 };
65
66 A a;
67 double AX = a.pow(1.0, 1);
68 __device__ double AY = a.pow(2.0, 2); // expected-error{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
69
70 const A ca;
71 const double CAX = ca.cpow(1.0, 1);
72 const __device__ double CAY = ca.cpow(2.0, 2);
73