73 lines · plain
1// RUN: %clang_cc1 %s -triple x86_64-linux-unknown -fsyntax-only -o - -verify2// RUN: %clang_cc1 %s -fcuda-is-device -triple nvptx -fsyntax-only -o - -verify3 4#include "Inputs/cuda.h"5 6// Check that we get an error if we try to call a __device__ function from a7// module initializer.8 9struct 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}}12 __device__ S() {}13 // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}14};15 16S s;17// expected-error@-1 {{no matching constructor for initialization of 'S'}}18 19struct T {20 __host__ __device__ T() {}21};22T t; // No error, this is OK.23 24struct 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}}27 __host__ U() {}28 // expected-note@-1 {{candidate constructor not viable: requires 0 arguments, but 1 was provided}}29 __device__ U(int) {}30 // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}31};32U u(42);33// expected-error@-1 {{no matching constructor for initialization of 'U'}}34 35__device__ int device_fn() { return 42; }36// expected-note@-1 {{candidate function not viable: call to __device__ function from __host__ function}}37int 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.41double pow(double, double);42 43__device__ double pow(double, int);44 45double 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 48constexpr double cpow(double, double) { return 1.0; }49 50constexpr __device__ double cpow(double, int) { return 2.0; }51 52const double CX = cpow(1.0, 1);53const __device__ double CY = cpow(2.0, 2);54 55struct A {56 double pow(double, double);57 58 __device__ double pow(double, int);59 60 constexpr double cpow(double, double) const { return 1.0; }61 62 constexpr __device__ double cpow(double, int) const { return 1.0; }63 64};65 66A a;67double 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 70const A ca;71const double CAX = ca.cpow(1.0, 1);72const __device__ double CAY = ca.cpow(2.0, 2);73