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.
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283963
91177308-0d34-0410-b5e6-
96231b3b80d8
// (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.
--- /dev/null
+// 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__
+}
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.
-// 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
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() {