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

Justin Lebar via cfe-commits Tue, 19 Jan 2016 14:22:56 -0800

jlebar created this revision.
jlebar added a reviewer: tra.
jlebar added subscribers: cfe-commits, jhen, echristo.


http://reviews.llvm.org/D16331

Files:
  include/clang/Basic/DiagnosticSemaKinds.td
  lib/Sema/SemaExpr.cpp
  test/SemaCUDA/vararg.cu

Index: test/SemaCUDA/vararg.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/vararg.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
+}
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp
+++ lib/Sema/SemaExpr.cpp
@@ -11715,9 +11715,8 @@
   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,16 @@
   Expr *OrigExpr = E;
   bool IsMS = false;
 
+  // CUDA device code does not support varargs.
+  if (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.)
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -6422,6 +6422,8 @@
 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 "

Index: test/SemaCUDA/vararg.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/vararg.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
+}
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp
+++ lib/Sema/SemaExpr.cpp
@@ -11715,9 +11715,8 @@
   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,16 @@
   Expr *OrigExpr = E;
   bool IsMS = false;
 
+  // CUDA device code does not support varargs.
+  if (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.)
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -6422,6 +6422,8 @@
 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 "

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

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

Reply via email to