[CUDA] Disallow exceptions in device code.

[tra]

One question, LGTM otherwise.

clang/lib/Sema/SemaExprCXX.cpp

688

Do you need/want to check returned result?

[jlebar]

clang/lib/Sema/SemaExprCXX.cpp

688

We could, and we could return ExprError here, but I thought it would make sense to do the same thing that we do right above, for -fno-exceptions: Continue parsing as normal, since we in fact *can* understand what the user is trying to do. In theory (certainly for try/catch) this will let us emit better errors elsewhere.

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_va_arg_in_device : Error<
   "CUDA device code does not support va_arg">;
 def err_alias_not_supported_on_nvptx : Error<"CUDA does not support aliases">;
 def err_cuda_unattributed_constexpr_cannot_overload_device : Error<
   "constexpr function '%0' without __host__ or __device__ attributes cannot "
   "overload __device__ function with same signature.  Add a __host__ "
   "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.">;
 def err_shared_var_init : Error<
     "initialization is not supported for __shared__ variables.">;
 def err_device_static_local_var : Error<
     "Within a __device__/__global__ function, "
     "only __shared__ variables may be marked \"static\"">;

clang/include/clang/Sema/Sema.h

   ///
   /// If the call is allowed in semantically-correct programs, but only if it's
   /// never codegen'ed (CFP_WrongSide), creates a deferred diagnostic to be
   /// emitted if and when the caller is codegen'ed, and returns true.
   ///
   /// 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.
   void EraseUnwantedCUDAMatches(const FunctionDecl *Caller,
                                 SmallVectorImpl<FunctionDecl *> &Matches);
   void EraseUnwantedCUDAMatches(const FunctionDecl *Caller,
                                 SmallVectorImpl<DeclAccessPair> &Matches);
   void EraseUnwantedCUDAMatches(

clang/lib/Sema/SemaCUDA.cpp

if (Pref == Sema::CFP_WrongSide) {
     Caller->addDeferredDiag({Callee->getLocation(), std::move(NotePD)});
 
     // This is not immediately an error, so return true.  The deferred errors
     // will be emitted if and when Caller is codegen'ed.
     return true;
   }
   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;
+}

clang/lib/Sema/SemaExprCXX.cpp

 
 ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex, 
                                bool IsThrownVarInScope) {
   // Don't report an error if 'throw' is used in system headers.
   if (!getLangOpts().CXXExceptions &&
       !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";
 
   if (Ex && !Ex->isTypeDependent()) {
     QualType ExceptionObjectTy = Context.getExceptionObjectType(Ex->getType());
     if (CheckCXXThrowOperand(OpLoc, ExceptionObjectTy, Ex))
       return ExprError();
 

clang/lib/Sema/SemaStmt.cpp

 /// handlers and creates a try statement from them.
 StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock,
                                   ArrayRef<Stmt *> Handlers) {
   // Don't report an error if 'try' is used in system headers.
   if (!getLangOpts().CXXExceptions &&
       !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";
 
   sema::FunctionScopeInfo *FSI = getCurFunction();
 
   // C++ try is incompatible with SEH __try.
   if (!getLangOpts().Borland && FSI->FirstSEHTryLoc.isValid()) {
     Diag(TryLoc, diag::err_mixing_cxx_try_seh_try);

clang/test/SemaCUDA/exceptions-host-device.cu

+// 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(); }

clang/test/SemaCUDA/exceptions.cu

+// 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'}}
+}