test/SemaCUDA/reference-to-kernel-fn.cu

   1 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify \
   2 // RUN:   -verify-ignore-unexpected=note %s
   3 // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify \
   4 // RUN:   -verify-ignore-unexpected=note -DDEVICE %s
   5
   6 // Check that we can reference (get a function pointer to) a __global__
   7 // function from the host side, but not the device side.  (We don't yet support
   8 // device-side kernel launches.)
   9
  10 #include "Inputs/cuda.h"
  11
  12 struct Dummy {};
  13
  14 __global__ void kernel() {}
  15
  16 typedef void (*fn_ptr_t)();
  17
  18 __host__ __device__ fn_ptr_t get_ptr_hd() {
  19   return kernel;
  20 #ifdef DEVICE
  21   // expected-error@-2 {{reference to __global__ function}}
  22 #endif
  23 }
  24 __host__ fn_ptr_t get_ptr_h() {
  25   return kernel;
  26 }
  27 __device__ fn_ptr_t get_ptr_d() {
  28   return kernel;  // expected-error {{reference to __global__ function}}
  29 }