From 6d0bdd7405b36e6823250610cff42512dee14119 Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Wed, 28 Sep 2016 22:45:54 +0000 Subject: [PATCH] [CUDA] Disallow exceptions in device code. Reviewers: tra Subscribers: cfe-commits, jhen Differential Revision: https://reviews.llvm.org/D25036 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@282646 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/DiagnosticSemaKinds.td | 3 ++ include/clang/Sema/Sema.h | 10 ++++++ lib/Sema/SemaCUDA.cpp | 24 ++++++++++++++ lib/Sema/SemaExprCXX.cpp | 4 +++ lib/Sema/SemaStmt.cpp | 4 +++ test/SemaCUDA/exceptions-host-device.cu | 38 ++++++++++++++++++++++ test/SemaCUDA/exceptions.cu | 21 ++++++++++++ 7 files changed, 104 insertions(+) create mode 100644 test/SemaCUDA/exceptions-host-device.cu create mode 100644 test/SemaCUDA/exceptions.cu diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td index af484d4003..36b1ebd229 100644 --- a/include/clang/Basic/DiagnosticSemaKinds.td +++ b/include/clang/Basic/DiagnosticSemaKinds.td @@ -6702,6 +6702,9 @@ def err_cuda_unattributed_constexpr_cannot_overload_device : Error< "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.">; diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h index 1ae2004a6f..2734e79652 100644 --- a/include/clang/Sema/Sema.h +++ b/include/clang/Sema/Sema.h @@ -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. diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp index 6f94e54a0c..b1939a1715 100644 --- a/lib/Sema/SemaCUDA.cpp +++ b/lib/Sema/SemaCUDA.cpp @@ -515,3 +515,27 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) { } return true; } + +bool Sema::CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy) { + assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); + FunctionDecl *CurFn = dyn_cast(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; +} diff --git a/lib/Sema/SemaExprCXX.cpp b/lib/Sema/SemaExprCXX.cpp index a0bf324a71..42888badd2 100644 --- a/lib/Sema/SemaExprCXX.cpp +++ b/lib/Sema/SemaExprCXX.cpp @@ -683,6 +683,10 @@ ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex, !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"; diff --git a/lib/Sema/SemaStmt.cpp b/lib/Sema/SemaStmt.cpp index 977d7447d7..901f875840 100644 --- a/lib/Sema/SemaStmt.cpp +++ b/lib/Sema/SemaStmt.cpp @@ -3644,6 +3644,10 @@ StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock, !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"; diff --git a/test/SemaCUDA/exceptions-host-device.cu b/test/SemaCUDA/exceptions-host-device.cu new file mode 100644 index 0000000000..3f323469ac --- /dev/null +++ b/test/SemaCUDA/exceptions-host-device.cu @@ -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(); } diff --git a/test/SemaCUDA/exceptions.cu b/test/SemaCUDA/exceptions.cu new file mode 100644 index 0000000000..75586bc4a9 --- /dev/null +++ b/test/SemaCUDA/exceptions.cu @@ -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'}} +} -- 2.40.0