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(); +} _______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
