[CUDA] Add -fcuda-allow-variadic-functions.

Summary:
Turns out the variadic function checking added in r258643 was too strict
for some existing users; give them an escape valve.  When
-fcuda-allow-variadic-functions is passed, the front-end makes no
attempt to disallow C-style variadic functions.  Calls to va_arg are
still not allowed.

Reviewers: tra

Subscribers: cfe-commits, jhen, echristo, bkramer

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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@258822 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/LangOptions.def b/include/clang/Basic/LangOptions.def
index cc70d6246c03603f596cc882809e08357a7e3c00..cd1eb9f9d978b620b930ab2ce58a75f634eaefce 100644 (file)
--- a/include/clang/Basic/LangOptions.def
+++ b/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/include/clang/Driver/CC1Options.td b/include/clang/Driver/CC1Options.td
index 837fbff801b951ccc2363101c996a8a4bbc0e8bd..2a966d856dde6212d2f032bbf7f6efdc3ed144e4 100644 (file)
--- a/include/clang/Driver/CC1Options.td
+++ b/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/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp
index d0b8b27b05252358110e24fc75f516ef0d455b90..a202996a79acaf571ea4e2e2767fc97225eda9e5 100644 (file)
--- a/lib/Frontend/CompilerInvocation.cpp
+++ b/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/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index e3365b559a774f3ebb8766f4c1b342f75176733a..166bd8ca24e5ff9810175355be8e844e262a0bbc 100644 (file)
--- a/lib/Sema/SemaDecl.cpp
+++ b/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/test/SemaCUDA/vararg.cu b/test/SemaCUDA/vararg.cu
index f3fe26376076062eca11ee44c1ce486e1d6540b4..c14d4891453abd316d6012443d7f90fa69dc1729 100644 (file)
--- a/test/SemaCUDA/vararg.cu
+++ b/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
 }