clang/test/SemaCUDA/kernel-call.cu

   1 // RUN: %clang_cc1 -fsyntax-only -verify %s
   2
   3 #include "Inputs/cuda.h"
   4
   5 __global__ void g1(int x) {}
   6
   7 template <typename T> void t1(T arg) {
   8   g1<<<arg, arg>>>(1);
   9 }
  10
  11 void h1(int x) {}
  12 int h2(int x) { return 1; }
  13
  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 {
  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 };