diff options
| -rw-r--r-- | clang/include/clang/Basic/LangOptions.def | 1 | ||||
| -rw-r--r-- | clang/include/clang/Driver/CC1Options.td | 2 | ||||
| -rw-r--r-- | clang/lib/Frontend/CompilerInvocation.cpp | 3 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaDecl.cpp | 8 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/vararg.cu | 21 |
5 files changed, 26 insertions, 9 deletions
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index cc70d6246c0..cd1eb9f9d97 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -171,6 +171,7 @@ 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(CUDATargetOverloads, 1, 0, "Enable function overloads based on CUDA target attributes") +LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "Allow variadic functions in CUDA device code") 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 837fbff801b..2a966d856dd 100644 --- a/clang/include/clang/Driver/CC1Options.td +++ b/clang/include/clang/Driver/CC1Options.td @@ -678,6 +678,8 @@ def fcuda_include_gpubinary : Separate<["-"], "fcuda-include-gpubinary">, HelpText<"Incorporate CUDA device-side binary into host object file.">; def fcuda_target_overloads : Flag<["-"], "fcuda-target-overloads">, HelpText<"Enable function overloads based on CUDA target attributes.">; +def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">, + HelpText<"Allow variadic functions in CUDA device code.">; //===----------------------------------------------------------------------===// // OpenMP Options diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index d0b8b27b052..a202996a79a 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -1521,6 +1521,9 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, if (Args.hasArg(OPT_fcuda_target_overloads)) Opts.CUDATargetOverloads = 1; + if (Args.hasArg(OPT_fcuda_allow_variadic_functions)) + Opts.CUDAAllowVariadicFunctions = 1; + if (Opts.ObjC1) { if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) { StringRef value = arg->getValue(); diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index e3365b559a7..166bd8ca24e 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -8290,9 +8290,11 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC, } // Variadic functions, other than a *declaration* of printf, are not allowed - // in device-side CUDA code. - if (NewFD->isVariadic() && (NewFD->hasAttr<CUDADeviceAttr>() || - NewFD->hasAttr<CUDAGlobalAttr>()) && + // in device-side CUDA code, unless someone passed + // -fcuda-allow-variadic-functions. + if (!getLangOpts().CUDAAllowVariadicFunctions && NewFD->isVariadic() && + (NewFD->hasAttr<CUDADeviceAttr>() || + NewFD->hasAttr<CUDAGlobalAttr>()) && !(II && II->isStr("printf") && NewFD->isExternC() && !D.isFunctionDefinition())) { Diag(NewFD->getLocation(), diag::err_variadic_device_fn); diff --git a/clang/test/SemaCUDA/vararg.cu b/clang/test/SemaCUDA/vararg.cu index f3fe2637607..c14d4891453 100644 --- a/clang/test/SemaCUDA/vararg.cu +++ b/clang/test/SemaCUDA/vararg.cu @@ -1,8 +1,11 @@ // REQUIRES: x86-registered-target // REQUIRES: nvptx-registered-target // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \ -// RUN: -verify -DEXPECT_ERR %s -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s +// RUN: -verify -DEXPECT_VA_ARG_ERR -DEXPECT_VARARG_ERR %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \ +// RUN: -fcuda-allow-variadic-functions -verify -DEXPECT_VA_ARG_ERR %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify \ +// RUN: -DEXPECT_VARARG_ERR %s #include <stdarg.h> #include "Inputs/cuda.h" @@ -10,7 +13,7 @@ __device__ void foo() { va_list list; va_arg(list, int); -#ifdef EXPECT_ERR +#ifdef EXPECT_VA_ARG_ERR // expected-error@-2 {{CUDA device code does not support va_arg}} #endif } @@ -28,15 +31,21 @@ __device__ void baz() { } __device__ void vararg(const char* x, ...) {} -// expected-error@-1 {{CUDA device code does not support variadic functions}} +#ifdef EXPECT_VARARG_ERR +// expected-error@-2 {{CUDA device code does not support variadic functions}} +#endif extern "C" __device__ int printf(const char* fmt, ...); // OK, special case. // Definition of printf not allowed. extern "C" __device__ int printf(const char* fmt, ...) { return 0; } -// expected-error@-1 {{CUDA device code does not support variadic functions}} +#ifdef EXPECT_VARARG_ERR +// expected-error@-2 {{CUDA device code does not support variadic functions}} +#endif namespace ns { __device__ int printf(const char* fmt, ...); -// expected-error@-1 {{CUDA device code does not support variadic functions}} +#ifdef EXPECT_VARARG_ERR +// expected-error@-2 {{CUDA device code does not support variadic functions}} +#endif } |
