From: Justin Lebar <jlebar@google.com>
Date: Wed, 14 Sep 2016 21:50:11 +0000 (+0000)
Subject: [CUDA] Add test checking our ability to take a function pointer to a __global__ funct... 
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=7c90bb73d0b53c126efb84b9acc4c0fc3b8500c5;p=clang

[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
---

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