r264106 - [CUDA] Don't allow templated variadic functions.

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Tue Mar 22 17:06:19 2016
New Revision: 264106

URL: http://llvm.org/viewvc/llvm-project?rev=264106&view=rev
Log:
[CUDA] Don't allow templated variadic functions.

Reviewers: tra

Subscribers: cfe-commits

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

Modified:
    cfe/trunk/lib/Sema/SemaDecl.cpp
    cfe/trunk/test/SemaCUDA/vararg.cu

Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=264106&r1=264105&r2=264106&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Mar 22 17:06:19 2016
@@ -8341,6 +8341,28 @@ Sema::ActOnFunctionDeclarator(Scope *S,
         isExplicitSpecialization || isFunctionTemplateSpecialization);
   }
 
+  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, 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);
+    }
+  }
+
   if (getLangOpts().CPlusPlus) {
     if (FunctionTemplate) {
       if (NewFD->isInvalidDecl())
@@ -8390,28 +8412,6 @@ Sema::ActOnFunctionDeclarator(Scope *S,
 
   MarkUnusedFileScopedDecl(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, 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);
-    }
-  }
-
   // 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.

Modified: cfe/trunk/test/SemaCUDA/vararg.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/vararg.cu?rev=264106&r1=264105&r2=264106&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/vararg.cu (original)
+++ cfe/trunk/test/SemaCUDA/vararg.cu Tue Mar 22 17:06:19 2016
@@ -35,6 +35,12 @@ __device__ void vararg(const char* x, ..
 // expected-error at -2 {{CUDA device code does not support variadic functions}}
 #endif
 
+template <typename T>
+__device__ void vararg(T t, ...) {}
+#ifdef EXPECT_VARARG_ERR
+// expected-error at -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.