[CUDA] Bail, rather than crash, on va_arg in device code.

Reviewers: tra

Subscribers: echristo, jhen, cfe-commits

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

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

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 6e652d782ff8ff099fecd015dc23ed3d24a6a20a..a276e4bf4176f321b2e967e93dea32bbfa9b8601 100644 (file)
--- 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_va_arg_in_device : Error<
+  "CUDA device code does not support va_arg">;
 
 def warn_non_pod_vararg_with_format_string : Warning<
   "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "

diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp
index f8537d3692b03b4ad64fad3c199486451551fbb2..008a2fc4c456dbd589d0b0a27d48e521a62592d2 100644 (file)
--- a/lib/Sema/SemaExpr.cpp
+++ b/lib/Sema/SemaExpr.cpp
@@ -11715,9 +11715,8 @@ ExprResult Sema::ActOnBlockStmtExpr(SourceLocation CaretLoc,
   return Result;
 }
 
-ExprResult Sema::ActOnVAArg(SourceLocation BuiltinLoc,
-                                        Expr *E, ParsedType Ty,
-                                        SourceLocation RPLoc) {
+ExprResult Sema::ActOnVAArg(SourceLocation BuiltinLoc, Expr *E, ParsedType Ty,
+                            SourceLocation RPLoc) {
   TypeSourceInfo *TInfo;
   GetTypeFromParser(Ty, &TInfo);
   return BuildVAArgExpr(BuiltinLoc, E, TInfo, RPLoc);
@@ -11729,6 +11728,15 @@ ExprResult Sema::BuildVAArgExpr(SourceLocation BuiltinLoc,
   Expr *OrigExpr = E;
   bool IsMS = false;
 
+  // CUDA device code does not support varargs.
+  if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
+    if (const FunctionDecl *F = dyn_cast<FunctionDecl>(CurContext)) {
+      CUDAFunctionTarget T = IdentifyCUDATarget(F);
+      if (T == CFT_Global || T == CFT_Device || T == CFT_HostDevice)
+        return ExprError(Diag(E->getLocStart(), diag::err_va_arg_in_device));
+    }
+  }
+
   // It might be a __builtin_ms_va_list. (But don't ever mark a va_arg()
   // as Microsoft ABI on an actual Microsoft platform, where
   // __builtin_ms_va_list and __builtin_va_list are the same.)

diff --git a/test/SemaCUDA/va-arg.cu b/test/SemaCUDA/va-arg.cu
new file mode 100644 (file)
index 0000000..aad9b74
--- /dev/null
+++ b/test/SemaCUDA/va-arg.cu
@@ -0,0 +1,28 @@
+// 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
+}