r235049 - Create a frontend flag to disable CUDA cross-target call checks

[Eli Bendersky]

Author: eliben
Date: Wed Apr 15 17:27:06 2015
New Revision: 235049

URL: http://llvm.org/viewvc/llvm-project?rev=235049&view=rev
Log:
Create a frontend flag to disable CUDA cross-target call checks

For CUDA source, Sema checks that the targets of call expressions make sense
(e.g. a host function can't call a device function).

Adding a flag that lets us skip this check. Motivation: for source-to-source
translation tools that have to accept code that's not strictly kosher CUDA but
is still accepted by nvcc. The source-to-source translation tool can then fix
the code and leave calls that are semantically valid for the actual compilation
stage.

Differential Revision: http://reviews.llvm.org/D9036


Added:
    cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu
Modified:
    cfe/trunk/include/clang/Basic/LangOptions.def
    cfe/trunk/include/clang/Driver/CC1Options.td
    cfe/trunk/lib/Frontend/CompilerInvocation.cpp
    cfe/trunk/lib/Sema/SemaCUDA.cpp

Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=235049&r1=235048&r2=235049&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Wed Apr 15 17:27:06 2015
@@ -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")

Modified: cfe/trunk/include/clang/Driver/CC1Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/CC1Options.td?rev=235049&r1=235048&r2=235049&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/CC1Options.td (original)
+++ cfe/trunk/include/clang/Driver/CC1Options.td Wed Apr 15 17:27:06 2015
@@ -610,6 +610,9 @@ def fcuda_is_device : Flag<["-"], "fcuda
 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]
 

Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=235049&r1=235048&r2=235049&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Wed Apr 15 17:27:06 2015
@@ -1375,6 +1375,9 @@ static void ParseLangArgs(LangOptions &O
   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();

Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=235049&r1=235048&r2=235049&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Apr 15 17:27:06 2015
@@ -62,6 +62,11 @@ Sema::CUDAFunctionTarget Sema::IdentifyC
 
 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);
 

Added: cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu?rev=235049&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu (added)
+++ cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu Wed Apr 15 17:27:06 2015
@@ -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();
+}