Author: jlebar Date: Wed Sep 28 17:45:54 2016 New Revision: 282646 URL: http://llvm.org/viewvc/llvm-project?rev=282646&view=rev Log: [CUDA] Disallow exceptions in device code.
Reviewers: tra Subscribers: cfe-commits, jhen Differential Revision: https://reviews.llvm.org/D25036 Added: cfe/trunk/test/SemaCUDA/exceptions-host-device.cu cfe/trunk/test/SemaCUDA/exceptions.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/SemaExprCXX.cpp cfe/trunk/lib/Sema/SemaStmt.cpp Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=282646&r1=282645&r2=282646&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Wed Sep 28 17:45:54 2016 @@ -6702,6 +6702,9 @@ def err_cuda_unattributed_constexpr_cann "attribute, or build with -fno-cuda-host-device-constexpr.">; def note_cuda_conflicting_device_function_declared_here : Note< "conflicting __device__ function declared here">; +def err_cuda_device_exceptions : Error< + "cannot use '%0' in " + "%select{__device__|__global__|__host__|__host__ __device__}1 function %2">; def err_dynamic_var_init : Error< "dynamic initialization is not supported for " "__device__, __constant__, and __shared__ variables.">; Modified: cfe/trunk/include/clang/Sema/Sema.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=282646&r1=282645&r2=282646&view=diff ============================================================================== --- cfe/trunk/include/clang/Sema/Sema.h (original) +++ cfe/trunk/include/clang/Sema/Sema.h Wed Sep 28 17:45:54 2016 @@ -9245,6 +9245,16 @@ public: /// Otherwise, returns true without emitting any diagnostics. bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee); + /// Check whether a 'try' or 'throw' expression is allowed within the current + /// context, and raise an error or create a deferred error, as appropriate. + /// + /// 'try' and 'throw' are never allowed in CUDA __device__ functions, and are + /// allowed in __host__ __device__ functions only if those functions are never + /// codegen'ed for the device. + /// + /// ExprTy should be the string "try" or "throw", as appropriate. + bool CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy); + /// 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=282646&r1=282645&r2=282646&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaCUDA.cpp (original) +++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Sep 28 17:45:54 2016 @@ -515,3 +515,27 @@ bool Sema::CheckCUDACall(SourceLocation } return true; } + +bool Sema::CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy) { + assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); + FunctionDecl *CurFn = dyn_cast<FunctionDecl>(CurContext); + if (!CurFn) + return true; + CUDAFunctionTarget Target = IdentifyCUDATarget(CurFn); + + // Raise an error immediately if this is a __global__ or __device__ function. + // If it's a __host__ __device__ function, enqueue a deferred error which will + // be emitted if the function is codegen'ed for device. + if (Target == CFT_Global || Target == CFT_Device) { + Diag(Loc, diag::err_cuda_device_exceptions) << ExprTy << Target << CurFn; + return false; + } + if (Target == CFT_HostDevice && getLangOpts().CUDAIsDevice) { + PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()}; + ErrPD.Reset(diag::err_cuda_device_exceptions); + ErrPD << ExprTy << Target << CurFn; + CurFn->addDeferredDiag({Loc, std::move(ErrPD)}); + return false; + } + return true; +} Modified: cfe/trunk/lib/Sema/SemaExprCXX.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExprCXX.cpp?rev=282646&r1=282645&r2=282646&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaExprCXX.cpp (original) +++ cfe/trunk/lib/Sema/SemaExprCXX.cpp Wed Sep 28 17:45:54 2016 @@ -683,6 +683,10 @@ ExprResult Sema::BuildCXXThrow(SourceLoc !getSourceManager().isInSystemHeader(OpLoc)) Diag(OpLoc, diag::err_exceptions_disabled) << "throw"; + // Exceptions aren't allowed in CUDA device code. + if (getLangOpts().CUDA) + CheckCUDAExceptionExpr(OpLoc, "throw"); + if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope()) Diag(OpLoc, diag::err_omp_simd_region_cannot_use_stmt) << "throw"; Modified: cfe/trunk/lib/Sema/SemaStmt.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaStmt.cpp?rev=282646&r1=282645&r2=282646&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaStmt.cpp (original) +++ cfe/trunk/lib/Sema/SemaStmt.cpp Wed Sep 28 17:45:54 2016 @@ -3644,6 +3644,10 @@ StmtResult Sema::ActOnCXXTryBlock(Source !getSourceManager().isInSystemHeader(TryLoc)) Diag(TryLoc, diag::err_exceptions_disabled) << "try"; + // Exceptions aren't allowed in CUDA device code. + if (getLangOpts().CUDA) + CheckCUDAExceptionExpr(TryLoc, "try"); + if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope()) Diag(TryLoc, diag::err_omp_simd_region_cannot_use_stmt) << "try"; Added: cfe/trunk/test/SemaCUDA/exceptions-host-device.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/exceptions-host-device.cu?rev=282646&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/exceptions-host-device.cu (added) +++ cfe/trunk/test/SemaCUDA/exceptions-host-device.cu Wed Sep 28 17:45:54 2016 @@ -0,0 +1,38 @@ +// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -verify %s -S -o /dev/null +// RUN: %clang_cc1 -fcxx-exceptions -verify -DHOST %s -S -o /dev/null + +#include "Inputs/cuda.h" + +// Check that it's an error to use 'try' and 'throw' from a __host__ __device__ +// function if and only if it's codegen'ed for device. + +#ifdef HOST +// expected-no-diagnostics +#endif + +__host__ __device__ void hd1() { + throw NULL; + try {} catch(void*) {} +#ifndef HOST + // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function 'hd1'}} + // expected-error@-3 {{cannot use 'try' in __host__ __device__ function 'hd1'}} +#endif +} + +// No error, never instantiated on device. +inline __host__ __device__ void hd2() { + throw NULL; + try {} catch(void*) {} +} +void call_hd2() { hd2(); } + +// Error, instantiated on device. +inline __host__ __device__ void hd3() { + throw NULL; + try {} catch(void*) {} +#ifndef HOST + // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function 'hd3'}} + // expected-error@-3 {{cannot use 'try' in __host__ __device__ function 'hd3'}} +#endif +} +__device__ void call_hd3() { hd3(); } Added: cfe/trunk/test/SemaCUDA/exceptions.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/exceptions.cu?rev=282646&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/exceptions.cu (added) +++ cfe/trunk/test/SemaCUDA/exceptions.cu Wed Sep 28 17:45:54 2016 @@ -0,0 +1,21 @@ +// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify %s + +#include "Inputs/cuda.h" + +void host() { + throw NULL; + try {} catch(void*) {} +} +__device__ void device() { + throw NULL; + // expected-error@-1 {{cannot use 'throw' in __device__ function 'device'}} + try {} catch(void*) {} + // expected-error@-1 {{cannot use 'try' in __device__ function 'device'}} +} +__global__ void kernel() { + throw NULL; + // expected-error@-1 {{cannot use 'throw' in __global__ function 'kernel'}} + try {} catch(void*) {} + // expected-error@-1 {{cannot use 'try' in __global__ function 'kernel'}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits