r282647 - [CUDA] Disallow variable-length arrays in CUDA device code.

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Wed Sep 28 17:45:58 2016
New Revision: 282647

URL: http://llvm.org/viewvc/llvm-project?rev=282647&view=rev
Log:
[CUDA] Disallow variable-length arrays in CUDA device code.

Reviewers: tra

Subscribers: cfe-commits, jhen

Differential Revision: https://reviews.llvm.org/D25050

Added:
    cfe/trunk/test/SemaCUDA/vla-host-device.cu
    cfe/trunk/test/SemaCUDA/vla.cu
Modified:
    cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
    cfe/trunk/include/clang/Sema/Sema.h
    cfe/trunk/lib/Sema/SemaCUDA.cpp
    cfe/trunk/lib/Sema/SemaType.cpp

Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=282647&r1=282646&r2=282647&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Wed Sep 28 17:45:58 2016
@@ -6713,6 +6713,10 @@ def err_shared_var_init : Error<
 def err_device_static_local_var : Error<
     "Within a __device__/__global__ function, "
     "only __shared__ variables may be marked \"static\"">;
+def err_cuda_vla : Error<
+    "cannot use variable-length arrays in "
+    "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
+
 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 "

Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=282647&r1=282646&r2=282647&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Wed Sep 28 17:45:58 2016
@@ -9255,6 +9255,8 @@ public:
   /// ExprTy should be the string "try" or "throw", as appropriate.
   bool CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy);
 
+  bool CheckCUDAVLA(SourceLocation Loc);
+
   /// Finds a function in \p Matches with highest calling priority
   /// from \p Caller context and erases all functions with lower
   /// calling priority.

Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=282647&r1=282646&r2=282647&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Sep 28 17:45:58 2016
@@ -539,3 +539,23 @@ bool Sema::CheckCUDAExceptionExpr(Source
   }
   return true;
 }
+
+bool Sema::CheckCUDAVLA(SourceLocation Loc) {
+  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
+  FunctionDecl *CurFn = dyn_cast<FunctionDecl>(CurContext);
+  if (!CurFn)
+    return true;
+  CUDAFunctionTarget Target = IdentifyCUDATarget(CurFn);
+  if (Target == CFT_Global || Target == CFT_Device) {
+    Diag(Loc, diag::err_cuda_vla) << Target;
+    return false;
+  }
+  if (Target == CFT_HostDevice && getLangOpts().CUDAIsDevice) {
+    PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()};
+    ErrPD.Reset(diag::err_cuda_vla);
+    ErrPD << Target;
+    CurFn->addDeferredDiag({Loc, std::move(ErrPD)});
+    return false;
+  }
+  return true;
+}

Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=282647&r1=282646&r2=282647&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Wed Sep 28 17:45:58 2016
@@ -2241,6 +2241,10 @@ QualType Sema::BuildArrayType(QualType T
     Diag(Loc, diag::err_opencl_vla);
     return QualType();
   }
+  // CUDA device code doesn't support VLAs.
+  if (getLangOpts().CUDA && T->isVariableArrayType() && !CheckCUDAVLA(Loc))
+    return QualType();
+
   // If this is not C99, extwarn about VLA's and C99 array size modifiers.
   if (!getLangOpts().C99) {
     if (T->isVariableArrayType()) {

Added: cfe/trunk/test/SemaCUDA/vla-host-device.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/vla-host-device.cu?rev=282647&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/vla-host-device.cu (added)
+++ cfe/trunk/test/SemaCUDA/vla-host-device.cu Wed Sep 28 17:45:58 2016
@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 -fcuda-is-device -verify -S %s -o /dev/null
+// RUN: %clang_cc1 -verify -DHOST %s -S -o /dev/null
+
+#include "Inputs/cuda.h"
+
+#ifdef HOST
+// expected-no-diagnostics
+#endif
+
+__host__ __device__ void hd(int n) {
+  int x[n];
+#ifndef HOST
+  // expected-error at -2 {{cannot use variable-length arrays in __host__ __device__ functions}}
+#endif
+}
+
+// No error because never codegen'ed for device.
+__host__ __device__ inline void hd_inline(int n) {
+  int x[n];
+}
+void call_hd_inline() { hd_inline(42); }

Added: cfe/trunk/test/SemaCUDA/vla.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/vla.cu?rev=282647&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/vla.cu (added)
+++ cfe/trunk/test/SemaCUDA/vla.cu Wed Sep 28 17:45:58 2016
@@ -0,0 +1,12 @@
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -verify -DHOST %s
+
+#include "Inputs/cuda.h"
+
+void host(int n) {
+  int x[n];
+}
+
+__device__ void device(int n) {
+  int x[n];  // expected-error {{cannot use variable-length arrays in __device__ functions}}
+}