summaryrefslogtreecommitdiffstats
path: root/clang
diff options
context:
space:
mode:
Diffstat (limited to 'clang')
-rw-r--r--clang/include/clang/Basic/LangOptions.def1
-rw-r--r--clang/include/clang/Driver/CC1Options.td3
-rw-r--r--clang/lib/Frontend/CompilerInvocation.cpp3
-rw-r--r--clang/lib/Sema/SemaCUDA.cpp5
-rw-r--r--clang/test/SemaCUDA/function-target-disabled-check.cu26
5 files changed, 38 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index f9c372c439b..bd90a2c2d74 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -162,6 +162,7 @@ LANGOPT(CUDA , 1, 0, "CUDA")
LANGOPT(OpenMP , 1, 0, "OpenMP support")
LANGOPT(CUDAIsDevice , 1, 0, "Compiling for CUDA device")
LANGOPT(CUDAAllowHostCallsFromHostDevice, 1, 0, "Allow host device functions to call host functions")
+LANGOPT(CUDADisableTargetCallChecks, 1, 0, "Disable checks for call targets (host, device, etc.)")
LANGOPT(AssumeSaneOperatorNew , 1, 1, "implicit __attribute__((malloc)) for C++'s new operators")
LANGOPT(SizedDeallocation , 1, 0, "enable sized deallocation functions")
diff --git a/clang/include/clang/Driver/CC1Options.td b/clang/include/clang/Driver/CC1Options.td
index b6b37c32d2e..12b4c655740 100644
--- a/clang/include/clang/Driver/CC1Options.td
+++ b/clang/include/clang/Driver/CC1Options.td
@@ -610,6 +610,9 @@ def fcuda_is_device : Flag<["-"], "fcuda-is-device">,
def fcuda_allow_host_calls_from_host_device : Flag<["-"],
"fcuda-allow-host-calls-from-host-device">,
HelpText<"Allow host device functions to call host functions">;
+def fcuda_disable_target_call_checks : Flag<["-"],
+ "fcuda-disable-target-call-checks">,
+ HelpText<"Disable all cross-target (host, device, etc.) call checks in CUDA">;
} // let Flags = [CC1Option]
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 968f212dda9..da1a088097e 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -1375,6 +1375,9 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
if (Args.hasArg(OPT_fcuda_allow_host_calls_from_host_device))
Opts.CUDAAllowHostCallsFromHostDevice = 1;
+ if (Args.hasArg(OPT_fcuda_disable_target_call_checks))
+ Opts.CUDADisableTargetCallChecks = 1;
+
if (Opts.ObjC1) {
if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) {
StringRef value = arg->getValue();
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 60338212c69..5973500826e 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -62,6 +62,11 @@ Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
+ // The CUDADisableTargetCallChecks short-circuits this check: we assume all
+ // cross-target calls are valid.
+ if (getLangOpts().CUDADisableTargetCallChecks)
+ return false;
+
CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
CalleeTarget = IdentifyCUDATarget(Callee);
diff --git a/clang/test/SemaCUDA/function-target-disabled-check.cu b/clang/test/SemaCUDA/function-target-disabled-check.cu
new file mode 100644
index 00000000000..979d4edbf89
--- /dev/null
+++ b/clang/test/SemaCUDA/function-target-disabled-check.cu
@@ -0,0 +1,26 @@
+// Test that we can disable cross-target call checks in Sema with the
+// -fcuda-disable-target-call-checks flag. Without this flag we'd get a bunch
+// of errors here, since there are invalid cross-target calls present.
+
+// RUN: %clang_cc1 -fsyntax-only -verify %s -fcuda-disable-target-call-checks
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -fcuda-disable-target-call-checks
+
+// expected-no-diagnostics
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+
+__attribute__((host)) void h1();
+
+__attribute__((device)) void d1() {
+ h1();
+}
+
+__attribute__((host)) void h2() {
+ d1();
+}
+
+__attribute__((global)) void g1() {
+ h2();
+}
OpenPOWER on IntegriCloud