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
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
