[CUDA] Add Sema::CUDADiagBuilder and Sema::CUDADiagIf{Device,Host}Code().

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.

Reviewers: rnk

Subscribers: cfe-commits, tra

Differential Revision: https://reviews.llvm.org/D25139

llvm-svn: 284143

2 files changed, 108 insertions, 19 deletions

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 59199aea7b8..ff7b3422a94 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6734,7 +6734,7 @@ 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.">;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 166f23b9768..1417b23b9ae 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -9245,6 +9245,100 @@ public:
   /// 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,
@@ -9253,8 +9347,18 @@ public:
     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 {
@@ -9295,9 +9399,9 @@ public:
 
 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.
@@ -9316,21 +9420,6 @@ 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);
-
-  /// 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.
   ///