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

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Tue Jan 19 18:27:00 2016
New Revision: 258264

URL: http://llvm.org/viewvc/llvm-project?rev=258264&view=rev
Log:
[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

Added:
    cfe/trunk/test/SemaCUDA/va-arg.cu
Modified:
    cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
    cfe/trunk/lib/Sema/SemaExpr.cpp

Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=258264&r1=258263&r2=258264&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Tue Jan 19 18:27:00 2016
@@ -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 "

Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=258264&r1=258263&r2=258264&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Tue Jan 19 18:27:00 2016
@@ -11715,9 +11715,8 @@ ExprResult Sema::ActOnBlockStmtExpr(Sour
   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(SourceLo
   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.)

Added: cfe/trunk/test/SemaCUDA/va-arg.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/va-arg.cu?rev=258264&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/va-arg.cu (added)
+++ cfe/trunk/test/SemaCUDA/va-arg.cu Tue Jan 19 18:27:00 2016
@@ -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 at -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
+}