From f3d02c144a250d213f6e2882eea6d1e3ca5a995c Mon Sep 17 00:00:00 2001
From: Justin Lebar <jlebar@google.com>
Date: Wed, 12 Oct 2016 01:30:08 +0000
Subject: [PATCH] [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.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283963 91177308-0d34-0410-b5e6-96231b3b80d8
---
 lib/Sema/SemaCUDA.cpp                   |  3 +--
 test/SemaCUDA/function-overload-hd.cu   | 28 +++++++++++++++++++++++++
 test/SemaCUDA/function-overload.cu      | 13 +-----------
 test/SemaCUDA/reference-to-kernel-fn.cu | 15 ++++++-------
 4 files changed, 38 insertions(+), 21 deletions(-)
 create mode 100644 test/SemaCUDA/function-overload-hd.cu

diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp
index 9e101d16da..2a66124080 100644
--- a/lib/Sema/SemaCUDA.cpp
+++ b/lib/Sema/SemaCUDA.cpp
@@ -120,8 +120,7 @@ Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
   // (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.
diff --git a/test/SemaCUDA/function-overload-hd.cu b/test/SemaCUDA/function-overload-hd.cu
new file mode 100644
index 0000000000..76f5a2824a
--- /dev/null
+++ b/test/SemaCUDA/function-overload-hd.cu
@@ -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@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif
+  g<<<0,0>>>();
+#if defined(__CUDA_ARCH__)
+  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif  // __CUDA_ARCH__
+}
diff --git a/test/SemaCUDA/function-overload.cu b/test/SemaCUDA/function-overload.cu
index 6071dbb274..11e8bae126 100644
--- a/test/SemaCUDA/function-overload.cu
+++ b/test/SemaCUDA/function-overload.cu
@@ -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@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif
-  g();
-  g<<<0,0>>>();
-#if !defined(__CUDA_ARCH__)
-  // expected-error@-3 {{call to global function g not configured}}
-#else
-  // expected-error@-5 {{no matching function for call to 'g'}}
-  // expected-error@-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.
diff --git a/test/SemaCUDA/reference-to-kernel-fn.cu b/test/SemaCUDA/reference-to-kernel-fn.cu
index 29efcfa2fa..2d25dde94f 100644
--- a/test/SemaCUDA/reference-to-kernel-fn.cu
+++ b/test/SemaCUDA/reference-to-kernel-fn.cu
@@ -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@-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}}
+  // 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@-2 {{reference to __global__ function}}
 #endif
 }
 __host__ fn_ptr_t get_ptr_h() {
-- 
2.40.0