[CUDA] Do a better job at detecting wrong-side calls.

[jlebar]

Summary

Move CheckCUDACall from ActOnCallExpr and BuildDeclRefExpr to
DiagnoseUseOfDecl. This lets us catch some edge cases we were missing,
specifically around class operators.

This necessitates a few other changes:

• Avoid emitting duplicate deferred diags in CheckCUDACall.
  
  Previously we'd carefully placed our call to CheckCUDACall such that it would only ever run once for a particular callsite. But now this isn't the case.
• Emit deferred diagnostics from a template specialization/instantiation's primary template, in addition to from the specialization/instantiation itself. DiagnoseUseOfDecl ends up putting the deferred diagnostics on the template, rather than the specialization, so we need to check both.

[jlebar]

I'm sorry to put more stuff on your plate, Richard, but you already have context on this, so you may be the best person to review this. Unfortunately people keep hitting the bug fixed here (and being confused), so it's maybe not as much of an edge case as I thought.

FWIW this patch is not actually as big as it looks at first: The ActOnCallExpr changes are just moving things out of an "Impl" function, nothing interesting there.

clang/include/clang/Sema/Sema.h

bool IsAllowedCUDACall(const FunctionDecl *Caller,
     return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
   }
 
   /// 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;
+
+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 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.
+  /// - 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.
+  /// - Otherwise, returns true without emitting any diagnostics.
   bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee);
 
   /// 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,

clang/lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
   // of the diagnostics just yet.  Emitting an error during codegen stops
   // further codegen, and we want to display as many deferred diags as possible.
   // We'll emit the now twice-deferred diags at the very end of codegen.
   //
   // (If a function has both error and non-error diags, we don't emit the
   // non-error diags here, because order can be significant, e.g. with notes
   // that follow errors.)
   auto Diags = D->takeDeferredDiags();
+  if (auto *Templ = D->getPrimaryTemplate()) {
+    auto TemplDiags = Templ->getAsFunction()->takeDeferredDiags();
+    Diags.insert(Diags.end(), TemplDiags.begin(), TemplDiags.end());
+  }
   bool HasError = llvm::any_of(Diags, [this](const PartialDiagnosticAt &PDAt) {
     return getDiags().getDiagnosticLevel(PDAt.second.getDiagID(), PDAt.first) >=
            DiagnosticsEngine::Error;
   });
   if (HasError) {
     DeferredDiags.insert(DeferredDiags.end(),
                          std::make_move_iterator(Diags.begin()),
                          std::make_move_iterator(Diags.end()));

clang/lib/Sema/SemaCUDA.cpp

bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
 
   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;
   }
