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

   1 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify=host \
   2 // RUN:   -verify-ignore-unexpected=note %s
   3 // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify=dev \
   4 // RUN:   -verify-ignore-unexpected=note %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 // host-no-diagnostics
  11
  12 #include "Inputs/cuda.h"
  13
  14 struct Dummy {};
  15
  16 __global__ void kernel() {}
  17
  18 typedef void (*fn_ptr_t)();
  19
  20 __host__ __device__ fn_ptr_t get_ptr_hd() {
  21   return kernel;
  22   // dev-error@-1 {{reference to __global__ function}}
  23 }
  24 __host__ fn_ptr_t get_ptr_h() {
  25   return kernel;
  26 }
  27 __device__ fn_ptr_t get_ptr_d() {
  28   return kernel;  // dev-error {{reference to __global__ function}}
  29 }