<p dir="ltr">Yes, I have yet to figure out when that REQUIRES comment is, erm, required.  Sorry.  :-/</p>
<p dir="ltr">I believe Art checked in a fix.</p>
<div class="gmail_extra"><br><div class="gmail_quote">On Oct 11, 2016 7:22 PM, "Chandler Carruth" <<a href="mailto:chandlerc@gmail.com">chandlerc@gmail.com</a>> wrote:<br type="attribution"><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div dir="ltr"><div class="gmail_quote"><div dir="ltr">On Tue, Oct 11, 2016 at 6:39 PM Justin Lebar via cfe-commits <<a href="mailto:cfe-commits@lists.llvm.org" target="_blank">cfe-commits@lists.llvm.org</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Author: jlebar<br class="m_-3704955923428731676gmail_msg"><br>Date: Tue Oct 11 20:30:08 2016<br class="m_-3704955923428731676gmail_msg"><br>New Revision: 283963<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>URL: <a href="http://llvm.org/viewvc/llvm-project?rev=283963&view=rev" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project?rev=283963&view=rev</a><br class="m_-3704955923428731676gmail_msg"><br>Log:<br class="m_-3704955923428731676gmail_msg"><br>[CUDA] Make touching a kernel from a __host__ __device__ function a deferred error.<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>Previously, this was an immediate, don't pass go, don't collect $200<br class="m_-3704955923428731676gmail_msg"><br>error.  But this precludes us from writing code like<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>  __host__ __device__ void launch_kernel() {<br class="m_-3704955923428731676gmail_msg"><br>    kernel<<<...>>>();<br class="m_-3704955923428731676gmail_msg"><br>  }<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>Such code isn't wrong, following our notions of right and wrong in CUDA,<br class="m_-3704955923428731676gmail_msg"><br>unless it's codegen'ed.<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>Added:<br class="m_-3704955923428731676gmail_msg"><br>    cfe/trunk/test/SemaCUDA/<a href="http://function-overload-hd.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">functi<wbr>on-overload-hd.cu</a></blockquote><div><br></div><div>This test fails for buildbots that don't configure the NVPTX backend:</div><div><a href="http://lab.llvm.org:8011/builders/clang-cmake-armv7-a15/builds/15830/steps/ninja%20check%201/logs/FAIL%3A%20Clang%3A%3Afunction-overload-hd.cu" target="_blank">http://lab.llvm.org:8011/<wbr>builders/clang-cmake-armv7-<wbr>a15/builds/15830/steps/ninja%<wbr>20check%201/logs/FAIL%3A%<wbr>20Clang%3A%3Afunction-<wbr>overload-hd.cu</a> </div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Modified:<br class="m_-3704955923428731676gmail_msg"><br>    cfe/trunk/lib/Sema/SemaCUDA.<wbr>cpp<br class="m_-3704955923428731676gmail_msg"><br>    cfe/trunk/test/SemaCUDA/<a href="http://function-overload.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">functi<wbr>on-overload.cu</a><br class="m_-3704955923428731676gmail_msg"><br>    cfe/trunk/test/SemaCUDA/<a href="http://reference-to-kernel-fn.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">refere<wbr>nce-to-kernel-fn.cu</a><br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>Modified: cfe/trunk/lib/Sema/SemaCUDA.<wbr>cpp<br class="m_-3704955923428731676gmail_msg"><br>URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=283963&r1=283962&r2=283963&view=diff" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project/cfe/trunk/lib/Sema/<wbr>SemaCUDA.cpp?rev=283963&r1=<wbr>283962&r2=283963&view=diff</a><br class="m_-3704955923428731676gmail_msg"><br>==============================<wbr>==============================<wbr>==================<br class="m_-3704955923428731676gmail_msg"><br>--- cfe/trunk/lib/Sema/SemaCUDA.<wbr>cpp (original)<br class="m_-3704955923428731676gmail_msg"><br>+++ cfe/trunk/lib/Sema/SemaCUDA.<wbr>cpp Tue Oct 11 20:30:08 2016<br class="m_-3704955923428731676gmail_msg"><br>@@ -120,8 +120,7 @@ Sema::IdentifyCUDAPreference(<wbr>const Funct<br class="m_-3704955923428731676gmail_msg"><br>   // (a) Can't call global from some contexts until we support CUDA's<br class="m_-3704955923428731676gmail_msg"><br>   // dynamic parallelism.<br class="m_-3704955923428731676gmail_msg"><br>   if (CalleeTarget == CFT_Global &&<br class="m_-3704955923428731676gmail_msg"><br>-      (CallerTarget == CFT_Global || CallerTarget == CFT_Device ||<br class="m_-3704955923428731676gmail_msg"><br>-       (CallerTarget == CFT_HostDevice && getLangOpts().CUDAIsDevice)))<br class="m_-3704955923428731676gmail_msg"><br>+      (CallerTarget == CFT_Global || CallerTarget == CFT_Device))<br class="m_-3704955923428731676gmail_msg"><br>     return CFP_Never;<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>   // (b) Calling HostDevice is OK for everyone.<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>Added: cfe/trunk/test/SemaCUDA/<a href="http://function-overload-hd.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">functi<wbr>on-overload-hd.cu</a><br class="m_-3704955923428731676gmail_msg"><br>URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload-hd.cu?rev=283963&view=auto" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project/cfe/trunk/test/<wbr>SemaCUDA/function-overload-hd.<wbr>cu?rev=283963&view=auto</a><br class="m_-3704955923428731676gmail_msg"><br>==============================<wbr>==============================<wbr>==================<br class="m_-3704955923428731676gmail_msg"><br>--- cfe/trunk/test/SemaCUDA/<a href="http://function-overload-hd.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">functi<wbr>on-overload-hd.cu</a> (added)<br class="m_-3704955923428731676gmail_msg"><br>+++ cfe/trunk/test/SemaCUDA/<a href="http://function-overload-hd.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">functi<wbr>on-overload-hd.cu</a> Tue Oct 11 20:30:08 2016<br class="m_-3704955923428731676gmail_msg"><br>@@ -0,0 +1,28 @@<br class="m_-3704955923428731676gmail_msg"><br>+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -S -o /dev/null -verify \<br class="m_-3704955923428731676gmail_msg"><br>+// RUN:   -verify-ignore-unexpected=<wbr>note %s<br class="m_-3704955923428731676gmail_msg"><br>+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -S -o /dev/null -fcuda-is-device \<br class="m_-3704955923428731676gmail_msg"><br>+// RUN:   -verify -verify-ignore-unexpected=note %s<br class="m_-3704955923428731676gmail_msg"><br>+<br class="m_-3704955923428731676gmail_msg"><br>+#include "Inputs/cuda.h"<br class="m_-3704955923428731676gmail_msg"><br>+<br class="m_-3704955923428731676gmail_msg"><br>+// FIXME: Merge into <a href="http://function-overload.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">function-overload.cu</a> once deferred errors can be emitted<br class="m_-3704955923428731676gmail_msg"><br>+// when non-deferred errors are present.<br class="m_-3704955923428731676gmail_msg"><br>+<br class="m_-3704955923428731676gmail_msg"><br>+#if !defined(__CUDA_ARCH__)<br class="m_-3704955923428731676gmail_msg"><br>+//expected-no-diagnostics<br class="m_-3704955923428731676gmail_msg"><br>+#endif<br class="m_-3704955923428731676gmail_msg"><br>+<br class="m_-3704955923428731676gmail_msg"><br>+typedef void (*GlobalFnPtr)();  // __global__ functions must return void.<br class="m_-3704955923428731676gmail_msg"><br>+<br class="m_-3704955923428731676gmail_msg"><br>+__global__ void g() {}<br class="m_-3704955923428731676gmail_msg"><br>+<br class="m_-3704955923428731676gmail_msg"><br>+__host__ __device__ void hd() {<br class="m_-3704955923428731676gmail_msg"><br>+  GlobalFnPtr fp_g = g;<br class="m_-3704955923428731676gmail_msg"><br>+#if defined(__CUDA_ARCH__)<br class="m_-3704955923428731676gmail_msg"><br>+  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}<br class="m_-3704955923428731676gmail_msg"><br>+#endif<br class="m_-3704955923428731676gmail_msg"><br>+  g<<<0,0>>>();<br class="m_-3704955923428731676gmail_msg"><br>+#if defined(__CUDA_ARCH__)<br class="m_-3704955923428731676gmail_msg"><br>+  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}<br class="m_-3704955923428731676gmail_msg"><br>+#endif  // __CUDA_ARCH__<br class="m_-3704955923428731676gmail_msg"><br>+}<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>Modified: cfe/trunk/test/SemaCUDA/<a href="http://function-overload.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">functi<wbr>on-overload.cu</a><br class="m_-3704955923428731676gmail_msg"><br>URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=283963&r1=283962&r2=283963&view=diff" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project/cfe/trunk/test/<wbr>SemaCUDA/function-overload.cu?<wbr>rev=283963&r1=283962&r2=<wbr>283963&view=diff</a><br class="m_-3704955923428731676gmail_msg"><br>==============================<wbr>==============================<wbr>==================<br class="m_-3704955923428731676gmail_msg"><br>--- cfe/trunk/test/SemaCUDA/<a href="http://function-overload.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">functi<wbr>on-overload.cu</a> (original)<br class="m_-3704955923428731676gmail_msg"><br>+++ cfe/trunk/test/SemaCUDA/<a href="http://function-overload.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">functi<wbr>on-overload.cu</a> Tue Oct 11 20:30:08 2016<br class="m_-3704955923428731676gmail_msg"><br>@@ -181,18 +181,7 @@ __host__ __device__ void hostdevicef() {<br class="m_-3704955923428731676gmail_msg"><br>   CurrentFnPtr fp_cdh = cdh;<br class="m_-3704955923428731676gmail_msg"><br>   CurrentReturnTy ret_cdh = cdh();<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>-  GlobalFnPtr fp_g = g;<br class="m_-3704955923428731676gmail_msg"><br>-#if defined(__CUDA_ARCH__)<br class="m_-3704955923428731676gmail_msg"><br>-  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}<br class="m_-3704955923428731676gmail_msg"><br>-#endif<br class="m_-3704955923428731676gmail_msg"><br>-  g();<br class="m_-3704955923428731676gmail_msg"><br>-  g<<<0,0>>>();<br class="m_-3704955923428731676gmail_msg"><br>-#if !defined(__CUDA_ARCH__)<br class="m_-3704955923428731676gmail_msg"><br>-  // expected-error@-3 {{call to global function g not configured}}<br class="m_-3704955923428731676gmail_msg"><br>-#else<br class="m_-3704955923428731676gmail_msg"><br>-  // expected-error@-5 {{no matching function for call to 'g'}}<br class="m_-3704955923428731676gmail_msg"><br>-  // expected-error@-5 {{reference to __global__ function 'g' in __host__ __device__ function}}<br class="m_-3704955923428731676gmail_msg"><br>-#endif  // __CUDA_ARCH__<br class="m_-3704955923428731676gmail_msg"><br>+  g(); // expected-error {{call to global function g not configured}}<br class="m_-3704955923428731676gmail_msg"><br> }<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br> // Test for address of overloaded function resolution in the global context.<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>Modified: cfe/trunk/test/SemaCUDA/<a href="http://reference-to-kernel-fn.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">refere<wbr>nce-to-kernel-fn.cu</a><br class="m_-3704955923428731676gmail_msg"><br>URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/reference-to-kernel-fn.cu?rev=283963&r1=283962&r2=283963&view=diff" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project/cfe/trunk/test/<wbr>SemaCUDA/reference-to-kernel-<wbr>fn.cu?rev=283963&r1=283962&r2=<wbr>283963&view=diff</a><br class="m_-3704955923428731676gmail_msg"><br>==============================<wbr>==============================<wbr>==================<br class="m_-3704955923428731676gmail_msg"><br>--- cfe/trunk/test/SemaCUDA/<a href="http://reference-to-kernel-fn.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">refere<wbr>nce-to-kernel-fn.cu</a> (original)<br class="m_-3704955923428731676gmail_msg"><br>+++ cfe/trunk/test/SemaCUDA/<a href="http://reference-to-kernel-fn.cu" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">refere<wbr>nce-to-kernel-fn.cu</a> Tue Oct 11 20:30:08 2016<br class="m_-3704955923428731676gmail_msg"><br>@@ -1,5 +1,7 @@<br class="m_-3704955923428731676gmail_msg"><br>-// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s<br class="m_-3704955923428731676gmail_msg"><br>-// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify -DDEVICE %s<br class="m_-3704955923428731676gmail_msg"><br>+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify \<br class="m_-3704955923428731676gmail_msg"><br>+// RUN:   -verify-ignore-unexpected=<wbr>note %s<br class="m_-3704955923428731676gmail_msg"><br>+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify \<br class="m_-3704955923428731676gmail_msg"><br>+// RUN:   -verify-ignore-unexpected=<wbr>note -DDEVICE %s<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br> // Check that we can reference (get a function pointer to) a __global__<br class="m_-3704955923428731676gmail_msg"><br> // function from the host side, but not the device side.  (We don't yet support<br class="m_-3704955923428731676gmail_msg"><br>@@ -10,17 +12,16 @@<br class="m_-3704955923428731676gmail_msg"><br> struct Dummy {};<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br> __global__ void kernel() {}<br class="m_-3704955923428731676gmail_msg"><br>-// expected-note@-1 {{declared here}}<br class="m_-3704955923428731676gmail_msg"><br>-#ifdef DEVICE<br class="m_-3704955923428731676gmail_msg"><br>-// expected-note@-3 {{declared here}}<br class="m_-3704955923428731676gmail_msg"><br>-#endif<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br> typedef void (*fn_ptr_t)();<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br> __host__ __device__ fn_ptr_t get_ptr_hd() {<br class="m_-3704955923428731676gmail_msg"><br>   return kernel;<br class="m_-3704955923428731676gmail_msg"><br> #ifdef DEVICE<br class="m_-3704955923428731676gmail_msg"><br>-  // expected-error@-2 {{reference to __global__ function}}<br class="m_-3704955923428731676gmail_msg"><br>+  // This emits a deferred error on the device, but we don't catch it in this<br class="m_-3704955923428731676gmail_msg"><br>+  // file because the non-deferred error below precludes this.<br class="m_-3704955923428731676gmail_msg"><br>+<br class="m_-3704955923428731676gmail_msg"><br>+  // FIXME-expected-error@-2 {{reference to __global__ function}}<br class="m_-3704955923428731676gmail_msg"><br> #endif<br class="m_-3704955923428731676gmail_msg"><br> }<br class="m_-3704955923428731676gmail_msg"><br> __host__ fn_ptr_t get_ptr_h() {<br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br><br class="m_-3704955923428731676gmail_msg"><br>______________________________<wbr>_________________<br class="m_-3704955923428731676gmail_msg"><br>cfe-commits mailing list<br class="m_-3704955923428731676gmail_msg"><br><a href="mailto:cfe-commits@lists.llvm.org" class="m_-3704955923428731676gmail_msg" target="_blank">cfe-commits@lists.llvm.org</a><br class="m_-3704955923428731676gmail_msg"><br><a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits" rel="noreferrer" class="m_-3704955923428731676gmail_msg" target="_blank">http://lists.llvm.org/cgi-bin/<wbr>mailman/listinfo/cfe-commits</a><br class="m_-3704955923428731676gmail_msg"><br></blockquote></div></div>
</blockquote></div></div>