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@-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}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
