~shimon/dismaltc

ref: 461e0617cb03f1d5027af4fbb3da0722791022e9 dismaltc/clang/disable-test/SemaCUDA/function-template-overload.cu -rw-r--r-- 6.8 KiB
461e0617Itoh Shimon re-enable tests. 1 year, 6 months ago
                                                                                
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
104
105
106
107
108
109
110
111
112
113
// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s

#include "Inputs/cuda.h"

struct HType {}; // expected-note-re 6 {{candidate constructor {{.*}} not viable: no known conversion from 'DType'}}
struct DType {}; // expected-note-re 6 {{candidate constructor {{.*}} not viable: no known conversion from 'HType'}}
struct HDType {};

template <typename T> __host__ HType overload_h_d(T a) { return HType(); }
// expected-note@-1 2 {{candidate template ignored: could not match 'HType' against 'DType'}}
// expected-note@-2 2 {{candidate template ignored: target attributes do not match}}
template <typename T> __device__ DType overload_h_d(T a) { return DType(); }
// expected-note@-1 2 {{candidate template ignored: could not match 'DType' against 'HType'}}
// expected-note@-2 2 {{candidate template ignored: target attributes do not match}}

// Check explicit instantiation.
template  __device__ __host__ DType overload_h_d(int a); // There's no HD template...
// 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}}
template  __device__ __host__ HType overload_h_d(int a); // There's no HD template...
// 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}}
template  __device__ DType overload_h_d(int a); // OK. instantiates D
template  __host__ HType overload_h_d(int a); // OK. instantiates H

// Check explicit specialization.
template  <> __device__ __host__ DType overload_h_d(long a); // There's no HD template...
// expected-error@-1 {{no function template matches function template specialization 'overload_h_d'}}
template  <> __device__ __host__ HType overload_h_d(long a); // There's no HD template...
// expected-error@-1 {{no function template matches function template specialization 'overload_h_d'}}
template  <> __device__ DType overload_h_d(long a); // OK. instantiates D
template  <> __host__ HType overload_h_d(long a); // OK. instantiates H


// Can't overload HD template with H or D template, though
// non-template functions are OK.
template <typename T> __host__ __device__ HDType overload_hd(T a) { return HDType(); }
// expected-note@-1 {{previous declaration is here}}
// expected-note@-2 2 {{candidate template ignored: could not match 'HDType' against 'HType'}}
template <typename T> __device__ HDType overload_hd(T a);
// expected-error@-1 {{__device__ function 'overload_hd' cannot overload __host__ __device__ function 'overload_hd'}}
__device__ HDType overload_hd(int a); // OK.

// Verify that target attributes are taken into account when we
// explicitly specialize or instantiate function templates.
template <> __host__ HType overload_hd(int a);
// expected-error@-1 {{no function template matches function template specialization 'overload_hd'}}
template __host__ HType overload_hd(long a);
// 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}}
__host__ HType overload_hd(int a); // OK

template <typename T> __host__ T overload_h(T a); // expected-note {{previous declaration is here}}
template <typename T> __host__ __device__ T overload_h(T a);
// expected-error@-1 {{__host__ __device__ function 'overload_h' cannot overload __host__ function 'overload_h'}}
template <typename T> __device__ T overload_h(T a); // OK. D can overload H.

template <typename T> __host__ HType overload_h_d2(T a) { return HType(); }
template <typename T> __host__ __device__ HDType overload_h_d2(T a) { return HDType(); }
template <typename T1, typename T2 = int> __device__ DType overload_h_d2(T1 a) { T1 x; T2 y; return DType(); }

// constexpr functions are implicitly HD, but explicit
// instantiation/specialization must use target attributes as written.
template <typename T> constexpr T overload_ce_implicit_hd(T a) { return a+1; }
// expected-note@-1 3 {{candidate template ignored: target attributes do not match}}

// These will not match the template.
template __host__ __device__ int overload_ce_implicit_hd(int a);
// 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}}
template <> __host__ __device__ long overload_ce_implicit_hd(long a);
// expected-error@-1 {{no function template matches function template specialization 'overload_ce_implicit_hd'}}
template <> __host__ __device__ constexpr long overload_ce_implicit_hd(long a);
// expected-error@-1 {{no function template matches function template specialization 'overload_ce_implicit_hd'}}

// These should work, because template matching ignores the implicit
// HD attributes the compiler gives to constexpr functions/templates,
// so 'overload_ce_implicit_hd' template will match __host__ functions
// only.
template __host__ int overload_ce_implicit_hd(int a);
template <> __host__ long overload_ce_implicit_hd(long a);

template float overload_ce_implicit_hd(float a);
template <> float* overload_ce_implicit_hd(float *a);
template <> constexpr double overload_ce_implicit_hd(double a) { return a + 3.0; };

__host__ void hf() {
  overload_hd(13);
  overload_ce_implicit_hd('h');        // Implicitly instantiated
  overload_ce_implicit_hd(1.0f);       // Explicitly instantiated
  overload_ce_implicit_hd(2.0);        // Explicitly specialized

  HType h = overload_h_d(10);
  HType h2i = overload_h_d2<int>(11);
  HType h2ii = overload_h_d2<int>(12);

  // These should be implicitly instantiated from __host__ template returning HType.
  DType d = overload_h_d(20);          // expected-error {{no viable conversion from 'HType' to 'DType'}}
  DType d2i = overload_h_d2<int>(21);  // expected-error {{no viable conversion from 'HType' to 'DType'}}
  DType d2ii = overload_h_d2<int>(22); // expected-error {{no viable conversion from 'HType' to 'DType'}}
}
__device__ void df() {
  overload_hd(23);
  overload_ce_implicit_hd('d');        // Implicitly instantiated
  overload_ce_implicit_hd(1.0f);       // Explicitly instantiated
  overload_ce_implicit_hd(2.0);        // Explicitly specialized

  // These should be implicitly instantiated from __device__ template returning DType.
  HType h = overload_h_d(10);          // expected-error {{no viable conversion from 'DType' to 'HType'}}
  HType h2i = overload_h_d2<int>(11);  // expected-error {{no viable conversion from 'DType' to 'HType'}}
  HType h2ii = overload_h_d2<int>(12); // expected-error {{no viable conversion from 'DType' to 'HType'}}

  DType d = overload_h_d(20);
  DType d2i = overload_h_d2<int>(21);
  DType d2ii = overload_h_d2<int>(22);
}