Yes, I have yet to figure out when that REQUIRES comment is, erm, required. Sorry. :-/
I believe Art checked in a fix. On Oct 11, 2016 7:22 PM, "Chandler Carruth" <chandl...@gmail.com> wrote: > On Tue, Oct 11, 2016 at 6:39 PM Justin Lebar via cfe-commits < > cfe-commits@lists.llvm.org> wrote: > >> 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 > > > This test fails for buildbots that don't configure the NVPTX backend: > 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 > > >> 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@-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__ >> >> +} >> >> >> >> 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@-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. >> >> >> >> 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@-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() { >> >> >> >> >> >> _______________________________________________ >> >> cfe-commits mailing list >> >> cfe-commits@lists.llvm.org >> >> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits >> >>
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
