[Sema] Produce diagnostics when unavailable aligned allocation/deallocation functions are called

[ahatanak]

Summary

Deallocation functions are called from deleting destructors that are declared virtual. Issue an error if an aligned deallocation function is selected but is not available.

[EricWF]

I think we should sink the DiagnoseUnavailableAlignedAllocation into DiagnoseUseOfDecl, and then adds call's in the few places they're needed. Most paths we care about already pass through there.

include/clang/Sema/Sema.h

5169

DiagnoseUnavailableAlignedAllocation.

Also, I think we can drop the IsDelete parameter and instead deduce it using FD.getDeclName().getCXXOverloadedOperator().

[ahatanak]

Sink diagnoseUnavailableAlignedAllocation into DiagnoseUseOfDecl and add calls to it in a few other places so that diagnostics are produced for calls to aligned operator and builtin operator new/delete in addition to aligned deallocation functions called from deleting destructors.

I had to make changes to two test cases (dr2xx.cpp and call-host-fn-from-device.cu) that have nothing to do with aligned allocation/deallocation functions. The warning in dr2xx.cpp seems correct to me judging from the comment left few lines above it ("We're also missing the -Wused-but-marked-unused "). I'm not sure about the diagnostic in call-host-fn-from-device.cu. The original comment says the sized delete function annotated with __device__ is called, but it seems that the non-sized __host__ function is being called, in which case I think the diagnostic is correct.

[ahatanak]

ping.

@jlebar, is the change I made to call-host-fn-from-device.cu correct?

[jlebar]

@jlebar, is the change I made to call-host-fn-from-device.cu correct?

I don't think so -- that's a change in overloading behavior afaict.

[ahatanak]

The original comment said "call sized device delete even though host has preferable non-sized version", but it seems that the non-sized host version 'T::operator delete(void*)' is being called in the IR, not the sized device version of delete. Is that a bug in overload resolution?

This is the command I used:

"clang -cc1 -internal-isystem --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - test.cu"

$ cat test.cu

#define __constant__ __attribute__((constant))
#define __device__ __attribute__((device))
#define __host__ __attribute__((host))

typedef __SIZE_TYPE__ size_t;

struct T {
  __host__ void operator delete(void*);
  __device__ void operator delete(void*, size_t);
};

__host__ __device__ void class_specific_delete(T *t) {
  delete t;
}

[ahatanak]

I mean ToT clang (without my patch applied) seems to select the non-sized host version 'T::operator delete(void*)'.

[jlebar]

In D47757#1142886, @ahatanak wrote:

I mean ToT clang (without my patch applied) seems to select the non-sized host version 'T::operator delete(void*)'.

OK, if this is just making an error out of something which previously silently didn't work (and should result in a compile error further down the line when we try to and can't resolve that function), then this is totally fine.

[ahatanak]

Yes, that is the case. It doesn't change overloading resolution, it is just producing a diagnostic.

[ahatanak]

ping

[ahatanak]

ping

[ahatanak]

ping

[ahatanak]

ping

[jlebar]

+tra in the hopes that perhaps he's comfortable reviewing this (sorry that I'm not).

[tra]

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

88

The C++ magic is way above my paygrade, but as far as CUDA goes this is a regression, compared to what nvcc does. This code in NVCC produced a warning and clang should not error out at this point in time either as
it's not an error to call a host function from HD unless we use HD in a host function, and we would not know how it's used until later. I think the error should be postponed until codegen.

[rsmith]

lib/Sema/SemaExprCXX.cpp

3463

