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
|
// RUN: %clang_cc1 -fsyntax-only -verify %s
#include "Inputs/cuda.h"
__global__ void g1(int x) {}
template <typename T> void t1(T arg) {
g1<<<arg, arg>>>(1);
}
void h1(int x) {}
int h2(int x) { return 1; }
int main(void) {
g1<<<1, 1>>>(42);
g1(42); // expected-error {{call to global function 'g1' not configured}}
g1<<<1>>>(42); // expected-error {{too few execution configuration arguments to kernel function call}}
g1<<<1, 1, 0, 0, 0>>>(42); // expected-error {{too many execution configuration arguments to kernel function call}}
t1(1);
h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function 'h1'}}
int (*fp)(int) = h2;
fp<<<1, 1>>>(42); // expected-error {{must have void return type}}
g1<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
}
// Make sure we can call static member kernels.
template <typename > struct a0 {
template <typename T> static __global__ void Call(T);
};
struct a1 {
template <typename T> static __global__ void Call(T);
};
template <typename T> struct a2 {
static __global__ void Call(T);
};
struct a3 {
static __global__ void Call(int);
static __global__ void Call(void*);
};
struct b {
template <typename c> void d0(c arg) {
a0<c>::Call<<<0, 0>>>(arg);
a1::Call<<<0,0>>>(arg);
a2<c>::Call<<<0,0>>>(arg);
a3::Call<<<0, 0>>>(arg);
}
void d1(void* arg) {
a0<void*>::Call<<<0, 0>>>(arg);
a1::Call<<<0,0>>>(arg);
a2<void*>::Call<<<0,0>>>(arg);
a3::Call<<<0, 0>>>(arg);
}
void e() { d0(1); }
};
|
