From: Justin Lebar <jlebar@google.com>
Date: Sat, 23 Jan 2016 21:28:17 +0000 (+0000)
Subject: [CUDA] Disallow variadic functions other than printf in device code.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=5496a8c53db4003c03e7e1341e4ec9b67e54563a;p=clang

[CUDA] Disallow variadic functions other than printf in device code.

Reviewers: tra

Subscribers: cfe-commits, echristo, jhen

Differential Revision: http://reviews.llvm.org/D16484

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@258643 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index e9e877eacf..180d3d18db 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6423,6 +6423,8 @@ def warn_kern_is_method : Extension<
 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">;
diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index 16ad13db90..6adbc3fab6 100644
--- a/lib/Sema/SemaDecl.cpp
+++ b/lib/Sema/SemaDecl.cpp
@@ -8276,18 +8276,26 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
 
   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);
-
-          Context.setcudaConfigureCallDecl(NewFD);
-        }
-      }
-  
+  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);
+    }
+  }
+
   // 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.
diff --git a/test/SemaCUDA/va-arg.cu b/test/SemaCUDA/vararg.cu
similarity index 55%
rename from test/SemaCUDA/va-arg.cu
rename to test/SemaCUDA/vararg.cu
index aad9b74c2d..f3fe263760 100644
--- a/test/SemaCUDA/va-arg.cu
+++ b/test/SemaCUDA/vararg.cu
@@ -2,7 +2,7 @@
 // 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
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
 
 #include <stdarg.h>
 #include "Inputs/cuda.h"
@@ -26,3 +26,17 @@ __device__ void baz() {
   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}}
+}