1 // RUN: %clang_cc1 -fsyntax-only -verify %s
2
3 #include "Inputs/cuda.h"
4
g1(int x)5 __global__ void g1(int x) {}
6
t1(T arg)7 template <typename T> void t1(T arg) {
8 g1<<<arg, arg>>>(1);
9 }
10
h1(int x)11 void h1(int x) {}
h2(int x)12 int h2(int x) { return 1; }
13
main(void)14 int main(void) {
15 g1<<<1, 1>>>(42);
16 g1(42); // expected-error {{call to global function 'g1' not configured}}
17 g1<<<1>>>(42); // expected-error {{too few execution configuration arguments to kernel function call}}
18 g1<<<1, 1, 0, 0, 0>>>(42); // expected-error {{too many execution configuration arguments to kernel function call}}
19
20 t1(1);
21
22 h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function 'h1'}}
23
24 int (*fp)(int) = h2;
25 fp<<<1, 1>>>(42); // expected-error {{must have void return type}}
26
27 g1<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
28 }
29
30 // Make sure we can call static member kernels.
31 template <typename > struct a0 {
32 template <typename T> static __global__ void Call(T);
33 };
34 struct a1 {
35 template <typename T> static __global__ void Call(T);
36 };
37 template <typename T> struct a2 {
38 static __global__ void Call(T);
39 };
40 struct a3 {
41 static __global__ void Call(int);
42 static __global__ void Call(void*);
43 };
44
45 struct b {
d0b46 template <typename c> void d0(c arg) {
47 a0<c>::Call<<<0, 0>>>(arg);
48 a1::Call<<<0,0>>>(arg);
49 a2<c>::Call<<<0,0>>>(arg);
50 a3::Call<<<0, 0>>>(arg);
51 }
d1b52 void d1(void* arg) {
53 a0<void*>::Call<<<0, 0>>>(arg);
54 a1::Call<<<0,0>>>(arg);
55 a2<void*>::Call<<<0,0>>>(arg);
56 a3::Call<<<0, 0>>>(arg);
57 }
eb58 void e() { d0(1); }
59 };
60