1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
|
// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=dev,com %s \
// RUN: -std=c++11 -fgpu-defer-diag
// RUN: %clang_cc1 -fsyntax-only -verify=host,com %s \
// RUN: -std=c++11 -fgpu-defer-diag
// RUN: %clang_cc1 -fopenmp -fsyntax-only -verify=host,com %s \
// RUN: -std=c++11 -fgpu-defer-diag
// With -fgpu-defer-diag, clang defers overloading resolution induced
// diagnostics when the full candidates set include host device
// functions or wrong-sided candidates. This roughly matches nvcc's
// behavior.
#include "Inputs/cuda.h"
// When callee is called by a host function with integer arguments, there is an error for ambiguity.
// It should be deferred since it involves wrong-sided candidates.
__device__ void callee(int);
__host__ void callee(float); // host-note {{candidate function}}
__host__ void callee(double); // host-note {{candidate function}}
// When callee2 is called by a device function without arguments, there is an error for 'no matching function'.
// It should be deferred since it involves wrong-sided candidates.
__host__ void callee2(); // dev-note{{candidate function not viable: call to __host__ function from __device__ function}}
// When callee3 is called by a device function without arguments, there is an error for 'no matching function'.
// It should be deferred since it involves wrong-sided candidates.
__host__ void callee3(); // dev-note{{candidate function not viable: call to __host__ function from __device__ function}}
__device__ void callee3(int); // dev-note{{candidate function not viable: requires 1 argument, but 0 were provided}}
// When callee4 is called by a host or device function without arguments, there is an error for 'no matching function'.
// It should be immediate since it involves no wrong-sided candidates (it is not a viable candiate due to signature).
__host__ void callee4(int); // com-note 2{{candidate function not viable: requires 1 argument, but 0 were provided}}
// When callee5 is called by a host function with integer arguments, there is an error for ambiguity.
// It should be immediate since it involves no wrong-sided candidates.
__host__ void callee5(float); // com-note {{candidate function}}
__host__ void callee5(double); // com-note {{candidate function}}
// When '<<` operator is called by a device function, there is error for 'invalid operands'.
// It should be deferred since it involves wrong-sided candidates.
struct S {
__host__ S &operator <<(int i); // dev-note {{candidate function not viable}}
};
__host__ void hf() {
callee(1); // host-error {{call to 'callee' is ambiguous}}
callee2();
callee3();
callee4(); // com-error {{no matching function for call to 'callee4'}}
callee5(1); // com-error {{call to 'callee5' is ambiguous}}
S s;
s << 1;
undeclared_func(); // com-error {{use of undeclared identifier 'undeclared_func'}}
}
__device__ void df() {
callee(1);
callee2(); // dev-error {{no matching function for call to 'callee2'}}
callee3(); // dev-error {{no matching function for call to 'callee3'}}
callee4(); // com-error {{no matching function for call to 'callee4'}}
S s;
s << 1; // dev-error {{invalid operands to binary expression}}
}
struct A { int x; typedef int isA; };
struct B { int x; };
// This function is invalid for A and B by SFINAE.
// This fails to substitue for A but no diagnostic
// should be emitted.
template<typename T, typename T::foo* = nullptr>
__host__ __device__ void sfinae(T t) { // host-note {{candidate template ignored: substitution failure [with T = B]}}
t.x = 1;
}
// This function is defined for A only by SFINAE.
// Calling it with A should succeed, with B should fail.
// The error should not be deferred since it happens in
// file scope.
template<typename T, typename T::isA* = nullptr>
__host__ __device__ void sfinae(T t) { // host-note {{candidate template ignored: substitution failure [with T = B]}}
t.x = 1;
}
void test_sfinae() {
sfinae(A());
sfinae(B()); // host-error{{no matching function for call to 'sfinae'}}
}
// Make sure throw is diagnosed in OpenMP parallel region in host function.
void test_openmp() {
#pragma omp parallel for
for (int i = 0; i < 10; i++) {
throw 1;
}
}
// If a syntax error causes a function not declared, it cannot
// be deferred.
inline __host__ __device__ void bad_func() { // com-note {{to match this '{'}}
// com-error {{expected '}'}}
|
