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 kernel()16__global__ void kernel() {} 17 18 typedef void (*fn_ptr_t)(); 19 get_ptr_hd()20__host__ __device__ fn_ptr_t get_ptr_hd() { 21 return kernel; 22 // dev-error@-1 {{reference to __global__ function}} 23 } get_ptr_h()24__host__ fn_ptr_t get_ptr_h() { 25 return kernel; 26 } get_ptr_d()27__device__ fn_ptr_t get_ptr_d() { 28 return kernel; // dev-error {{reference to __global__ function}} 29 } 30