Are we also missing a MarkFunctionReferenced call here? (I don't think it matters much for the predefined new/delete, since they can't be inline or templated, but it's still wrong to not mark the function the builtin will call as referenced.)

test/CXX/drs/dr2xx.cpp

721–722

These are -> This is

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

88

We're in -fcuda-is-device mode, so IIUC it's correct to reject a call to a host function here (because __host__ __device__ is treated as basically meaning __device__ in that mode for the purpose of checking whether a call is valid), right?

However, the comment suggests that the intent was that this would instead call the device version. Did that actually previously happen (in which case this patch is somehow affecting overload resolution and should be fixed), or is the comment prior to this patch wrong and we were silently calling a host function from a device function (in which case this patch is fine, but we should add a FIXME here to select the device delete function if we think that's appropriate)?

[rsmith]

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

88

OK, I see from prior review comments (that phab is helpfully hiding from view) that this is just adding a diagnostic and the overload resolution behavior is unchanged. So I think this change is correct. @tra, can you confirm? My testing shows that

__host__ void f(); __host__ __device__ void g() { f(); }

is accepted by default but rejected in -fcuda-is-device mode, which is consistent with the behavior after this patch is applied.

[ahatanak]

Call MarkFunctionReferenced to mark the function __builtin_operator_new or __builtin_operator_delete will call as referenced.

[tra]

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

88

"-fcuda-is-device" does not necessarily mean that the host device function will be used.

In the example above the error is indeed correct, as g() is considered to be externally visible and we will attempt to generate code for it, and we can't call f() on device.

However, if you make it static, there should be no error:

__host__ void f(); 
static __host__ __device__ void g() { f(); }

CUDA is somewhat weird when it comes to what's considered available and what is not.
If you want some of the gory details, see D12453 and https://goo.gl/EXnymm

@jlebar has details on how we handle the errors in cases like that:
https://github.com/llvm-mirror/clang/blob/master/lib/Sema/SemaCUDA.cpp#L590

[ahatanak]

@tra and @rsmith: Can we move forward and fix the incorrect cuda diagnostics in a separate patch?

[tra]

In D47757#1204545, @ahatanak wrote:

@tra and @rsmith: Can we move forward and fix the incorrect cuda diagnostics in a separate patch?

Doing that in a separate patch is OK, provided that that patch will be committed along with this one.

It's a regression. There's a decent chance it breaks someone and this patch, if committed by itself, will end up being rolled back.

[ahatanak]

In D47757#1204561, @tra wrote:

It's a regression. There's a decent chance it breaks someone and this patch, if committed by itself, will end up being rolled back.

Is the regression you are referring to about the static function case? I don't see a difference between ToT clang and my patch in the diagnostics they produce when I compile the following code:

__host__ void f();
static __host__ __device__ void g() { f(); }
__host__ __device__ void g2() { g(); }

Both error out when -fcuda-is-device is provided. If I comment out the definition of g2, it compiles fine.

[tra]

In D47757#1204621, @ahatanak wrote:

In D47757#1204561, @tra wrote:

It's a regression. There's a decent chance it breaks someone and this patch, if committed by itself, will end up being rolled back.

Is the regression you are referring to about the static function case? I don't see a difference between ToT clang and my patch in the diagnostics they produce when I compile the following code:

__host__ void f();
static __host__ __device__ void g() { f(); }
__host__ __device__ void g2() { g(); }

Both error out when -fcuda-is-device is provided. If I comment out the definition of g2, it compiles fine.

The example above *is* expected to produce the error on device side, bacause g2() is externally visible, uses g(), which in turn uses host-only f().

I'm talking about a case where g() {f()} is present in the source code, but will not be codegen'ed on device side.

The code below is expected to compile. Note that g2() is host-only.

__host__ void f(); 
static __host__ __device__ void g() { f(); }
__host__ void g2() { g(); }

[ahatanak]

The code you showed does compile with or without -fcuda-is-device after applying my patch.

[tra]

Talked to @ahatanak over IRC. It appears that this patch may have exposed a preexisting bug.
Apparently delete t; in test/SemaCUDA/call-host-fn-from-device.cu does actually end up calling __host__ operator delete. It should've picked __device__ operator delete, but it does not, so reporting an error here appears to be correct.

It's visible in AST and the IR.

@rsmith -- the original change was done a while back in rL283830. I assume it worked at that time and wonder if it's a (possibly not-so-)recent regression.

[tra]

I've confirmed that the patch does not break anything in our CUDA code, so it's good to go as far as CUDA is concerned.

I'll fix the exposed CUDA issue in a separate patch.

[ahatanak]

In D47757#1211276, @tra wrote:

I've confirmed that the patch does not break anything in our CUDA code, so it's good to go as far as CUDA is concerned.

Thanks. @rsmith, do you have any other comments about the patch?

[ahatanak]

Rebase & ping.

I've reverted the changes I made to test/SemaCUDA/call-host-fn-from-device.cu since r342749 fixed the overload resolution bug.

[ahatanak]

ping

[rsmith]

lib/Sema/SemaExpr.cpp

53

Does this also need to be updated?

[ahatanak]

Check whether the declaration passed to Sema::CanUseDecl is an aligned allocation/deallocation function that is unavailable.

[vsapsai]

The change itself looks correct. Cannot really tell if you need to make changes in other places. For that I rely on Richard's opinion.

include/clang/Sema/Sema.h

ExprResult BuildCXXNew(SourceRange Range, bool UseGlobal,
                          SourceLocation PlacementRParen,
                          SourceRange TypeIdParens,
                          QualType AllocType,
                          TypeSourceInfo *AllocTypeInfo,
                          Expr *ArraySize,
                          SourceRange DirectInitRange,
                          Expr *Initializer);
 
+  /// Produce diagnostics if FD is an aligned allocation or deallocation
+  /// function that is unavailable.
+  void diagnoseUnavailableAlignedAllocation(const FunctionDecl &FD,
+                                            SourceLocation Loc);
+
   bool CheckAllocatedType(QualType AllocType, SourceLocation Loc,
                           SourceRange R);
 
   /// The scope in which to find allocation functions.
   enum AllocationFunctionScope {
     /// Only look for allocation functions in the global scope.
     AFS_Global,
     /// Only look for allocation functions in the scope of the

lib/Sema/SemaDeclCXX.cpp

if (FunctionDecl *OperatorDelete =
             // in the right order.
             Diag(Loc, diag::note_implicit_delete_this_in_destructor_here);
             return true;
           }
           ThisArg = This.get();
         }
       }
 
+      DiagnoseUseOfDecl(OperatorDelete, Loc);
       MarkFunctionReferenced(Loc, OperatorDelete);
       Destructor->setOperatorDelete(OperatorDelete, ThisArg);
     }
   }
 
   return false;
 }
 

