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
