Artem Belevich | ed12690 | 2016-12-07 19:27:16 +0000 | [diff] [blame] | 1 | // RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s |
| 2 | // RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s |
| 3 | |
| 4 | #include "Inputs/cuda.h" |
| 5 | |
| 6 | struct HType {}; // expected-note-re 6 {{candidate constructor {{.*}} not viable: no known conversion from 'DType'}} |
| 7 | struct DType {}; // expected-note-re 6 {{candidate constructor {{.*}} not viable: no known conversion from 'HType'}} |
| 8 | struct HDType {}; |
| 9 | |
| 10 | template <typename T> __host__ HType overload_h_d(T a) { return HType(); } |
| 11 | // expected-note@-1 2 {{candidate template ignored: could not match 'HType' against 'DType'}} |
| 12 | // expected-note@-2 2 {{candidate template ignored: target attributes do not match}} |
| 13 | template <typename T> __device__ DType overload_h_d(T a) { return DType(); } |
| 14 | // expected-note@-1 2 {{candidate template ignored: could not match 'DType' against 'HType'}} |
| 15 | // expected-note@-2 2 {{candidate template ignored: target attributes do not match}} |
| 16 | |
| 17 | // Check explicit instantiation. |
| 18 | template __device__ __host__ DType overload_h_d(int a); // There's no HD template... |
| 19 | // expected-error@-1 {{explicit instantiation of 'overload_h_d' does not refer to a function template, variable template, member function, member class, or static data member}} |
| 20 | template __device__ __host__ HType overload_h_d(int a); // There's no HD template... |
| 21 | // expected-error@-1 {{explicit instantiation of 'overload_h_d' does not refer to a function template, variable template, member function, member class, or static data member}} |
| 22 | template __device__ DType overload_h_d(int a); // OK. instantiates D |
| 23 | template __host__ HType overload_h_d(int a); // OK. instantiates H |
| 24 | |
| 25 | // Check explicit specialization. |
| 26 | template <> __device__ __host__ DType overload_h_d(long a); // There's no HD template... |
| 27 | // expected-error@-1 {{no function template matches function template specialization 'overload_h_d'}} |
| 28 | template <> __device__ __host__ HType overload_h_d(long a); // There's no HD template... |
| 29 | // expected-error@-1 {{no function template matches function template specialization 'overload_h_d'}} |
| 30 | template <> __device__ DType overload_h_d(long a); // OK. instantiates D |
| 31 | template <> __host__ HType overload_h_d(long a); // OK. instantiates H |
| 32 | |
| 33 | |
Artem Belevich | 6074136 | 2016-12-08 19:38:13 +0000 | [diff] [blame] | 34 | // Can't overload HD template with H or D template, though |
| 35 | // non-template functions are OK. |
Artem Belevich | ed12690 | 2016-12-07 19:27:16 +0000 | [diff] [blame] | 36 | template <typename T> __host__ __device__ HDType overload_hd(T a) { return HDType(); } |
| 37 | // expected-note@-1 {{previous declaration is here}} |
| 38 | // expected-note@-2 2 {{candidate template ignored: could not match 'HDType' against 'HType'}} |
| 39 | template <typename T> __device__ HDType overload_hd(T a); |
| 40 | // expected-error@-1 {{__device__ function 'overload_hd' cannot overload __host__ __device__ function 'overload_hd'}} |
| 41 | __device__ HDType overload_hd(int a); // OK. |
| 42 | |
| 43 | // Verify that target attributes are taken into account when we |
Alexander Kornienko | b8b94581 | 2018-04-06 15:14:32 +0000 | [diff] [blame] | 44 | // explicitly specialize or instantiate function templates. |
Artem Belevich | ed12690 | 2016-12-07 19:27:16 +0000 | [diff] [blame] | 45 | template <> __host__ HType overload_hd(int a); |
| 46 | // expected-error@-1 {{no function template matches function template specialization 'overload_hd'}} |
| 47 | template __host__ HType overload_hd(long a); |
| 48 | // expected-error@-1 {{explicit instantiation of 'overload_hd' does not refer to a function template, variable template, member function, member class, or static data member}} |
| 49 | __host__ HType overload_hd(int a); // OK |
| 50 | |
| 51 | template <typename T> __host__ T overload_h(T a); // expected-note {{previous declaration is here}} |
| 52 | template <typename T> __host__ __device__ T overload_h(T a); |
| 53 | // expected-error@-1 {{__host__ __device__ function 'overload_h' cannot overload __host__ function 'overload_h'}} |
| 54 | template <typename T> __device__ T overload_h(T a); // OK. D can overload H. |
| 55 | |
| 56 | template <typename T> __host__ HType overload_h_d2(T a) { return HType(); } |
| 57 | template <typename T> __host__ __device__ HDType overload_h_d2(T a) { return HDType(); } |
| 58 | template <typename T1, typename T2 = int> __device__ DType overload_h_d2(T1 a) { T1 x; T2 y; return DType(); } |
| 59 | |
Artem Belevich | 6074136 | 2016-12-08 19:38:13 +0000 | [diff] [blame] | 60 | // constexpr functions are implicitly HD, but explicit |
| 61 | // instantiation/specialization must use target attributes as written. |
| 62 | template <typename T> constexpr T overload_ce_implicit_hd(T a) { return a+1; } |
| 63 | // expected-note@-1 3 {{candidate template ignored: target attributes do not match}} |
| 64 | |
| 65 | // These will not match the template. |
| 66 | template __host__ __device__ int overload_ce_implicit_hd(int a); |
| 67 | // expected-error@-1 {{explicit instantiation of 'overload_ce_implicit_hd' does not refer to a function template, variable template, member function, member class, or static data member}} |
| 68 | template <> __host__ __device__ long overload_ce_implicit_hd(long a); |
| 69 | // expected-error@-1 {{no function template matches function template specialization 'overload_ce_implicit_hd'}} |
| 70 | template <> __host__ __device__ constexpr long overload_ce_implicit_hd(long a); |
| 71 | // expected-error@-1 {{no function template matches function template specialization 'overload_ce_implicit_hd'}} |
| 72 | |
| 73 | // These should work, because template matching ignores the implicit |
| 74 | // HD attributes the compiler gives to constexpr functions/templates, |
| 75 | // so 'overload_ce_implicit_hd' template will match __host__ functions |
| 76 | // only. |
| 77 | template __host__ int overload_ce_implicit_hd(int a); |
| 78 | template <> __host__ long overload_ce_implicit_hd(long a); |
| 79 | |
| 80 | template float overload_ce_implicit_hd(float a); |
| 81 | template <> float* overload_ce_implicit_hd(float *a); |
| 82 | template <> constexpr double overload_ce_implicit_hd(double a) { return a + 3.0; }; |
| 83 | |
Artem Belevich | ed12690 | 2016-12-07 19:27:16 +0000 | [diff] [blame] | 84 | __host__ void hf() { |
| 85 | overload_hd(13); |
Artem Belevich | 6074136 | 2016-12-08 19:38:13 +0000 | [diff] [blame] | 86 | overload_ce_implicit_hd('h'); // Implicitly instantiated |
| 87 | overload_ce_implicit_hd(1.0f); // Explicitly instantiated |
| 88 | overload_ce_implicit_hd(2.0); // Explicitly specialized |
Artem Belevich | ed12690 | 2016-12-07 19:27:16 +0000 | [diff] [blame] | 89 | |
| 90 | HType h = overload_h_d(10); |
| 91 | HType h2i = overload_h_d2<int>(11); |
| 92 | HType h2ii = overload_h_d2<int>(12); |
| 93 | |
| 94 | // These should be implicitly instantiated from __host__ template returning HType. |
Artem Belevich | 6074136 | 2016-12-08 19:38:13 +0000 | [diff] [blame] | 95 | DType d = overload_h_d(20); // expected-error {{no viable conversion from 'HType' to 'DType'}} |
| 96 | DType d2i = overload_h_d2<int>(21); // expected-error {{no viable conversion from 'HType' to 'DType'}} |
Artem Belevich | ed12690 | 2016-12-07 19:27:16 +0000 | [diff] [blame] | 97 | DType d2ii = overload_h_d2<int>(22); // expected-error {{no viable conversion from 'HType' to 'DType'}} |
| 98 | } |
| 99 | __device__ void df() { |
| 100 | overload_hd(23); |
Artem Belevich | 6074136 | 2016-12-08 19:38:13 +0000 | [diff] [blame] | 101 | overload_ce_implicit_hd('d'); // Implicitly instantiated |
| 102 | overload_ce_implicit_hd(1.0f); // Explicitly instantiated |
| 103 | overload_ce_implicit_hd(2.0); // Explicitly specialized |
Artem Belevich | ed12690 | 2016-12-07 19:27:16 +0000 | [diff] [blame] | 104 | |
| 105 | // These should be implicitly instantiated from __device__ template returning DType. |
Artem Belevich | 6074136 | 2016-12-08 19:38:13 +0000 | [diff] [blame] | 106 | HType h = overload_h_d(10); // expected-error {{no viable conversion from 'DType' to 'HType'}} |
| 107 | HType h2i = overload_h_d2<int>(11); // expected-error {{no viable conversion from 'DType' to 'HType'}} |
Artem Belevich | ed12690 | 2016-12-07 19:27:16 +0000 | [diff] [blame] | 108 | HType h2ii = overload_h_d2<int>(12); // expected-error {{no viable conversion from 'DType' to 'HType'}} |
| 109 | |
| 110 | DType d = overload_h_d(20); |
| 111 | DType d2i = overload_h_d2<int>(21); |
| 112 | DType d2ii = overload_h_d2<int>(22); |
| 113 | } |