lib/Sema/SemaExpr.cpp

 #include "clang/Sema/SemaInternal.h"
 #include "clang/Sema/Template.h"
 #include "llvm/Support/ConvertUTF.h"
 using namespace clang;
 using namespace sema;
 
 /// Determine whether the use of this declaration is valid, without
 /// emitting diagnostics.
 bool Sema::CanUseDecl(NamedDecl *D, bool TreatUnavailableAsInvalid) {
   // See if this is an auto-typed variable whose initializer we are parsing.
   if (ParsingInitForAutoVars.count(D))
     return false;
 
   // See if this is a deleted function.
   if (FunctionDecl *FD = dyn_cast<FunctionDecl>(D)) {
     if (FD->isDeleted())
       return false;
if (Pos != SuppressedDiagnostics.end()) {
       // diagnostics again.
       Pos->second.clear();
     }
 
     // C++ [basic.start.main]p3:
     //   The function 'main' shall not be used within a program.
     if (cast<FunctionDecl>(D)->isMain())
       Diag(Loc, diag::ext_main_used);
+
+    diagnoseUnavailableAlignedAllocation(*cast<FunctionDecl>(D), Loc);
   }
 
   // See if this is an auto-typed variable whose initializer we are parsing.
   if (ParsingInitForAutoVars.count(D)) {
     if (isa<BindingDecl>(D)) {
       Diag(Loc, diag::err_binding_cannot_appear_in_own_initializer)
         << D->getDeclName();
     } else {

lib/Sema/SemaExprCXX.cpp

assert(isa<InitListExpr>(Init) &&
            "Shouldn't create list CXXConstructExprs for arrays.");
     return true;
   }
   return false;
 }
 
 // Emit a diagnostic if an aligned allocation/deallocation function that is not
 // implemented in the standard library is selected.
-static void diagnoseUnavailableAlignedAllocation(const FunctionDecl &FD,
-                                                 SourceLocation Loc, bool IsDelete,
-                                                 Sema &S) {
-  if (!S.getLangOpts().AlignedAllocationUnavailable)
+void Sema::diagnoseUnavailableAlignedAllocation(const FunctionDecl &FD,
+                                                SourceLocation Loc) {
+  if (!getLangOpts().AlignedAllocationUnavailable)
     return;
 
   // Return if there is a definition.
   if (FD.isDefined())
     return;
 
   bool IsAligned = false;
   if (FD.isReplaceableGlobalAllocationFunction(&IsAligned) && IsAligned) {
-    const llvm::Triple &T = S.getASTContext().getTargetInfo().getTriple();
+    const llvm::Triple &T = getASTContext().getTargetInfo().getTriple();
     StringRef OSName = AvailabilityAttr::getPlatformNameSourceSpelling(
-        S.getASTContext().getTargetInfo().getPlatformName());
-
-    S.Diag(Loc, diag::warn_aligned_allocation_unavailable)
+        getASTContext().getTargetInfo().getPlatformName());
+    OverloadedOperatorKind Kind = FD.getDeclName().getCXXOverloadedOperator();
+    bool IsDelete = Kind == OO_Delete || Kind == OO_Array_Delete;
+    Diag(Loc, diag::warn_aligned_allocation_unavailable)
          << IsDelete << FD.getType().getAsString() << OSName
          << alignedAllocMinVersion(T.getOS()).getAsString();
-    S.Diag(Loc, diag::note_silence_unligned_allocation_unavailable);
+    Diag(Loc, diag::note_silence_unligned_allocation_unavailable);
   }
 }
 
 ExprResult
 Sema::BuildCXXNew(SourceRange Range, bool UseGlobal,
                   SourceLocation PlacementLParen,
                   MultiExprArg PlacementArgs,
                   SourceLocation PlacementRParen,
if (!AllocType->isDependentType() &&
     Initializer = FullInit.get();
   }
 
   // Mark the new and delete operators as referenced.
   if (OperatorNew) {
     if (DiagnoseUseOfDecl(OperatorNew, StartLoc))
       return ExprError();
     MarkFunctionReferenced(StartLoc, OperatorNew);
-    diagnoseUnavailableAlignedAllocation(*OperatorNew, StartLoc, false, *this);
   }
   if (OperatorDelete) {
     if (DiagnoseUseOfDecl(OperatorDelete, StartLoc))
       return ExprError();
     MarkFunctionReferenced(StartLoc, OperatorDelete);
-    diagnoseUnavailableAlignedAllocation(*OperatorDelete, StartLoc, true, *this);
   }
 
   // C++0x [expr.new]p17:
   //   If the new expression creates an array of objects of class type,
   //   access and ambiguity control are done for the destructor.
   QualType BaseAllocType = Context.getBaseElementType(AllocType);
   if (ArraySize && !BaseAllocType->isDependentType()) {
     if (const RecordType *BaseRecordType = BaseAllocType->getAs<RecordType>()) {
if (!Ex.get()->isTypeDependent()) {
     if (PointeeRD) {
       if (CXXDestructorDecl *Dtor = LookupDestructor(PointeeRD)) {
         CheckDestructorAccess(Ex.get()->getExprLoc(), Dtor,
                               PDiag(diag::err_access_dtor) << PointeeElem);
         IsVirtualDelete = Dtor->isVirtual();
       }
     }
 
-    diagnoseUnavailableAlignedAllocation(*OperatorDelete, StartLoc, true,
-                                         *this);
+    DiagnoseUseOfDecl(OperatorDelete, StartLoc);
 
     // Convert the operand to the type of the first parameter of operator
     // delete. This is only necessary if we selected a destroying operator
     // delete that we are going to call (non-virtually); converting to void*
     // is trivial and left to AST consumers to handle.
     QualType ParamType = OperatorDelete->getParamDecl(0)->getType();
     if (!IsVirtualDelete && !ParamType->getPointeeType()->isVoidType()) {
       Qualifiers Qs = Pointee.getQualifiers();
Sema::SemaBuiltinOperatorNewDeleteOverloaded(ExprResult TheCallResult,
   DeclareGlobalNewDelete();
 
   FunctionDecl *OperatorNewOrDelete = nullptr;
   if (resolveBuiltinNewDeleteOverload(*this, TheCall, IsDelete,
                                       OperatorNewOrDelete))
     return ExprError();
   assert(OperatorNewOrDelete && "should be found");
 
+  DiagnoseUseOfDecl(OperatorNewOrDelete, TheCall->getExprLoc());
+
   TheCall->setType(OperatorNewOrDelete->getReturnType());
   for (unsigned i = 0; i != TheCall->getNumArgs(); ++i) {
     QualType ParamTy = OperatorNewOrDelete->getParamDecl(i)->getType();
     InitializedEntity Entity =
         InitializedEntity::InitializeParameter(Context, ParamTy, false);
     ExprResult Arg = PerformCopyInitialization(
         Entity, TheCall->getArg(i)->getLocStart(), TheCall->getArg(i));
     if (Arg.isInvalid())

test/CXX/drs/dr2xx.cpp

#pragma clang diagnostic warning "-Wused-but-marked-unused"
   // translation unit.
   // We're also missing the -Wused-but-marked-unused diagnostic here.
   struct A {
     inline void *operator new(size_t) __attribute__((unused));
     inline void operator delete(void*) __attribute__((unused));
     A() {}
   };
 
   // FIXME: These are ill-formed, with a required diagnostic, for the same
   // reason.
   struct B {
     inline void operator delete(void*) __attribute__((unused));
     ~B() {}
   };
   struct C {
     inline void operator delete(void*) __attribute__((unused));
-    virtual ~C() {}
+    virtual ~C() {} // expected-warning {{'operator delete' was marked unused but was used}}
   };
 
   struct D {
     inline void operator delete(void*) __attribute__((unused));
   };
   void h() { C::operator delete(0); } // expected-warning {{marked unused but was used}}
 
 #pragma clang diagnostic pop

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

 }
 
 __host__ __device__ void explicit_destructor(S *s) {
   s->~S();
   // expected-error@-1 {{reference to __host__ function '~S' in __host__ __device__ function}}
 }
 
 __host__ __device__ void class_specific_delete(T *t, U *u) {
-  delete t; // ok, call sized device delete even though host has preferable non-sized version
+  delete t; // expected-error {{reference to __host__ function 'operator delete' in __host__ __device__ function}}
   delete u; // ok, call non-sized HD delete rather than sized D delete
 }
 
 __host__ __device__ void hd_member_fn() {
   T t;
   // Necessary to trigger an error on T::hd.  It's (implicitly) inline, so
   // isn't codegen'ed until we call it.
   t.hd();

test/SemaCXX/unavailable_aligned_allocation.cpp

 #else
 // expected-error@-16 {{aligned allocation function of type 'void *(unsigned long, enum std::align_val_t)' is only available on macOS 10.13 or newer}}}
 // expected-error@-17 {{aligned deallocation function of type 'void (void *, enum std::align_val_t) noexcept' is only available on macOS 10.13 or newer}}}
 #endif
 
 // expected-note@-20 2 {{if you supply your own aligned allocation functions}}
 #endif
 
-// No errors if user-defined aligned allocation functions are available.
+// Test that diagnostics are produced when an unavailable aligned deallocation
+// function is called from a deleting destructor.
+struct alignas(256) OveralignedS2 {
+  int a[4];
+  virtual ~OveralignedS2();
+};
+
+OveralignedS2::~OveralignedS2() {}
+
+#ifdef NO_ERRORS
+// expected-no-diagnostics
+#else
+#if defined(IOS)
+// expected-error@-6 {{aligned deallocation function of type 'void (void *, enum std::align_val_t) noexcept' is only available on iOS 11 or newer}}}
+// expected-note@-7 {{if you supply your own aligned allocation functions}}
+#elif defined(TVOS)
+// expected-error@-9 {{aligned deallocation function of type 'void (void *, enum std::align_val_t) noexcept' is only available on tvOS 11 or newer}}}
+// expected-note@-10 {{if you supply your own aligned allocation functions}}
+#elif defined(WATCHOS)
+// expected-error@-12 {{aligned deallocation function of type 'void (void *, enum std::align_val_t) noexcept' is only available on watchOS 4 or newer}}}
+// expected-note@-13 {{if you supply your own aligned allocation functions}}
+#else
+// expected-error@-15 {{aligned deallocation function of type 'void (void *, enum std::align_val_t) noexcept' is only available on macOS 10.13 or newer}}}
+// expected-note@-16 {{if you supply your own aligned allocation functions}}
+#endif
+#endif
+
+void testExplicitOperatorNewDelete() {
+  void *p = operator new(128);
+  operator delete(p);
+  p = operator new[](128);
+  operator delete[](p);
+  p = __builtin_operator_new(128);
+  __builtin_operator_delete(p);
+}
+
+void testExplicitOperatorNewDeleteOveraligned() {
+  void *p = operator new(128, (std::align_val_t)64);
+  operator delete(p, (std::align_val_t)64);
+  p = operator new[](128, (std::align_val_t)64);
+  operator delete[](p, (std::align_val_t)64);
+  p = __builtin_operator_new(128, (std::align_val_t)64);
+  __builtin_operator_delete(p, (std::align_val_t)64);
+}
+
+#ifdef NO_ERRORS
+// expected-no-diagnostics
+#else
+// expected-error@-11 {{aligned allocation function of type 'void *(unsigned long, enum std::align_val_t)' is only available on}}
+// expected-note@-12 {{if you supply your own aligned allocation functions}}
+
+// expected-error@-13 {{aligned deallocation function of type 'void (void *, enum std::align_val_t) noexcept' is only available on}}
+// expected-note@-14 {{if you supply your own aligned allocation functions}}
+
+// expected-error@-15 {{aligned allocation function of type 'void *(unsigned long, enum std::align_val_t)' is only available on}}
+// expected-note@-16 {{if you supply your own aligned allocation functions}}
+
+// expected-error@-17 {{aligned deallocation function of type 'void (void *, enum std::align_val_t) noexcept' is only available on}}
+// expected-note@-18 {{if you supply your own aligned allocation functions}}
+
+// expected-error@-19 {{aligned allocation function of type 'void *(unsigned long, enum std::align_val_t)' is only available on}}
+// expected-note@-20 {{if you supply your own aligned allocation functions}}
+
+// expected-error@-21 {{aligned deallocation function of type 'void (void *, enum std::align_val_t) noexcept' is only available on}}
+// expected-note@-22 {{if you supply your own aligned allocation functions}}
+#endif
+
 void *operator new(std::size_t __sz, std::align_val_t) {
   static char array[256];
   return &array;
 }
 
 void operator delete(void *p, std::align_val_t) {
 }
 
 void testOveraligned2() {
   auto p = new ((std::align_val_t)8) OveralignedS;
   delete p;
 }