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

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Tue Jan 26 11:47:20 2016
New Revision: 258822

URL: http://llvm.org/viewvc/llvm-project?rev=258822&view=rev
Log:
[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

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/SemaDecl.cpp
    cfe/trunk/test/SemaCUDA/vararg.cu

Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=258822&r1=258821&r2=258822&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Tue Jan 26 11:47:20 2016
@@ -171,6 +171,7 @@ LANGOPT(CUDAIsDevice      , 1, 0, "Compi
 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")

Modified: cfe/trunk/include/clang/Driver/CC1Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/CC1Options.td?rev=258822&r1=258821&r2=258822&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/CC1Options.td (original)
+++ cfe/trunk/include/clang/Driver/CC1Options.td Tue Jan 26 11:47:20 2016
@@ -678,6 +678,8 @@ def fcuda_include_gpubinary : Separate<[
   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

Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=258822&r1=258821&r2=258822&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Tue Jan 26 11:47:20 2016
@@ -1521,6 +1521,9 @@ static void ParseLangArgs(LangOptions &O
   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();

Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=258822&r1=258821&r2=258822&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Jan 26 11:47:20 2016
@@ -8290,9 +8290,11 @@ Sema::ActOnFunctionDeclarator(Scope *S,
     }
 
     // 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);

Modified: cfe/trunk/test/SemaCUDA/vararg.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/vararg.cu?rev=258822&r1=258821&r2=258822&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/vararg.cu (original)
+++ cfe/trunk/test/SemaCUDA/vararg.cu Tue Jan 26 11:47:20 2016
@@ -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 at -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 at -1 {{CUDA device code does not support variadic functions}}
+#ifdef EXPECT_VARARG_ERR
+// expected-error at -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 at -1 {{CUDA device code does not support variadic functions}}
+#ifdef EXPECT_VARARG_ERR
+// expected-error at -2 {{CUDA device code does not support variadic functions}}
+#endif
 
 namespace ns {
 __device__ int printf(const char* fmt, ...);
-// expected-error at -1 {{CUDA device code does not support variadic functions}}
+#ifdef EXPECT_VARARG_ERR
+// expected-error at -2 {{CUDA device code does not support variadic functions}}
+#endif
 }