diff options
author | Justin Lebar <jlebar@google.com> | 2016-09-14 21:50:11 +0000 |
---|---|---|
committer | Justin Lebar <jlebar@google.com> | 2016-09-14 21:50:11 +0000 |
commit | 7c90bb73d0b53c126efb84b9acc4c0fc3b8500c5 (patch) | |
tree | dbeb33f1f04f011b98e4388f5532912b18d889a9 /test/SemaCUDA | |
parent | 5eb95c4c284486351e3ed0fdad011acf41540c8b (diff) |
[CUDA] Add test checking our ability to take a function pointer to a __global__ function on the host side.
Summary: This functionality is used by Thrust.
Reviewers: tra
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D24581
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@281543 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/SemaCUDA')
-rw-r--r-- | test/SemaCUDA/reference-to-kernel-fn.cu | 31 |
1 files changed, 31 insertions, 0 deletions
diff --git a/test/SemaCUDA/reference-to-kernel-fn.cu b/test/SemaCUDA/reference-to-kernel-fn.cu new file mode 100644 index 0000000000..29efcfa2fa --- /dev/null +++ b/test/SemaCUDA/reference-to-kernel-fn.cu @@ -0,0 +1,31 @@ +// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s +// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify -DDEVICE %s + +// Check that we can reference (get a function pointer to) a __global__ +// function from the host side, but not the device side. (We don't yet support +// device-side kernel launches.) + +#include "Inputs/cuda.h" + +struct Dummy {}; + +__global__ void kernel() {} +// expected-note@-1 {{declared here}} +#ifdef DEVICE +// expected-note@-3 {{declared here}} +#endif + +typedef void (*fn_ptr_t)(); + +__host__ __device__ fn_ptr_t get_ptr_hd() { + return kernel; +#ifdef DEVICE + // expected-error@-2 {{reference to __global__ function}} +#endif +} +__host__ fn_ptr_t get_ptr_h() { + return kernel; +} +__device__ fn_ptr_t get_ptr_d() { + return kernel; // expected-error {{reference to __global__ function}} +} |