[PATCH] D24581: [CUDA] Add test checking our ability to take a function pointer to a __global__ function on the host side.

Justin Lebar via cfe-commits cfe-commits at lists.llvm.org
Wed Sep 14 14:58:46 PDT 2016


This revision was automatically updated to reflect the committed changes.
Closed by commit rL281543: [CUDA] Add test checking our ability to take a function pointer to a… (authored by jlebar).

Changed prior to commit:
  https://reviews.llvm.org/D24581?vs=71423&id=71446#toc

Repository:
  rL LLVM

https://reviews.llvm.org/D24581

Files:
  cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu

Index: cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu
+++ cfe/trunk/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 at -1 {{declared here}}
+#ifdef DEVICE
+// expected-note at -3 {{declared here}}
+#endif
+
+typedef void (*fn_ptr_t)();
+
+__host__ __device__ fn_ptr_t get_ptr_hd() {
+  return kernel;
+#ifdef DEVICE
+  // expected-error at -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}}
+}


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D24581.71446.patch
Type: text/x-patch
Size: 1107 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160914/d0ed6918/attachment.bin>


More information about the cfe-commits mailing list