Hi rnk,
nvcc allows host device functions to call host functions with only a warning
being produced (host device functions calling device functions is an error in
nvcc). This nvcc feature (calling host functions from host device functions) is
used by some existing GPU code. Add an option to clang to allow similar
behavior. This does not affect code generation and trying to call a host
function from the GPU is still an error. We are investigating a more complete
solution that would avoid this but this is a first step to allow tools
analyzing GPU code to accept the same code as nvcc does.
http://reviews.llvm.org/D7841
Files:
include/clang/Basic/DiagnosticSemaKinds.td
include/clang/Basic/LangOptions.def
include/clang/Driver/CC1Options.td
lib/Frontend/CompilerInvocation.cpp
lib/Sema/SemaCUDA.cpp
test/SemaCUDA/function-target.cu
EMAIL PREFERENCES
http://reviews.llvm.org/settings/panel/emailpreferences/
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -6067,6 +6067,9 @@
def err_ref_bad_target : Error<
"reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
"function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">;
+def warn_host_device_function_calling_host_function : Warning<
+ "calling __host__ function %0 from __host__ __device__ function %1 can lead to runtime errors">,
+ InGroup<CudaCompat>;
def warn_non_pod_vararg_with_format_string : Warning<
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
Index: include/clang/Basic/LangOptions.def
===================================================================
--- include/clang/Basic/LangOptions.def
+++ include/clang/Basic/LangOptions.def
@@ -160,6 +160,7 @@
LANGOPT(CUDA , 1, 0, "CUDA")
LANGOPT(OpenMP , 1, 0, "OpenMP support")
LANGOPT(CUDAIsDevice , 1, 0, "Compiling for CUDA device")
+LANGOPT(CUDAHostDeviceFunctionsCallingHostFunctions, 1, 0, "Allow host device functions to call host functions")
LANGOPT(AssumeSaneOperatorNew , 1, 1, "implicit __attribute__((malloc)) for C++'s new operators")
LANGOPT(SizedDeallocation , 1, 0, "enable sized deallocation functions")
Index: include/clang/Driver/CC1Options.td
===================================================================
--- include/clang/Driver/CC1Options.td
+++ include/clang/Driver/CC1Options.td
@@ -608,6 +608,9 @@
def fcuda_is_device : Flag<["-"], "fcuda-is-device">,
HelpText<"Generate code for CUDA device">;
+def fcuda_host_device_functions_calling_host_functions : Flag<["-"],
+ "fcuda-host-device-functions-calling-host-functions">,
+ HelpText<"Allow host device functions to call host functions">;
} // let Flags = [CC1Option]
Index: lib/Frontend/CompilerInvocation.cpp
===================================================================
--- lib/Frontend/CompilerInvocation.cpp
+++ lib/Frontend/CompilerInvocation.cpp
@@ -253,7 +253,7 @@
for (unsigned i = 0, e = checkers.size(); i != e; ++i)
Opts.CheckersControlList.push_back(std::make_pair(checkers[i], enable));
}
-
+
// Go through the analyzer configuration options.
for (arg_iterator it = Args.filtered_begin(OPT_analyzer_config),
ie = Args.filtered_end(); it != ie; ++it) {
@@ -1393,6 +1393,9 @@
if (Args.hasArg(OPT_fcuda_is_device))
Opts.CUDAIsDevice = 1;
+ if (Args.hasArg(OPT_fcuda_host_device_functions_calling_host_functions))
+ Opts.CUDAHostDeviceFunctionsCallingHostFunctions = 1;
+
if (Opts.ObjC1) {
if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) {
StringRef value = arg->getValue();
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -92,9 +92,21 @@
if (Caller->isImplicit()) return false;
bool InDeviceMode = getLangOpts().CUDAIsDevice;
- if ((InDeviceMode && CalleeTarget != CFT_Device) ||
- (!InDeviceMode && CalleeTarget != CFT_Host))
+ if (!InDeviceMode && CalleeTarget != CFT_Host)
+ return true;
+ if (InDeviceMode && CalleeTarget != CFT_Device) {
+ // Allow host device functions to call host functions if explicitly
+ // requested.
+ if (CalleeTarget == CFT_Host &&
+ getLangOpts().CUDAHostDeviceFunctionsCallingHostFunctions) {
+ Diag(Caller->getLocation(),
+ diag::warn_host_device_function_calling_host_function)
+ << Callee->getNameAsString() << Caller->getNameAsString();
+ return false;
+ }
+
return true;
+ }
}
return false;
Index: test/SemaCUDA/function-target.cu
===================================================================
--- test/SemaCUDA/function-target.cu
+++ test/SemaCUDA/function-target.cu
@@ -1,5 +1,7 @@
-// RUN: %clang_cc1 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -fsyntax-only -verify %s -DTEST_HOST
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DTEST_DEVICE
+// RUN: %clang_cc1 -fsyntax-only -fcuda-host-device-functions-calling-host-functions -verify %s -DTEST_WARN_HD -DTEST_HOST
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-host-device-functions-calling-host-functions -verify %s -DTEST_WARN_HD -DTEST_DEVICE
#include "Inputs/cuda.h"
@@ -32,14 +34,21 @@
d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}}
}
-// Expected 0-1 as in one of host/device side compilation it is an error, while
-// not in the other
-__host__ void hd1h(void); // expected-note 0-1 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-__device__ void hd1d(void); // expected-note 0-1 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+__host__ void hd1h(void);
+#if defined(TEST_DEVICE) && !defined(TEST_WARN_HD)
+// expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
+__device__ void hd1d(void);
+#ifdef TEST_HOST
+// expected-note@-2 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#endif
__host__ void hd1hg(void);
__device__ void hd1dg(void);
#ifdef __CUDA_ARCH__
-__host__ void hd1hig(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+__host__ void hd1hig(void);
+#if defined(TEST_DEVICE) && !defined(TEST_WARN_HD)
+// expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
#else
__device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
#endif
@@ -47,10 +56,18 @@
__global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
__host__ __device__ void hd1(void) {
- // Expected 0-1 as in one of host/device side compilation it is an error,
- // while not in the other
- hd1d(); // expected-error 0-1 {{no matching function}}
- hd1h(); // expected-error 0-1 {{no matching function}}
+#if defined(TEST_WARN_HD) && defined(TEST_DEVICE)
+// expected-warning@-2 {{calling __host__ function hd1h from __host__ __device__ function hd1}}
+// expected-warning@-3 {{calling __host__ function hd1hig from __host__ __device__ function hd1}}
+#endif
+ hd1d();
+#ifdef TEST_HOST
+// expected-error@-2 {{no matching function}}
+#endif
+ hd1h();
+#if defined(TEST_DEVICE) && !defined(TEST_WARN_HD)
+// expected-error@-2 {{no matching function}}
+#endif
// No errors as guarded
#ifdef __CUDA_ARCH__
@@ -63,8 +80,11 @@
#ifndef __CUDA_ARCH__
hd1dig(); // expected-error {{no matching function}}
#else
- hd1hig(); // expected-error {{no matching function}}
+ hd1hig();
+#ifndef TEST_WARN_HD
+// expected-error@-2 {{no matching function}}
#endif
+#endif
hd1hd();
hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