-  if (Pref == Sema::CFP_WrongSide) {
+
+  // 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()};

clang/lib/Sema/SemaExpr.cpp

if (FD->isDeleted()) {
       return true;
     }
 
     // If the function has a deduced return type, and we can't deduce it,
     // then we can't use it either.
     if (getLangOpts().CPlusPlus14 && FD->getReturnType()->isUndeducedType() &&
         DeduceReturnType(FD, Loc))
       return true;
+
+    if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD))
+      return true;
   }
 
   // [OpenMP 4.0], 2.15 declare reduction Directive, Restrictions
   // Only the variables omp_in and omp_out are allowed in the combiner.
   // Only the variables omp_priv and omp_orig are allowed in the
   // initializer-clause.
   auto *DRD = dyn_cast<OMPDeclareReductionDecl>(CurContext);
   if (LangOpts.OpenMP && DRD && !CurContext->containsDecl(D) &&

 
 /// BuildDeclRefExpr - Build an expression that references a
 /// declaration that does not require a closure capture.
 ExprResult
 Sema::BuildDeclRefExpr(ValueDecl *D, QualType Ty, ExprValueKind VK,
                        const DeclarationNameInfo &NameInfo,
                        const CXXScopeSpec *SS, NamedDecl *FoundD,
                        const TemplateArgumentListInfo *TemplateArgs) {
-  if (getLangOpts().CUDA)
-    if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(D))
-      if (!CheckCUDACall(NameInfo.getLoc(), Callee))
-        return ExprError();
-
   bool RefersToCapturedVariable =
       isa<VarDecl>(D) &&
       NeedToCaptureVariable(cast<VarDecl>(D), NameInfo.getLoc());
 
   DeclRefExpr *E;
   if (isa<VarTemplateSpecializationDecl>(D)) {
     VarTemplateSpecializationDecl *VarSpec =
         cast<VarTemplateSpecializationDecl>(D);

 static bool isNumberOfArgsValidForCall(Sema &S, const FunctionDecl *Callee,
                                        std::size_t NumArgs) {
   if (S.TooManyArguments(Callee->getNumParams(), NumArgs,
                          /*PartialOverloading=*/false))
     return Callee->isVariadic();
   return Callee->getMinRequiredArguments() <= NumArgs;
 }
 
-static ExprResult ActOnCallExprImpl(Sema &S, Scope *Scope, Expr *Fn,
-                                    SourceLocation LParenLoc,
-                                    MultiExprArg ArgExprs,
-                                    SourceLocation RParenLoc, Expr *ExecConfig,
-                                    bool IsExecConfig) {
+/// ActOnCallExpr - Handle a call to Fn with the specified array of arguments.
+/// This provides the location of the left/right parens and a list of comma
+/// locations.
+ExprResult Sema::ActOnCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc,
+                               MultiExprArg ArgExprs, SourceLocation RParenLoc,
+                               Expr *ExecConfig, bool IsExecConfig) {
   // Since this might be a postfix expression, get rid of ParenListExprs.
-  ExprResult Result = S.MaybeConvertParenListExprToParenExpr(Scope, Fn);
+  ExprResult Result = MaybeConvertParenListExprToParenExpr(Scope, Fn);
   if (Result.isInvalid()) return ExprError();
   Fn = Result.get();
 
-  if (checkArgsForPlaceholders(S, ArgExprs))
+  if (checkArgsForPlaceholders(*this, ArgExprs))
     return ExprError();
 
-  if (S.getLangOpts().CPlusPlus) {
+  if (getLangOpts().CPlusPlus) {
     // If this is a pseudo-destructor expression, build the call immediately.
     if (isa<CXXPseudoDestructorExpr>(Fn)) {
       if (!ArgExprs.empty()) {
         // Pseudo-destructor calls should not have any arguments.
-        S.Diag(Fn->getLocStart(), diag::err_pseudo_dtor_call_with_args)
+        Diag(Fn->getLocStart(), diag::err_pseudo_dtor_call_with_args)
             << FixItHint::CreateRemoval(
                    SourceRange(ArgExprs.front()->getLocStart(),
                                ArgExprs.back()->getLocEnd()));
       }
 
-      return new (S.Context)
-          CallExpr(S.Context, Fn, None, S.Context.VoidTy, VK_RValue, RParenLoc);
+      return new (Context)
+          CallExpr(Context, Fn, None, Context.VoidTy, VK_RValue, RParenLoc);
     }
-    if (Fn->getType() == S.Context.PseudoObjectTy) {
-      ExprResult result = S.CheckPlaceholderExpr(Fn);
+    if (Fn->getType() == Context.PseudoObjectTy) {
+      ExprResult result = CheckPlaceholderExpr(Fn);
       if (result.isInvalid()) return ExprError();
       Fn = result.get();
     }
 
     // Determine whether this is a dependent call inside a C++ template,
     // in which case we won't do any semantic analysis now.
     bool Dependent = false;
     if (Fn->isTypeDependent())
       Dependent = true;
     else if (Expr::hasAnyTypeDependentArguments(ArgExprs))
       Dependent = true;
 
     if (Dependent) {
       if (ExecConfig) {
-        return new (S.Context) CUDAKernelCallExpr(
-            S.Context, Fn, cast<CallExpr>(ExecConfig), ArgExprs,
-            S.Context.DependentTy, VK_RValue, RParenLoc);
+        return new (Context) CUDAKernelCallExpr(
+            Context, Fn, cast<CallExpr>(ExecConfig), ArgExprs,
+            Context.DependentTy, VK_RValue, RParenLoc);
       } else {
-        return new (S.Context)
-            CallExpr(S.Context, Fn, ArgExprs, S.Context.DependentTy, VK_RValue,
-                     RParenLoc);
+        return new (Context) CallExpr(
+            Context, Fn, ArgExprs, Context.DependentTy, VK_RValue, RParenLoc);
       }
     }
 
     // Determine whether this is a call to an object (C++ [over.call.object]).
     if (Fn->getType()->isRecordType())
-      return S.BuildCallToObjectOfClassType(Scope, Fn, LParenLoc, ArgExprs,
-                                            RParenLoc);
+      return BuildCallToObjectOfClassType(Scope, Fn, LParenLoc, ArgExprs,
 
-    if (Fn->getType() == S.Context.UnknownAnyTy) {
-      ExprResult result = rebuildUnknownAnyFunction(S, Fn);
+    if (Fn->getType() == Context.UnknownAnyTy) {
+      ExprResult result = rebuildUnknownAnyFunction(*this, Fn);
       if (result.isInvalid()) return ExprError();
       Fn = result.get();
     }
 
-    if (Fn->getType() == S.Context.BoundMemberTy) {
-      return S.BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                         RParenLoc);
+    if (Fn->getType() == Context.BoundMemberTy) {
+      return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
     }
   }
 
   // Check for overloaded calls.  This can happen even in C due to extensions.
-  if (Fn->getType() == S.Context.OverloadTy) {
+  if (Fn->getType() == Context.OverloadTy) {
     OverloadExpr::FindResult find = OverloadExpr::find(Fn);
 
     // We aren't supposed to apply this logic for if there'Scope an '&'
     // involved.
     if (!find.HasFormOfMemberPointer) {
       OverloadExpr *ovl = find.Expression;
       if (UnresolvedLookupExpr *ULE = dyn_cast<UnresolvedLookupExpr>(ovl))
-        return S.BuildOverloadedCallExpr(
+        return BuildOverloadedCallExpr(
             Scope, Fn, ULE, LParenLoc, ArgExprs, RParenLoc, ExecConfig,
             /*AllowTypoCorrection=*/true, find.IsAddressOfOperand);
-      return S.BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
-                                         RParenLoc);
+      return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
     }
   }
 
   // If we're directly calling a function, get the appropriate declaration.
-  if (Fn->getType() == S.Context.UnknownAnyTy) {
-    ExprResult result = rebuildUnknownAnyFunction(S, Fn);
+  if (Fn->getType() == Context.UnknownAnyTy) {
+    ExprResult result = rebuildUnknownAnyFunction(*this, Fn);
     if (result.isInvalid()) return ExprError();
     Fn = result.get();
   }
 
   Expr *NakedFn = Fn->IgnoreParens();
 
   bool CallingNDeclIndirectly = false;
   NamedDecl *NDecl = nullptr;
   if (UnaryOperator *UnOp = dyn_cast<UnaryOperator>(NakedFn)) {
     if (UnOp->getOpcode() == UO_AddrOf) {
       CallingNDeclIndirectly = true;
       NakedFn = UnOp->getSubExpr()->IgnoreParens();
     }
   }
 
   if (isa<DeclRefExpr>(NakedFn)) {
     NDecl = cast<DeclRefExpr>(NakedFn)->getDecl();
 
     FunctionDecl *FDecl = dyn_cast<FunctionDecl>(NDecl);
     if (FDecl && FDecl->getBuiltinID()) {
       // Rewrite the function decl for this builtin by replacing parameters
       // with no explicit address space with the address space of the arguments
       // in ArgExprs.
       if ((FDecl =
-               rewriteBuiltinFunctionDecl(&S, S.Context, FDecl, ArgExprs))) {
+               rewriteBuiltinFunctionDecl(this, Context, FDecl, ArgExprs))) {
         NDecl = FDecl;
         Fn = DeclRefExpr::Create(
-            S.Context, FDecl->getQualifierLoc(), SourceLocation(), FDecl, false,
+            Context, FDecl->getQualifierLoc(), SourceLocation(), FDecl, false,
             SourceLocation(), FDecl->getType(), Fn->getValueKind(), FDecl);
       }
     }
   } else if (isa<MemberExpr>(NakedFn))
     NDecl = cast<MemberExpr>(NakedFn)->getMemberDecl();
 
   if (FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(NDecl)) {
     if (CallingNDeclIndirectly &&
-        !S.checkAddressOfFunctionIsAvailable(FD, /*Complain=*/true,
-                                             Fn->getLocStart()))
+        !checkAddressOfFunctionIsAvailable(FD, /*Complain=*/true,
       return ExprError();
 
     // CheckEnableIf assumes that the we're passing in a sane number of args for
     // FD, but that doesn't always hold true here. This is because, in some
     // cases, we'll emit a diag about an ill-formed function call, but then
     // we'll continue on as if the function call wasn't ill-formed. So, if the
     // number of args looks incorrect, don't do enable_if checks; we should've
     // already emitted an error about the bad call.
     if (FD->hasAttr<EnableIfAttr>() &&
-        isNumberOfArgsValidForCall(S, FD, ArgExprs.size())) {
-      if (const EnableIfAttr *Attr = S.CheckEnableIf(FD, ArgExprs, true)) {
-        S.Diag(Fn->getLocStart(),
-               isa<CXXMethodDecl>(FD)
-                   ? diag::err_ovl_no_viable_member_function_in_call
-                   : diag::err_ovl_no_viable_function_in_call)
+        isNumberOfArgsValidForCall(*this, FD, ArgExprs.size())) {
+      if (const EnableIfAttr *Attr = CheckEnableIf(FD, ArgExprs, true)) {
+        Diag(Fn->getLocStart(),
             << FD << FD->getSourceRange();
-        S.Diag(FD->getLocation(),
-               diag::note_ovl_candidate_disabled_by_enable_if_attr)
+        Diag(FD->getLocation(),
             << Attr->getCond()->getSourceRange() << Attr->getMessage();
       }
     }
   }
 
-  return S.BuildResolvedCallExpr(Fn, NDecl, LParenLoc, ArgExprs, RParenLoc,
-                                 ExecConfig, IsExecConfig);
+  return BuildResolvedCallExpr(Fn, NDecl, LParenLoc, ArgExprs, RParenLoc,
 }
 
-/// ActOnCallExpr - Handle a call to Fn with the specified array of arguments.
-/// This provides the location of the left/right parens and a list of comma
-/// locations.
-ExprResult Sema::ActOnCallExpr(Scope *S, Expr *Fn, SourceLocation LParenLoc,
-                               MultiExprArg ArgExprs, SourceLocation RParenLoc,
-                               Expr *ExecConfig, bool IsExecConfig) {
-  ExprResult Ret = ActOnCallExprImpl(*this, S, Fn, LParenLoc, ArgExprs,
-                                     RParenLoc, ExecConfig, IsExecConfig);
-
-  // If appropriate, check that this is a valid CUDA call (and emit an error if
-  // the call is not allowed).
-  if (getLangOpts().CUDA && Ret.isUsable())
-    if (auto *Call = dyn_cast<CallExpr>(Ret.get()))
-      if (auto *FD = Call->getDirectCallee())
-        if (!CheckCUDACall(Call->getLocStart(), FD))
-          return ExprError();
-
-  return Ret;
-}
-
 /// ActOnAsTypeExpr - create a new asType (bitcast) from the arguments.
 ///
 /// __builtin_astype( value, dst type )
 ///
 ExprResult Sema::ActOnAsTypeExpr(Expr *E, ParsedType ParsedDestTy,
                                  SourceLocation BuiltinLoc,
                                  SourceLocation RParenLoc) {
   ExprValueKind VK = VK_RValue;

clang/test/SemaCUDA/Inputs/cuda.h

struct dim3 {
   __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
 };
 
 typedef struct cudaStream *cudaStream_t;
 
 int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
                       cudaStream_t stream = 0);
 
-// Device-side placement new overloads.
+// Host- and device-side placement new overloads.
+void *operator new(__SIZE_TYPE__, void *p) { return p; }
+void *operator new[](__SIZE_TYPE__, void *p) { return p; }
 __device__ void *operator new(__SIZE_TYPE__, void *p) { return p; }
 __device__ void *operator new[](__SIZE_TYPE__, void *p) { return p; }
 
 #endif // !__NVCC__

clang/test/SemaCUDA/call-host-fn-from-device.cu

 // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - -verify
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
 // are emitted during codegen.
 
 #include "Inputs/cuda.h"
 
 extern "C" void host_fn() {}
 // expected-note@-1 {{'host_fn' declared here}}
 // expected-note@-2 {{'host_fn' declared here}}
 // expected-note@-3 {{'host_fn' declared here}}
 // expected-note@-4 {{'host_fn' declared here}}
 // expected-note@-5 {{'host_fn' declared here}}
 // expected-note@-6 {{'host_fn' declared here}}
+// expected-note@-7 {{'host_fn' declared here}}
+
+struct Dummy {};
 
 struct S {
   S() {}
   // expected-note@-1 {{'S' declared here}}
   // expected-note@-2 {{'S' declared here}}
   ~S() { host_fn(); }
   // expected-note@-1 {{'~S' declared here}}
   int x;
 };
 
 struct T {
   __host__ __device__ void hd() { host_fn(); }
   // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
 
   // No error; this is (implicitly) inline and is never called, so isn't
   // codegen'ed.
   __host__ __device__ void hd2() { host_fn(); }
 
   __host__ __device__ void hd3();
 
   void h() {}
   // expected-note@-1 {{'h' declared here}}
+
+  void operator+();
+  // expected-note@-1 {{'operator+' declared here}}
+
+  void operator-(const T&) {}
+  // expected-note@-1 {{'operator-' declared here}}
+
+  operator Dummy() { return Dummy(); }
+  // expected-note@-1 {{'operator Dummy' declared here}}
 };
 
 __host__ __device__ void T::hd3() {
   host_fn();
   // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
 }
 
 template <typename T> __host__ __device__ void hd2() { host_fn(); }
__host__ __device__ void fn_ptr() {
   auto* ptr = &host_fn;
   // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
 }
 
 template <typename T>
 __host__ __device__ void fn_ptr_template() {
   auto* ptr = &host_fn;  // Not an error because the template isn't instantiated.
 }
+
+__host__ __device__ void unaryOp() {
+  T t;
+  (void) +t; // expected-error {{reference to __host__ function 'operator+' in __host__ __device__ function}}
+}
+
+__host__ __device__ void binaryOp() {
+  T t;
+  (void) (t - t); // expected-error {{reference to __host__ function 'operator-' in __host__ __device__ function}}
+}
+
+__host__ __device__ void implicitConversion() {
+  T t;
+  Dummy d = t; // expected-error {{reference to __host__ function 'operator Dummy' in __host__ __device__ function}}
+}
+
+template <typename T>
+struct TmplStruct {
+  template <typename U> __host__ __device__ void fn() {}
+};
+
+template <>
+template <>
+__host__ __device__ void TmplStruct<int>::fn<int>() { host_fn(); }
+// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
+
+__device__ void double_specialization() { TmplStruct<int>().fn<int>(); }