This revision was automatically updated to reflect the committed changes.
Closed by commit rL258643: [CUDA] Disallow variadic functions other than printf
in device code. (authored by jlebar).
Changed prior to commit:
http://reviews.llvm.org/D16484?vs=45736&id=45811#toc
Repository:
rL LLVM
http://reviews.llvm.org/D16484
Files:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/SemaCUDA/va-arg.cu
cfe/trunk/test/SemaCUDA/vararg.cu
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6423,6 +6423,8 @@
def warn_kern_is_inline : Warning<
"ignored 'inline' attribute on kernel function %0">,
InGroup<CudaCompat>;
+def err_variadic_device_fn : Error<
+ "CUDA device code does not support variadic functions">;
def err_va_arg_in_device : Error<
"CUDA device code does not support va_arg">;
def err_alias_not_supported_on_nvptx : Error<"CUDA does not support aliases">;
Index: cfe/trunk/test/SemaCUDA/vararg.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/vararg.cu
+++ cfe/trunk/test/SemaCUDA/vararg.cu
@@ -0,0 +1,42 @@
+// 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
+
+#include <stdarg.h>
+#include "Inputs/cuda.h"
+
+__device__ void foo() {
+ va_list list;
+ va_arg(list, int);
+#ifdef EXPECT_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}}
+
+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}}
+
+namespace ns {
+__device__ int printf(const char* fmt, ...);
+// expected-error@-1 {{CUDA device code does not support variadic functions}}
+}
Index: cfe/trunk/test/SemaCUDA/va-arg.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/va-arg.cu
+++ cfe/trunk/test/SemaCUDA/va-arg.cu
@@ -1,28 +0,0 @@
-// 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 %s
-
-#include <stdarg.h>
-#include "Inputs/cuda.h"
-
-__device__ void foo() {
- va_list list;
- va_arg(list, int);
-#ifdef EXPECT_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
-}
Index: cfe/trunk/lib/Sema/SemaDecl.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp
+++ cfe/trunk/lib/Sema/SemaDecl.cpp
@@ -8276,18 +8276,26 @@
MarkUnusedFileScopedDecl(NewFD);
- if (getLangOpts().CUDA)
- if (IdentifierInfo *II = NewFD->getIdentifier())
- if (!NewFD->isInvalidDecl() &&
- NewFD->getDeclContext()->getRedeclContext()->isTranslationUnit()) {
- if (II->isStr("cudaConfigureCall")) {
- if (!R->getAs<FunctionType>()->getReturnType()->isScalarType())
- Diag(NewFD->getLocation(), diag::err_config_scalar_return);
+ if (getLangOpts().CUDA) {
+ IdentifierInfo *II = NewFD->getIdentifier();
+ 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>()) &&
+ !(II && II->isStr("printf") && NewFD->isExternC() &&
+ !D.isFunctionDefinition())) {
+ Diag(NewFD->getLocation(), diag::err_variadic_device_fn);
+ }
+ }
- Context.setcudaConfigureCallDecl(NewFD);
- }
- }
-
// Here we have an function template explicit specialization at class scope.
// The actually specialization will be postponed to template instatiation
// time via the ClassScopeFunctionSpecializationDecl node.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
