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

[jlebar]

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.

[bkramer]

Looks good to me, but I don't really now CUDA.

[tra]

LGTM.

[mcrosier]

cfe/trunk/include/clang/Driver/CC1Options.td

681 ↗

(On Diff #46005)

AFAICT, these are customer facing flags, correct? If so, shouldn't these options be hidden from the help (via the HelpHidden flag AFAICT)?

[mcrosier]

cfe/trunk/include/clang/Driver/CC1Options.td

681 ↗

(On Diff #46005)

This might cause issues such as those reported in PR26317 and PR26318.

[tra]

cfe/trunk/include/clang/Driver/CC1Options.td

681 ↗

(On Diff #46005)

These are cc1-only options and do *not* show up in top-level --help.

[mcrosier]

cfe/trunk/include/clang/Driver/CC1Options.td

681 ↗

(On Diff #46005)

Ah, yes. You are correct.

[garymm]

Could you please add this to the documentation?
Could this be made the default? It seems like nvcc does this by default.

[tra]

In D16559#4410067, @garymm wrote:

Could you please add this to the documentation?
Could this be made the default? It seems like nvcc does this by default.

Clang already does that, though we only allow variadic functions that don't actually use the vararg arguments: https://reviews.llvm.org/D151359
It's sufficient to compile recent CUDA/libcu++ headers w/o errors.

include/clang/Basic/LangOptions.def

 LANGOPT(OpenMP            , 1, 0, "OpenMP support")
 LANGOPT(OpenMPUseTLS      , 1, 0, "Use TLS for threadprivates or runtime calls")
 LANGOPT(OpenMPIsDevice    , 1, 0, "Generate code only for OpenMP target device")
 
 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")
 LANGOPT(ConceptsTS , 1, 0, "enable C++ Extensions for Concepts")
 BENIGN_LANGOPT(ElideConstructors , 1, 1, "C++ copy constructor elision")
 BENIGN_LANGOPT(DumpRecordLayouts , 1, 0, "dumping the layout of IRgen'd records")
 BENIGN_LANGOPT(DumpRecordLayoutsSimple , 1, 0, "dumping the layout of IRgen'd records in a simple form")
 BENIGN_LANGOPT(DumpVTableLayouts , 1, 0, "dumping the layouts of emitted vtables")

include/clang/Driver/CC1Options.td

def fcuda_allow_host_calls_from_host_device : Flag<["-"],
   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">;
 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
 //===----------------------------------------------------------------------===//
 
 def fopenmp_is_device : Flag<["-"], "fopenmp-is-device">,
   HelpText<"Generate code only for an OpenMP target device.">;
 def omp_host_ir_file_path : Separate<["-"], "omp-host-ir-file-path">,

lib/Frontend/CompilerInvocation.cpp

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 (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();
       if (Opts.ObjCRuntime.tryParse(value))
         Diags.Report(diag::err_drv_unknown_objc_runtime) << value;
     }
 
     if (Args.hasArg(OPT_fobjc_gc_only))

lib/Sema/SemaDecl.cpp

if (II && II->isStr("cudaConfigureCall") && !NewFD->isInvalidDecl() &&
         NewFD->getDeclContext()->getRedeclContext()->isTranslationUnit()) {
       if (!R->getAs<FunctionType>()->getReturnType()->isScalarType())
         Diag(NewFD->getLocation(), diag::err_config_scalar_return);
 
       Context.setcudaConfigureCallDecl(NewFD);
     }
 
     // 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>() ||
         !(II && II->isStr("printf") && NewFD->isExternC() &&
           !D.isFunctionDefinition())) {
       Diag(NewFD->getLocation(), diag::err_variadic_device_fn);
     }
   }
 
   // Here we have an function template explicit specialization at class scope.
   // The actually specialization will be postponed to template instatiation

test/SemaCUDA/vararg.cu

 // 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"
 
 __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
 }
 
 void bar() {
   va_list list;
   va_arg(list, int);  // OK: host-only
 }
 
 __device__ void baz() {
 #if !defined(__CUDA_ARCH__)
   va_list list;
   va_arg(list, int);  // OK: only seen when compiling for host
 #endif
 }
 
 __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
 }