| 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 | } |
|---|
| 30 | |
|---|