[CUDA] Add Sema::CUDADiagBuilder and Sema::CUDADiagIfDeviceCode().

[jlebar]

Summary

Together these let you easily create diagnostics that

• are never emitted for host code
• are always emitted for device and global functions, and
• are emitted for host device functions iff these functions are codegen'ed.

At the moment there are only three diagnostics that need this treatment,
but I have more to add, and it's not sustainable to write code for emitting
every such diagnostic twice, and from a special wrapper in SemaCUDA.cpp.

While we're at it, don't emit the function name in
err_cuda_device_exceptions: It's not necessary to print it, and making
this work in the new framework in the face of a null value for
dyn_cast<FunctionDecl>(CurContext) isn't worth the effort.

[jlebar]

Add CUDADiagIfHostCode().

[jlebar]

Tweak API a bit.

Now we rely on an implicit conversion to bool. Which is not great, I know, but
in practice I think works better than an explicit named function.

[rnk]

clang/include/clang/Sema/Sema.h

9210 ↗

(On Diff #73171)

LLVM has a different enum naming convention: http://llvm.org/docs/CodingStandards.html#name-types-functions-variables-and-enumerators-properly

9238 ↗

(On Diff #73171)

I'm concerned that this usage pattern isn't going to be efficient because you build the complete diagnostic before calling the bool conversion operator to determine that it doesn't need to be emitted. I think you want to construct something more like:

if (isCUDADeviceCode())
  CUDADiag(...) << ...;

Otherwise you are going to construct and destruct a large number of diagnostics about language features that are forbidden in device code, but are legal in host code, and 99% of the TU is going to be host code that uses these illegal features.

9258 ↗

(On Diff #73171)

Remind me why we need to do this? Which arena is this stuff allocated in and where would I go to read more about it? My thought is that, if we don't construct very many of these, we should just allocate them in the usual ASTContext arena and let them live forever. It would be more consistent with our usual way of doing things.

[jlebar]

Address review comments, and rebase atop https://reviews.llvm.org/D24573.

[jlebar]

clang/include/clang/Sema/Sema.h

9238 ↗

(On Diff #73171)

I think the comment is misleading -- I tried to update it to resolve this misunderstanding. Does it make more sense now?

9258 ↗

(On Diff #73171)

These diagnostics live until the end of codegen, and so are destroyed after the ASTContext.

I am becoming increasingly displeased with emitting these errors during codegen. In particular, it makes it annoying to write tests. You cannot test deferred and immediate errors in the same test, because if you have any immediate errors, we never codegen, so we never emit the deferred ones. This will also be a suboptimal user experience.

The only serious alternative mooted thus far is emitting the deferred errors when a function is marked used. But this is going to emit deferred errors if e.g. two inline host+device functions have mutual recursion but are otherwise never touched (so don't need to be codegen'ed). I am not sure if this will be OK or not, but I need to look at it.

If we move the errors to being emitted earlier, we won't need to do this dance.

[rnk]

clang/include/clang/Sema/Sema.h

9258 ↗

(On Diff #73171)

The ASTContext should outlive IRgen, since the AST is allocated in its arena. Is there a separate diagnostic memory pool that I don't know about?

[jlebar]

Rebase atop https://reviews.llvm.org/D25260, which obviates the need for this
ugly PD allocation dance.

clang/include/clang/Sema/Sema.h

9258 ↗

(On Diff #73171)

You're right, this is a silly bug. Fixed in a separate patch, https://reviews.llvm.org/D25260 -- with that change we can do the natural thing here.

[jlebar]

If you're busy with other stuff, that's cool, this is not urgent, but pinging in case this got lost.

[rnk]

lgtm

clang/include/clang/Sema/Sema.h

9238 ↗

(On Diff #73171)

Actually, I just had to think about the problem more to realize why it has to be this way. This is for deferred diagnostics where you don't know if the code is device code yet. :)

cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td

 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">;
+  "%select{__device__|__global__|__host__|__host__ __device__}1 function">;
 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\"">;

cfe/trunk/include/clang/Sema/Sema.h

public:
   /// than zero, all functions encountered will be __host__ __device__.
   void PushForceCUDAHostDevice();
 
   /// Decrements our count of the number of times we've seen a pragma forcing
   /// functions to be __host__ __device__.  Returns false if the count is 0
   /// before incrementing, so you can emit an error.
   bool PopForceCUDAHostDevice();
 
+  /// Diagnostic builder for CUDA errors which may or may not be deferred.
+  ///
+  /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch)
+  /// which are not allowed to appear inside __device__ functions and are
+  /// allowed to appear in __host__ __device__ functions only if the host+device
+  /// function is never codegen'ed.
+  ///
+  /// To handle this, we use the notion of "deferred diagnostics", where we
+  /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed.
+  ///
+  /// This class lets you emit either a regular diagnostic, a deferred
+  /// diagnostic, or no diagnostic at all, according to an argument you pass to
+  /// its constructor, thus simplifying the process of creating these "maybe
+  /// deferred" diagnostics.
+  class CUDADiagBuilder {
+  public:
+    enum Kind {
+      /// Emit no diagnostics.
+      K_Nop,
+      /// Emit the diagnostic immediately (i.e., behave like Sema::Diag()).
+      K_Immediate,
+      /// Create a deferred diagnostic, which is emitted only if the function
+      /// it's attached to is codegen'ed.
+      K_Deferred
+    };
+
+    CUDADiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
+                    FunctionDecl *Fn, Sema &S);
+
+    /// Convertible to bool: True if we immediately emitted an error, false if
+    /// we didn't emit an error or we created a deferred error.
+    ///
+    /// Example usage:
+    ///
+    ///   if (CUDADiagBuilder(...) << foo << bar)
+    ///     return ExprError();
+    ///
+    /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
+    /// want to use these instead of creating a CUDADiagBuilder yourself.
+    operator bool() const { return ImmediateDiagBuilder.hasValue(); }
+
+    template <typename T>
+    friend const CUDADiagBuilder &operator<<(const CUDADiagBuilder &Diag,
+                                             const T &Value) {
+      if (Diag.ImmediateDiagBuilder.hasValue())
+        *Diag.ImmediateDiagBuilder << Value;
+      else if (Diag.PartialDiagInfo.hasValue())
+        Diag.PartialDiagInfo->PD << Value;
+      return Diag;
+    }
+
+  private:
+    struct PartialDiagnosticInfo {
+      PartialDiagnosticInfo(SourceLocation Loc, PartialDiagnostic PD,
+                            FunctionDecl *Fn)
+          : Loc(Loc), PD(std::move(PD)), Fn(Fn) {}
+
+      ~PartialDiagnosticInfo() { Fn->addDeferredDiag({Loc, std::move(PD)}); }
+
+      SourceLocation Loc;
+      PartialDiagnostic PD;
+      FunctionDecl *Fn;
+    };
+
+    // Invariant: At most one of these Optionals has a value.
+    // FIXME: Switch these to a Variant once that exists.
+    llvm::Optional<Sema::SemaDiagnosticBuilder> ImmediateDiagBuilder;
+    llvm::Optional<PartialDiagnosticInfo> PartialDiagInfo;
+  };
+
+  /// Creates a CUDADiagBuilder that emits the diagnostic if the current context
+  /// is "used as device code".
+  ///
+  /// - If CurContext is a __host__ function, does not emit any diagnostics.
+  /// - If CurContext is a __device__ or __global__ function, emits the
+  ///   diagnostics immediately.
+  /// - If CurContext is a __host__ __device__ function and we are compiling for
+  ///   the device, creates a deferred diagnostic which is emitted if and when
+  ///   the function is codegen'ed.
+  ///
+  /// Example usage:
+  ///
+  ///  // Variable-length arrays are not allowed in CUDA device code.
+  ///  if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget())
+  ///    return ExprError();
+  ///  // Otherwise, continue parsing as normal.
+  CUDADiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
+
+  /// Creates a CUDADiagBuilder that emits the diagnostic if the current context
+  /// is "used as host code".
+  ///
+  /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched.
+  CUDADiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID);
+
   enum CUDAFunctionTarget {
     CFT_Device,
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };
 
+  /// Determines whether the given function is a CUDA device/host/kernel/etc.
+  /// function.
+  ///
+  /// Use this rather than examining the function's attributes yourself -- you
+  /// will get it wrong.  Returns CFT_Host if D is null.
   CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
 
+  /// Gets the CUDA target for the current context.
+  CUDAFunctionTarget CurrentCUDATarget() {
+    return IdentifyCUDATarget(dyn_cast<FunctionDecl>(CurContext));
+  }
+
   // CUDA function call preference. Must be ordered numerically from
   // worst to best.
   enum CUDAFunctionPreference {
     CFP_Never,      // Invalid caller/callee combination.
     CFP_WrongSide,  // Calls from host-device to host or device
                     // function that do not match current compilation
                     // mode.
     CFP_HostDevice, // Any calls to host/device functions.
public:
 
   /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
   /// depending on FD and the current compilation settings.
   void maybeAddCUDAHostDeviceAttrs(Scope *S, FunctionDecl *FD,
                                    const LookupResult &Previous);
 
 private:
   /// Raw encodings of SourceLocations for which CheckCUDACall has emitted a
-  /// deferred "bad call" diagnostic.  We use this to avoid emitting the same
-  /// deferred diag twice.
-  llvm::DenseSet<unsigned> LocsWithCUDACallDeferredDiags;
+  /// (maybe deferred) "bad call" diagnostic.  We use this to avoid emitting the
+  /// same deferred diag twice.
+  llvm::DenseSet<unsigned> LocsWithCUDACallDiags;
 
 public:
   /// Check whether we're allowed to call Callee from the current context.
   ///
   /// - If the call is never allowed in a semantically-correct program
   ///   (CFP_Never), emits an error and returns false.
   ///
   /// - 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.
   ///
   ///   Will only create deferred diagnostics for a given SourceLocation once,
   ///   so you can safely call this multiple times without generating duplicate
   ///   deferred errors.
   ///
   /// - 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);
-
-  /// Check whether it's legal for us to create a variable-length array in the
-  /// current context.  Returns true if the VLA is OK; returns false and emits
-  /// an error otherwise.
-  bool CheckCUDAVLA(SourceLocation Loc);
-
   /// Set __device__ or __host__ __device__ attributes on the given lambda
   /// operator() method.
   ///
   /// CUDA lambdas declared inside __device__ or __global__ functions inherit
   /// the __device__ attribute.  Similarly, lambdas inside __host__ __device__
   /// functions become __host__ __device__ themselves.
   void CUDASetLambdaAttrs(CXXMethodDecl *Method);
 

cfe/trunk/lib/Sema/SemaCUDA.cpp

 
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/ExprCXX.h"
 #include "clang/Lex/Preprocessor.h"
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/Sema.h"
 #include "clang/Sema/SemaDiagnostic.h"
+#include "clang/Sema/SemaInternal.h"
 #include "clang/Sema/Template.h"
 #include "llvm/ADT/Optional.h"
 #include "llvm/ADT/SmallVector.h"
 using namespace clang;
 
 void Sema::PushForceCUDAHostDevice() {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   ForceCUDAHostDeviceDepth++;
ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
   MarkFunctionReferenced(LLLLoc, ConfigDecl);
 
   return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
                        /*IsExecConfig=*/true);
 }
 
 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
 Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
+  // Code that lives outside a function is run on the host.
+  if (D == nullptr)
+    return CFT_Host;
+
   if (D->hasAttr<CUDAInvalidTargetAttr>())
     return CFT_InvalidTarget;
 
   if (D->hasAttr<CUDAGlobalAttr>())
     return CFT_Global;
 
   if (D->hasAttr<CUDADeviceAttr>()) {
     if (D->hasAttr<CUDAHostAttr>())

 // | hd | g  | SS  | --  |(d/a)|
 // | hd | h  | SS  | WS  | (d) |
 // | hd | hd | HD  | HD  | (b) |
 
 Sema::CUDAFunctionPreference
 Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
                              const FunctionDecl *Callee) {
   assert(Callee && "Callee must be valid.");
+  CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
   CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
-  CUDAFunctionTarget CallerTarget =
-      (Caller != nullptr) ? IdentifyCUDATarget(Caller) : Sema::CFT_Host;
 
   // If one of the targets is invalid, the check always fails, no matter what
   // the other target is.
   if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
     return CFP_Never;
 
   // (a) Can't call global from some contexts until we support CUDA's
   // dynamic parallelism.
if (It != Previous.end()) {
     }
     return;
   }
 
   NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
-bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
-  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
-  assert(Callee && "Callee may not be null.");
-  FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
-  if (!Caller)
-    return true;
-
-  Sema::CUDAFunctionPreference Pref = IdentifyCUDAPreference(Caller, Callee);
-  if (Pref == Sema::CFP_Never) {
-    Diag(Loc, diag::err_ref_bad_target) << IdentifyCUDATarget(Callee) << Callee
-                                        << IdentifyCUDATarget(Caller);
-    Diag(Callee->getLocation(), diag::note_previous_decl) << Callee;
-    return false;
-  }
-
-  // Insert into LocsWithCUDADeferredDiags to avoid emitting duplicate deferred
-  // diagnostics for the same location.  Duplicate deferred diags are otherwise
-  // tricky to avoid, because, unlike with regular errors, sema checking
-  // proceeds unhindered when we omit a deferred diagnostic.
-  if (Pref == Sema::CFP_WrongSide &&
-      LocsWithCUDACallDeferredDiags.insert(Loc.getRawEncoding()).second) {
-    // We have to do this odd dance to create our PartialDiagnostic because we
-    // want its storage to be allocated with operator new, not in an arena.
-    PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()};
-    ErrPD.Reset(diag::err_ref_bad_target);
-    ErrPD << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
-    Caller->addDeferredDiag({Loc, std::move(ErrPD)});
-
-    PartialDiagnostic NotePD{PartialDiagnostic::NullDiagnostic()};
-    NotePD.Reset(diag::note_previous_decl);
-    NotePD << Callee;
-    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;
+Sema::CUDADiagBuilder::CUDADiagBuilder(Kind K, SourceLocation Loc,
+                                       unsigned DiagID, FunctionDecl *Fn,
+                                       Sema &S) {
+  switch (K) {
+  case K_Nop:
+    break;
+  case K_Immediate:
+    ImmediateDiagBuilder.emplace(S.Diag(Loc, DiagID));
+    break;
+  case K_Deferred:
+    assert(Fn && "Must have a function to attach the deferred diag to.");
+    PartialDiagInfo.emplace(Loc, S.PDiag(DiagID), Fn);
+    break;
   }
-  return true;
 }
 
-bool Sema::CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy) {
+Sema::CUDADiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
+                                                 unsigned DiagID) {
   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;
+  CUDADiagBuilder::Kind DiagKind;
+  switch (CurrentCUDATarget()) {
+  case CFT_Global:
+  case CFT_Device:
+    DiagKind = CUDADiagBuilder::K_Immediate;
+    break;
+  case CFT_HostDevice:
+    DiagKind = getLangOpts().CUDAIsDevice ? CUDADiagBuilder::K_Deferred
+                                          : CUDADiagBuilder::K_Nop;
+    break;
+  default:
+    DiagKind = CUDADiagBuilder::K_Nop;
   }
-  return true;
+  return CUDADiagBuilder(DiagKind, Loc, DiagID,
+                         dyn_cast<FunctionDecl>(CurContext), *this);
 }
 
-bool Sema::CheckCUDAVLA(SourceLocation Loc) {
+Sema::CUDADiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
+                                               unsigned DiagID) {
   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;
+  CUDADiagBuilder::Kind DiagKind;
+  switch (CurrentCUDATarget()) {
+  case CFT_Host:
+    DiagKind = CUDADiagBuilder::K_Immediate;
+    break;
+  case CFT_HostDevice:
+    DiagKind = getLangOpts().CUDAIsDevice ? CUDADiagBuilder::K_Nop
+                                          : CUDADiagBuilder::K_Deferred;
+    break;
+  default:
+    DiagKind = CUDADiagBuilder::K_Nop;
   }
-  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 CUDADiagBuilder(DiagKind, Loc, DiagID,
+                         dyn_cast<FunctionDecl>(CurContext), *this);
+
+bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
+  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
+  assert(Callee && "Callee may not be null.");
+  FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
+  if (!Caller)
     return true;
+
+  CUDADiagBuilder::Kind DiagKind;
+  switch (IdentifyCUDAPreference(Caller, Callee)) {
+  case CFP_Never:
+    DiagKind = CUDADiagBuilder::K_Immediate;
+    break;
+  case CFP_WrongSide:
+    assert(Caller && "WrongSide calls require a non-null caller");
+    DiagKind = CUDADiagBuilder::K_Deferred;
+    break;
+  default:
+    DiagKind = CUDADiagBuilder::K_Nop;
+  }
+
+  // Avoid emitting this error twice for the same location.  Using a hashtable
+  // like this is unfortunate, but because we must continue parsing as normal
+  // after encountering a deferred error, it's otherwise very tricky for us to
+  // ensure that we only emit this deferred error once.
+  if (!LocsWithCUDACallDiags.insert(Loc.getRawEncoding()).second)
+    return true;
+
+  bool IsImmediateErr =
+      CUDADiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
+      << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
+  CUDADiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl,
+                  Caller, *this)
+      << Callee;
+  return !IsImmediateErr;
 }
 
 void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
     return;
   FunctionDecl *CurFn = dyn_cast<FunctionDecl>(CurContext);
   if (!CurFn)

cfe/trunk/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");
+    CUDADiagIfDeviceCode(OpLoc, diag::err_cuda_device_exceptions)
+        << "throw" << CurrentCUDATarget();
 
   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();

cfe/trunk/lib/Sema/SemaStmt.cpp

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");
+    CUDADiagIfDeviceCode(TryLoc, diag::err_cuda_device_exceptions)
+        << "try" << CurrentCUDATarget();
 
   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()) {

cfe/trunk/lib/Sema/SemaType.cpp

QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM,
   }
 
   // OpenCL v1.2 s6.9.d: variable length arrays are not supported.
   if (getLangOpts().OpenCL && T->isVariableArrayType()) {
     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 (getLangOpts().CUDA && T->isVariableArrayType())
+    CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget();
 
   // If this is not C99, extwarn about VLA's and C99 array size modifiers.
   if (!getLangOpts().C99) {
     if (T->isVariableArrayType()) {
       // Prohibit the use of VLAs during template argument deduction.
       if (isSFINAEContext()) {
         Diag(Loc, diag::err_vla_in_sfinae);
         return QualType();

cfe/trunk/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'}}
+  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
+  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
 #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'}}
+  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
+  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
 #endif
 }
 __device__ void call_hd3() { hd3(); }

cfe/trunk/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'}}
+  // expected-error@-1 {{cannot use 'throw' in __device__ function}}
   try {} catch(void*) {}
-  // expected-error@-1 {{cannot use 'try' in __device__ function 'device'}}
+  // expected-error@-1 {{cannot use 'try' in __device__ function}}
 }
 __global__ void kernel() {
   throw NULL;
-  // expected-error@-1 {{cannot use 'throw' in __global__ function 'kernel'}}
+  // expected-error@-1 {{cannot use 'throw' in __global__ function}}
   try {} catch(void*) {}
-  // expected-error@-1 {{cannot use 'try' in __global__ function 'kernel'}}
+  // expected-error@-1 {{cannot use 'try' in __global__ function}}
 }