This revision was automatically updated to reflect the committed changes. jlebar marked an inline comment as done. Closed by commit rL282646: [CUDA] Disallow exceptions in device code. (authored by jlebar).
Changed prior to commit: https://reviews.llvm.org/D25036?vs=72872&id=72918#toc Repository: rL LLVM https://reviews.llvm.org/D25036 Files: 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 cfe/trunk/test/SemaCUDA/exceptions-host-device.cu cfe/trunk/test/SemaCUDA/exceptions.cu
Index: cfe/trunk/include/clang/Sema/Sema.h =================================================================== --- cfe/trunk/include/clang/Sema/Sema.h +++ cfe/trunk/include/clang/Sema/Sema.h @@ -9245,6 +9245,16 @@ /// 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. Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td @@ -6702,6 +6702,9 @@ "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.">; Index: cfe/trunk/test/SemaCUDA/exceptions-host-device.cu =================================================================== --- cfe/trunk/test/SemaCUDA/exceptions-host-device.cu +++ cfe/trunk/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(); } Index: cfe/trunk/test/SemaCUDA/exceptions.cu =================================================================== --- cfe/trunk/test/SemaCUDA/exceptions.cu +++ cfe/trunk/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'}} +} Index: cfe/trunk/lib/Sema/SemaCUDA.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaCUDA.cpp +++ cfe/trunk/lib/Sema/SemaCUDA.cpp @@ -515,3 +515,27 @@ } 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; +} Index: cfe/trunk/lib/Sema/SemaStmt.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaStmt.cpp +++ cfe/trunk/lib/Sema/SemaStmt.cpp @@ -3644,6 +3644,10 @@ !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"; Index: cfe/trunk/lib/Sema/SemaExprCXX.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaExprCXX.cpp +++ cfe/trunk/lib/Sema/SemaExprCXX.cpp @@ -683,6 +683,10 @@ !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";
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits