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

[Justin Lebar via cfe-commits]

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" <chandlerc at gmail.com> wrote:

> On Tue, Oct 11, 2016 at 6:39 PM Justin Lebar via cfe-commits <
> cfe-commits at 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 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() {
>>
>>
>>
>>
>>
>> _______________________________________________
>>
>> cfe-commits mailing list
>>
>> cfe-commits at lists.llvm.org
>>
>> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
>>
>>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20161011/fd32cbc7/attachment.html>