[HIP] Defer operator overloading errors

Although clang is able to defer overloading resolution
diagnostics for common functions. It does not defer
overloading resolution caused diagnostics for overloaded
operators.

This patch extends the existing deferred
diagnostic mechanism and defers a diagnostic caused
by overloaded operator.

Reviewed by: Artem Belevich

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

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -6771,7 +6771,7 @@ def warn_param_mismatched_alignment : Warning<
 def err_objc_object_assignment : Error<
   "cannot assign to class object (%0 invalid)">;
 def err_typecheck_invalid_operands : Error<
-  "invalid operands to binary expression (%0 and %1)">;
+  "invalid operands to binary expression (%0 and %1)">, Deferrable;
 def note_typecheck_invalid_operands_converted : Note<
   "%select{first|second}0 operand was implicitly converted to type %1">;
 def err_typecheck_logical_vector_expr_gnu_cpp_restrict : Error<

clang/include/clang/Sema/Overload.h

@@ -1048,9 +1048,6 @@ class Sema;
 
     void destroyCandidates();
 
-    /// Whether diagnostics should be deferred.
-    bool shouldDeferDiags(Sema &S, ArrayRef<Expr *> Args, SourceLocation OpLoc);
-
   public:
     OverloadCandidateSet(SourceLocation Loc, CandidateSetKind CSK,
                          OperatorRewriteInfo RewriteInfo = {})
@@ -1063,6 +1060,9 @@ class Sema;
     CandidateSetKind getKind() const { return Kind; }
     OperatorRewriteInfo getRewriteInfo() const { return RewriteInfo; }
 
+    /// Whether diagnostics should be deferred.
+    bool shouldDeferDiags(Sema &S, ArrayRef<Expr *> Args, SourceLocation OpLoc);
+
     /// Determine when this overload candidate will be new to the
     /// overload set.
     bool isNewCandidate(Decl *F, OverloadCandidateParamOrder PO =

clang/include/clang/Sema/Sema.h

@@ -1769,6 +1769,22 @@ public:
   /// Build a partial diagnostic.
   PartialDiagnostic PDiag(unsigned DiagID = 0); // in SemaInternal.h
 
+  /// Whether deferrable diagnostics should be deferred.
+  bool DeferDiags = false;
+
+  /// RAII class to control scope of DeferDiags.
+  class DeferDiagsRAII {
+    Sema &S;
+    bool SavedDeferDiags = false;
+
+  public:
+    DeferDiagsRAII(Sema &_S, bool DeferDiags)
+        : S(_S), SavedDeferDiags(S.DeferDiags) {
+      S.DeferDiags = DeferDiags;
+    }
+    ~DeferDiagsRAII() { S.DeferDiags = SavedDeferDiags; }
+  };
+
   /// Whether uncompilable error has occurred. This includes error happens
   /// in deferred diagnostics.
   bool hasUncompilableErrorOccurred() const;

clang/lib/Sema/Sema.cpp

@@ -1790,7 +1790,7 @@ Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, unsigned DiagID,
   bool IsError = Diags.getDiagnosticIDs()->isDefaultMappingAsError(DiagID);
   bool ShouldDefer = getLangOpts().CUDA && LangOpts.GPUDeferDiag &&
                      DiagnosticIDs::isDeferrable(DiagID) &&
-                     (DeferHint || !IsError);
+                     (DeferHint || DeferDiags || !IsError);
   auto SetIsLastErrorImmediate = [&](bool Flag) {
     if (IsError)
       IsLastErrorImmediate = Flag;

clang/lib/Sema/SemaOverload.cpp

@@ -11641,7 +11641,8 @@ bool OverloadCandidateSet::shouldDeferDiags(Sema &S, ArrayRef<Expr *> Args,
         CompleteCandidates(S, OCD_AllCandidates, Args, OpLoc, [](auto &Cand) {
           return (Cand.Viable == false &&
                   Cand.FailureKind == ovl_fail_bad_target) ||
-                 (Cand.Function->template hasAttr<CUDAHostAttr>() &&
+                 (Cand.Function &&
+                  Cand.Function->template hasAttr<CUDAHostAttr>() &&
                   Cand.Function->template hasAttr<CUDADeviceAttr>());
         });
     DeferHint = !WrongSidedCands.empty();
@@ -13820,6 +13821,8 @@ ExprResult Sema::CreateOverloadedBinOp(SourceLocation OpLoc,
       StringRef OpcStr = BinaryOperator::getOpcodeStr(Opc);
       auto Cands = CandidateSet.CompleteCandidates(*this, OCD_AllCandidates,
                                                    Args, OpLoc);
+      DeferDiagsRAII DDR(*this,
+                         CandidateSet.shouldDeferDiags(*this, Args, OpLoc));
       if (Args[0]->getType()->isRecordType() &&
           Opc >= BO_Assign && Opc <= BO_OrAssign) {
         Diag(OpLoc,  diag::err_ovl_no_viable_oper)

clang/test/SemaCUDA/deferred-oeverload.cu

@@ -5,6 +5,11 @@
 // RUN: %clang_cc1 -fopenmp -fsyntax-only -verify=host,com %s \
 // RUN:   -std=c++11 -fgpu-defer-diag
 
+// With -fgpu-defer-diag, clang defers overloading resolution induced
+// diagnostics when the full candidates set include host device
+// functions or wrong-sided candidates. This roughly matches nvcc's
+// behavior.
+
 #include "Inputs/cuda.h"
 
 // When callee is called by a host function with integer arguments, there is an error for ambiguity.
@@ -31,12 +36,20 @@ __host__ void callee4(int); // com-note 2{{candidate function not viable: requir
 __host__ void callee5(float); // com-note {{candidate function}}
 __host__ void callee5(double); // com-note {{candidate function}}
 
+// When '<<` operator is called by a device function, there is error for 'invalid operands'.
+// It should be deferred since it involves wrong-sided candidates.
+struct S {
+  __host__ S &operator <<(int i); // dev-note {{candidate function not viable}}
+};
+
 __host__ void hf() {
  callee(1); // host-error {{call to 'callee' is ambiguous}}
  callee2();
  callee3();
  callee4(); // com-error {{no matching function for call to 'callee4'}}
  callee5(1); // com-error {{call to 'callee5' is ambiguous}}
+ S s;
+ s << 1;
  undeclared_func(); // com-error {{use of undeclared identifier 'undeclared_func'}}
 }
 
@@ -45,6 +58,8 @@ __device__ void df() {
  callee2(); // dev-error {{no matching function for call to 'callee2'}}
  callee3(); // dev-error {{no matching function for call to 'callee3'}}
  callee4(); // com-error {{no matching function for call to 'callee4'}}
+ S s;
+ s << 1;    // dev-error {{invalid operands to binary expression}}
 }
 
 struct A { int x; typedef int isA; };