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

[tra]

Small nit. LGTM otherwise.

lib/Sema/SemaDecl.cpp

8291–8293 ↗

(On Diff #45730)

It would be nice to emit this diagnostics on one side of compilation only.

Nit - unneeded {}

[jlebar]

Add check for extern "C" printf -- we don't want to take just any printf!

lib/Sema/SemaDecl.cpp

8291–8293 ↗

(On Diff #45730)

Resolved irl.

cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td

def warn_host_calls_from_host_device : Warning<
   "calling __host__ function %0 from __host__ __device__ function %1 can lead to runtime errors">,
   InGroup<CudaCompat>;
 def warn_kern_is_method : Extension<
   "kernel function %0 is a member function; this may not be accepted by nvcc">,
   InGroup<CudaCompat>;
 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">;
 
 def warn_non_pod_vararg_with_format_string : Warning<
   "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
   "%select{function|block|method|constructor}2; expected type from format "
   "string was %3">, InGroup<NonPODVarargs>, DefaultError;

cfe/trunk/lib/Sema/SemaDecl.cpp

if (getLangOpts().OpenCLVersion >= 200) {
             D.setInvalidType();
           }
       }
     }
   }
 
   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() &&
 
-          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.
   if (isDependentClassScopeExplicitSpecialization) {
     ClassScopeFunctionSpecializationDecl *NewSpec =
                          ClassScopeFunctionSpecializationDecl::Create(
                                 Context, CurContext, SourceLocation(), 
                                 cast<CXXMethodDecl>(NewFD),

cfe/trunk/test/SemaCUDA/va-arg.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 %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
-}

cfe/trunk/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
+
+#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}}
+}