60 lines · plain
1// RUN: %clang_cc1 -fsyntax-only -verify %s2 3#include "Inputs/cuda.h"4 5__global__ void g1(int x) {}6 7template <typename T> void t1(T arg) {8 g1<<<arg, arg>>>(1);9}10 11void h1(int x) {}12int h2(int x) { return 1; }13 14int 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.31template <typename > struct a0 {32 template <typename T> static __global__ void Call(T);33};34struct a1 {35 template <typename T> static __global__ void Call(T);36};37template <typename T> struct a2 {38 static __global__ void Call(T);39};40struct a3 {41 static __global__ void Call(int);42 static __global__ void Call(void*);43};44 45struct b {46 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 }52 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 }58 void e() { d0(1); }59};60