Index: cfe/trunk/include/clang/Basic/LangOptions.def
===================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def
+++ cfe/trunk/include/clang/Basic/LangOptions.def
@@ -171,6 +171,7 @@
 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")
Index: cfe/trunk/include/clang/Driver/CC1Options.td
===================================================================
--- cfe/trunk/include/clang/Driver/CC1Options.td
+++ cfe/trunk/include/clang/Driver/CC1Options.td
@@ -678,6 +678,8 @@
   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
Index: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
===================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp
@@ -1521,6 +1521,9 @@
   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();
Index: cfe/trunk/lib/Sema/SemaDecl.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp
+++ cfe/trunk/lib/Sema/SemaDecl.cpp
@@ -8290,9 +8290,11 @@
     }
 
     // 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);
Index: cfe/trunk/test/SemaCUDA/vararg.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/vararg.cu
+++ cfe/trunk/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 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
 }