llvm-project

1
// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=dev,com %s \
2
// RUN:   -std=c++11 -fgpu-defer-diag
3
// RUN: %clang_cc1 -fsyntax-only -verify=host,com %s \
4
// RUN:   -std=c++11 -fgpu-defer-diag
5
// RUN: %clang_cc1 -fopenmp -fsyntax-only -verify=host,com %s \
6
// RUN:   -std=c++11 -fgpu-defer-diag
7

8
// With -fgpu-defer-diag, clang defers overloading resolution induced
9
// diagnostics when the full candidates set include host device
10
// functions or wrong-sided candidates. This roughly matches nvcc's
11
// behavior.
12

13
#include "Inputs/cuda.h"
14

15
// When callee is called by a host function with integer arguments, there is an error for ambiguity.
16
// It should be deferred since it involves wrong-sided candidates.
17
__device__ void callee(int);
18
__host__ void callee(float); // host-note {{candidate function}}
19
__host__ void callee(double); // host-note {{candidate function}}
20

21
// When callee2 is called by a device function without arguments, there is an error for 'no matching function'.
22
// It should be deferred since it involves wrong-sided candidates.
23
__host__ void callee2(); // dev-note{{candidate function not viable: call to __host__ function from __device__ function}}
24

25
// When callee3 is called by a device function without arguments, there is an error for 'no matching function'.
26
// It should be deferred since it involves wrong-sided candidates.
27
__host__ void callee3(); // dev-note{{candidate function not viable: call to __host__ function from __device__ function}}
28
__device__ void callee3(int); // dev-note{{candidate function not viable: requires 1 argument, but 0 were provided}}
29

30
// When callee4 is called by a host or device function without arguments, there is an error for 'no matching function'.
31
// It should be immediate since it involves no wrong-sided candidates (it is not a viable candiate due to signature).
32
__host__ void callee4(int); // com-note 2{{candidate function not viable: requires 1 argument, but 0 were provided}}
33

34
// When callee5 is called by a host function with integer arguments, there is an error for ambiguity.
35
// It should be immediate since it involves no wrong-sided candidates.
36
__host__ void callee5(float); // com-note {{candidate function}}
37
__host__ void callee5(double); // com-note {{candidate function}}
38

39
// When '<<` operator is called by a device function, there is error for 'invalid operands'.
40
// It should be deferred since it involves wrong-sided candidates.
41
struct S {
42
  __host__ S &operator <<(int i); // dev-note {{candidate function not viable}}
43
};
44

45
__host__ void hf() {
46
 callee(1); // host-error {{call to 'callee' is ambiguous}}
47
 callee2();
48
 callee3();
49
 callee4(); // com-error {{no matching function for call to 'callee4'}}
50
 callee5(1); // com-error {{call to 'callee5' is ambiguous}}
51
 S s;
52
 s << 1;
53
 undeclared_func(); // com-error {{use of undeclared identifier 'undeclared_func'}}
54
}
55

56
__device__ void df() {
57
 callee(1);
58
 callee2(); // dev-error {{no matching function for call to 'callee2'}}
59
 callee3(); // dev-error {{no matching function for call to 'callee3'}}
60
 callee4(); // com-error {{no matching function for call to 'callee4'}}
61
 S s;
62
 s << 1;    // dev-error {{invalid operands to binary expression}}
63
}
64

65
struct A { int x; typedef int isA; };
66
struct B { int x; };
67

68
// This function is invalid for A and B by SFINAE.
69
// This fails to substitue for A but no diagnostic
70
// should be emitted.
71
template<typename T, typename T::foo* = nullptr>
72
__host__ __device__ void sfinae(T t) { // host-note {{candidate template ignored: substitution failure [with T = B]}}
73
  t.x = 1;
74
}
75

76
// This function is defined for A only by SFINAE.
77
// Calling it with A should succeed, with B should fail.
78
// The error should not be deferred since it happens in
79
// file scope.
80

81
template<typename T, typename T::isA* = nullptr>
82
__host__ __device__ void sfinae(T t) { // host-note {{candidate template ignored: substitution failure [with T = B]}}
83
  t.x = 1;
84
}
85

86
void test_sfinae() {
87
  sfinae(A());
88
  sfinae(B()); // host-error{{no matching function for call to 'sfinae'}}
89
}
90

91
// Make sure throw is diagnosed in OpenMP parallel region in host function.
92
void test_openmp() {
93
  #pragma omp parallel for
94
  for (int i = 0; i < 10; i++) {
95
    throw 1;
96
  }
97
}
98

99
// If a syntax error causes a function not declared, it cannot
100
// be deferred.
101

102
inline __host__ __device__ void bad_func() { // com-note {{to match this '{'}}
103
// com-error {{expected '}'}}
104