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

[Chandler Carruth via cfe-commits]

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/20161012/466e02b9/attachment-0001.html>