This revision was automatically updated to reflect the committed changes.
Closed by commit rL258822: [CUDA] Add -fcuda-allow-variadic-functions.
(authored by jlebar).
Changed prior to commit:
http://reviews.llvm.org/D16559?vs=45915&id=46005#toc
Repository:
rL LLVM
http://reviews.llvm.org/D16559
Files:
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
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/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/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/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/test/SemaCUDA/vararg.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/vararg.cu
+++ cfe/trunk/test/SemaCUDA/vararg.cu
@@ -1,16 +1,19 @@
// 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
}
@@ -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
}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
