r283963 - [CUDA] Make touching a kernel from a __host__ __device__ function a deferred error.

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Tue Oct 11 20:30:08 2016
New Revision: 283963

URL: http://llvm.org/viewvc/llvm-project?rev=283963&view=rev
Log:
[CUDA] Make touching a kernel from a __host__ __device__ function a deferred error.

Previously, this was an immediate, don't pass go, don't collect $200
error.  But this precludes us from writing code like

  __host__ __device__ void launch_kernel() {
    kernel<<<...>>>();
  }

Such code isn't wrong, following our notions of right and wrong in CUDA,
unless it's codegen'ed.

Added:
    cfe/trunk/test/SemaCUDA/function-overload-hd.cu
Modified:
    cfe/trunk/lib/Sema/SemaCUDA.cpp
    cfe/trunk/test/SemaCUDA/function-overload.cu
    cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu

Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=283963&r1=283962&r2=283963&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Tue Oct 11 20:30:08 2016
@@ -120,8 +120,7 @@ Sema::IdentifyCUDAPreference(const Funct
   // (a) Can't call global from some contexts until we support CUDA's
   // dynamic parallelism.
   if (CalleeTarget == CFT_Global &&
-      (CallerTarget == CFT_Global || CallerTarget == CFT_Device ||
-       (CallerTarget == CFT_HostDevice && getLangOpts().CUDAIsDevice)))
+      (CallerTarget == CFT_Global || CallerTarget == CFT_Device))
     return CFP_Never;
 
   // (b) Calling HostDevice is OK for everyone.

Added: cfe/trunk/test/SemaCUDA/function-overload-hd.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload-hd.cu?rev=283963&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-overload-hd.cu (added)
+++ cfe/trunk/test/SemaCUDA/function-overload-hd.cu Tue Oct 11 20:30:08 2016
@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -S -o /dev/null -verify \
+// RUN:   -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -S -o /dev/null -fcuda-is-device \
+// RUN:   -verify -verify-ignore-unexpected=note %s
+
+#include "Inputs/cuda.h"
+
+// FIXME: Merge into function-overload.cu once deferred errors can be emitted
+// when non-deferred errors are present.
+
+#if !defined(__CUDA_ARCH__)
+//expected-no-diagnostics
+#endif
+
+typedef void (*GlobalFnPtr)();  // __global__ functions must return void.
+
+__global__ void g() {}
+
+__host__ __device__ void hd() {
+  GlobalFnPtr fp_g = g;
+#if defined(__CUDA_ARCH__)
+  // expected-error at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif
+  g<<<0,0>>>();
+#if defined(__CUDA_ARCH__)
+  // expected-error at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif  // __CUDA_ARCH__
+}

Modified: cfe/trunk/test/SemaCUDA/function-overload.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=283963&r1=283962&r2=283963&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-overload.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-overload.cu Tue Oct 11 20:30:08 2016
@@ -181,18 +181,7 @@ __host__ __device__ void hostdevicef() {
   CurrentFnPtr fp_cdh = cdh;
   CurrentReturnTy ret_cdh = cdh();
 
-  GlobalFnPtr fp_g = g;
-#if defined(__CUDA_ARCH__)
-  // expected-error at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif
-  g();
-  g<<<0,0>>>();
-#if !defined(__CUDA_ARCH__)
-  // expected-error at -3 {{call to global function g not configured}}
-#else
-  // expected-error at -5 {{no matching function for call to 'g'}}
-  // expected-error at -5 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif  // __CUDA_ARCH__
+  g(); // expected-error {{call to global function g not configured}}
 }
 
 // Test for address of overloaded function resolution in the global context.

Modified: cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu?rev=283963&r1=283962&r2=283963&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu (original)
+++ cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu Tue Oct 11 20:30:08 2016
@@ -1,5 +1,7 @@
-// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify -DDEVICE %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify \
+// RUN:   -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify \
+// RUN:   -verify-ignore-unexpected=note -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
@@ -10,17 +12,16 @@
 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}}
+  // This emits a deferred error on the device, but we don't catch it in this
+  // file because the non-deferred error below precludes this.
+
+  // FIXME-expected-error at -2 {{reference to __global__ function}}
 #endif
 }
 __host__ fn_ptr_t get_ptr_h() {