Merge branch 'main' into giselgetrounding

Author: JaydeepChauhan14

clang/Maintainers.rst

@@ -59,6 +59,8 @@ Clang MLIR generation
 | Bruno Cardoso Lopes
 | bruno.cardoso\@gmail.com (email), sonicsprawl (Discord), bcardosolopes (GitHub)
 
+| Henrich Lauko
+| henrich.lau\@gmail.com  (email), henrich.lauko (Discord), xlauko (GitHub)
 
 Analysis & CFG
 ~~~~~~~~~~~~~~

clang/include/clang/Basic/Attr.td

@@ -492,9 +492,22 @@ def TargetHasDLLImportExport : TargetSpec {
 def TargetItaniumCXXABI : TargetSpec {
   let CustomCode = [{ Target.getCXXABI().isItaniumFamily() }];
 }
+
 def TargetMicrosoftCXXABI : TargetArch<["x86", "x86_64", "arm", "thumb", "aarch64"]> {
   let CustomCode = [{ Target.getCXXABI().isMicrosoft() }];
 }
+
+// The target follows Microsoft record layout. Usually this happens in two
+// cases: 1. the target itself has Microsoft C++ ABI, e.g. x86_64 in MSVC
+// environment on Windows 2. an offloading target e.g. amdgcn or nvptx with
+// a host target in MSVC environment on Windows.
+def TargetMicrosoftRecordLayout : TargetArch<["x86", "x86_64", "arm", "thumb",
+                                              "aarch64", "amdgcn", "nvptx",
+                                              "nvptx64", "spirv", "spirv32",
+                                              "spirv64"]> {
+  let CustomCode = [{ Target.hasMicrosoftRecordLayout() }];
+}
+
 def TargetELF : TargetSpec {
   let ObjectFormats = ["ELF"];
 }
@@ -1789,7 +1802,7 @@ def Destructor : InheritableAttr {
   let Documentation = [CtorDtorDocs];
 }
 
-def EmptyBases : InheritableAttr, TargetSpecificAttr<TargetMicrosoftCXXABI> {
+def EmptyBases : InheritableAttr, TargetSpecificAttr<TargetMicrosoftRecordLayout> {
   let Spellings = [Declspec<"empty_bases">];
   let Subjects = SubjectList<[CXXRecord]>;
   let Documentation = [EmptyBasesDocs];
@@ -2021,7 +2034,7 @@ def Restrict : InheritableAttr {
   let Documentation = [RestrictDocs];
 }
 
-def LayoutVersion : InheritableAttr, TargetSpecificAttr<TargetMicrosoftCXXABI> {
+def LayoutVersion : InheritableAttr, TargetSpecificAttr<TargetMicrosoftRecordLayout> {
   let Spellings = [Declspec<"layout_version">];
   let Args = [UnsignedArgument<"Version">];
   let Subjects = SubjectList<[CXXRecord]>;
@@ -2239,7 +2252,7 @@ def NoUniqueAddress : InheritableAttr {
   let Spellings = [CXX11<"", "no_unique_address", 201803>, CXX11<"msvc", "no_unique_address", 201803>];
   let TargetSpecificSpellings = [
     TargetSpecificSpelling<TargetItaniumCXXABI, [CXX11<"", "no_unique_address", 201803>]>,
-    TargetSpecificSpelling<TargetMicrosoftCXXABI, [CXX11<"msvc", "no_unique_address", 201803>]>,
+    TargetSpecificSpelling<TargetMicrosoftRecordLayout, [CXX11<"msvc", "no_unique_address", 201803>]>,
   ];
   let Documentation = [NoUniqueAddressDocs];
 }

clang/include/clang/Basic/BuiltinsSPIRVCL.td

@@ -10,3 +10,6 @@ include "clang/Basic/BuiltinsSPIRVBase.td"
 
 def generic_cast_to_ptr_explicit
     : SPIRVBuiltin<"void*(void*, int)", [NoThrow, Const, CustomTypeChecking]>;
+def global_size : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def global_offset : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def subgroup_max_size : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;

clang/include/clang/Basic/BuiltinsSPIRVCommon.td

@@ -8,6 +8,16 @@
 
 include "clang/Basic/BuiltinsSPIRVBase.td"
 
+def num_workgroups : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def workgroup_size : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def workgroup_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def local_invocation_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def global_invocation_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def subgroup_size : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
+def num_subgroups : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
+def subgroup_id : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
+def subgroup_local_invocation_id : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
+
 def distance : SPIRVBuiltin<"void(...)", [NoThrow, Const]>;
 def length : SPIRVBuiltin<"void(...)", [NoThrow, Const]>;
 def smoothstep : SPIRVBuiltin<"void(...)", [NoThrow, Const, CustomTypeChecking]>;

clang/include/clang/Basic/TargetInfo.h

@@ -289,6 +289,8 @@ class TargetInfo : public TransferrableTargetInfo,
 
   std::optional<llvm::Triple> DarwinTargetVariantTriple;
 
+  bool HasMicrosoftRecordLayout = false;
+
   // TargetInfo Constructor.  Default initializes all fields.
   TargetInfo(const llvm::Triple &T);
 
@@ -1331,7 +1333,8 @@ class TargetInfo : public TransferrableTargetInfo,
   /// Apply changes to the target information with respect to certain
   /// language options which change the target configuration and adjust
   /// the language based on the target options where applicable.
-  virtual void adjust(DiagnosticsEngine &Diags, LangOptions &Opts);
+  virtual void adjust(DiagnosticsEngine &Diags, LangOptions &Opts,
+                      const TargetInfo *Aux);
 
   /// Initialize the map with the default set of target features for the
   /// CPU this should include all legal feature strings on the target.
@@ -1846,6 +1849,8 @@ class TargetInfo : public TransferrableTargetInfo,
 
   virtual void setAuxTarget(const TargetInfo *Aux) {}
 
+  bool hasMicrosoftRecordLayout() const { return HasMicrosoftRecordLayout; }
+
   /// Whether target allows debuginfo types for decl only variables/functions.
   virtual bool allowDebugInfoForExternalRef() const { return false; }
 

clang/lib/AST/ByteCode/Compiler.cpp

@@ -4912,6 +4912,18 @@ bool Compiler<Emitter>::VisitBuiltinCallExpr(const CallExpr *E,
   return true;
 }
 
+static const Expr *stripDerivedToBaseCasts(const Expr *E) {
+  if (const auto *PE = dyn_cast<ParenExpr>(E))
+    return stripDerivedToBaseCasts(PE->getSubExpr());
+
+  if (const auto *CE = dyn_cast<CastExpr>(E);
+      CE &&
+      (CE->getCastKind() == CK_DerivedToBase || CE->getCastKind() == CK_NoOp))
+    return stripDerivedToBaseCasts(CE->getSubExpr());
+
+  return E;
+}
+
 template <class Emitter>
 bool Compiler<Emitter>::VisitCallExpr(const CallExpr *E) {
   const FunctionDecl *FuncDecl = E->getDirectCallee();
@@ -4995,6 +5007,7 @@ bool Compiler<Emitter>::VisitCallExpr(const CallExpr *E) {
     }
   }
 
+  bool Devirtualized = false;
   std::optional<unsigned> CalleeOffset;
   // Add the (optional, implicit) This pointer.
   if (const auto *MC = dyn_cast<CXXMemberCallExpr>(E)) {
@@ -5013,8 +5026,26 @@ bool Compiler<Emitter>::VisitCallExpr(const CallExpr *E) {
         return false;
       if (!this->emitGetMemberPtrBase(E))
         return false;
-    } else if (!this->visit(MC->getImplicitObjectArgument())) {
-      return false;
+    } else {
+      const auto *InstancePtr = MC->getImplicitObjectArgument();
+      if (isa_and_nonnull<CXXDestructorDecl>(CompilingFunction) ||
+          isa_and_nonnull<CXXConstructorDecl>(CompilingFunction)) {
+        const auto *Stripped = stripDerivedToBaseCasts(InstancePtr);
+        if (isa<CXXThisExpr>(Stripped)) {
+          FuncDecl =
+              cast<CXXMethodDecl>(FuncDecl)->getCorrespondingMethodInClass(
+                  Stripped->getType()->getPointeeType()->getAsCXXRecordDecl());
+          Devirtualized = true;
+          if (!this->visit(Stripped))
+            return false;
+        } else {
+          if (!this->visit(InstancePtr))
+            return false;
+        }
+      } else {
+        if (!this->visit(InstancePtr))
+          return false;
+      }
     }
   } else if (const auto *PD =
                  dyn_cast<CXXPseudoDestructorExpr>(E->getCallee())) {
@@ -5060,7 +5091,7 @@ bool Compiler<Emitter>::VisitCallExpr(const CallExpr *E) {
 
     bool IsVirtual = false;
     if (const auto *MD = dyn_cast<CXXMethodDecl>(FuncDecl))
-      IsVirtual = MD->isVirtual();
+      IsVirtual = !Devirtualized && MD->isVirtual();
 
     // In any case call the function. The return value will end up on the stack
     // and if the function has RVO, we already have the pointer on the stack to
@@ -6027,6 +6058,8 @@ bool Compiler<Emitter>::visitFunc(const FunctionDecl *F) {
   // Classify the return type.
   ReturnType = this->classify(F->getReturnType());
 
+  this->CompilingFunction = F;
+
   if (const auto *Ctor = dyn_cast<CXXConstructorDecl>(F))
     return this->compileConstructor(Ctor);
   if (const auto *Dtor = dyn_cast<CXXDestructorDecl>(F))

clang/lib/AST/ByteCode/Compiler.h

@@ -448,6 +448,8 @@ class Compiler : public ConstStmtVisitor<Compiler<Emitter>, bool>,
   OptLabelTy ContinueLabel;
   /// Default case label.
   OptLabelTy DefaultLabel;
+
+  const FunctionDecl *CompilingFunction = nullptr;
 };
 
 extern template class Compiler<ByteCodeEmitter>;

clang/lib/AST/RecordLayoutBuilder.cpp

@@ -2458,15 +2458,6 @@ static bool mustSkipTailPadding(TargetCXXABI ABI, const CXXRecordDecl *RD) {
   llvm_unreachable("bad tail-padding use kind");
 }
 
-static bool isMsLayout(const ASTContext &Context) {
-  // Check if it's CUDA device compilation; ensure layout consistency with host.
-  if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice &&
-      Context.getAuxTargetInfo())
-    return Context.getAuxTargetInfo()->getCXXABI().isMicrosoft();
-
-  return Context.getTargetInfo().getCXXABI().isMicrosoft();
-}
-
 // This section contains an implementation of struct layout that is, up to the
 // included tests, compatible with cl.exe (2013).  The layout produced is
 // significantly different than those produced by the Itanium ABI.  Here we note
@@ -3399,7 +3390,7 @@ ASTContext::getASTRecordLayout(const RecordDecl *D) const {
 
   const ASTRecordLayout *NewEntry = nullptr;
 
-  if (isMsLayout(*this)) {
+  if (getTargetInfo().hasMicrosoftRecordLayout()) {
     if (const auto *RD = dyn_cast<CXXRecordDecl>(D)) {
       EmptySubobjectMap EmptySubobjects(*this, RD);
       MicrosoftRecordLayoutBuilder Builder(*this, &EmptySubobjects);
@@ -3656,7 +3647,8 @@ static void DumpRecordLayout(raw_ostream &OS, const RecordDecl *RD,
     bool HasOwnVBPtr = Layout.hasOwnVBPtr();
 
     // Vtable pointer.
-    if (CXXRD->isDynamicClass() && !PrimaryBase && !isMsLayout(C)) {
+    if (CXXRD->isDynamicClass() && !PrimaryBase &&
+        !C.getTargetInfo().hasMicrosoftRecordLayout()) {
       PrintOffset(OS, Offset, IndentLevel);
       OS << '(' << *RD << " vtable pointer)\n";
     } else if (HasOwnVFPtr) {
@@ -3754,7 +3746,7 @@ static void DumpRecordLayout(raw_ostream &OS, const RecordDecl *RD,
 
   PrintIndentNoOffset(OS, IndentLevel - 1);
   OS << "[sizeof=" << Layout.getSize().getQuantity();
-  if (CXXRD && !isMsLayout(C))
+  if (CXXRD && !C.getTargetInfo().hasMicrosoftRecordLayout())
     OS << ", dsize=" << Layout.getDataSize().getQuantity();
   OS << ", align=" << Layout.getAlignment().getQuantity();
   if (C.getTargetInfo().defaultsToAIXPowerAlignment())
@@ -3793,7 +3785,7 @@ void ASTContext::DumpRecordLayout(const RecordDecl *RD, raw_ostream &OS,
   OS << "\nLayout: ";
   OS << "<ASTRecordLayout\n";
   OS << "  Size:" << toBits(Info.getSize()) << "\n";
-  if (!isMsLayout(*this))
+  if (!getTargetInfo().hasMicrosoftRecordLayout())
     OS << "  DataSize:" << toBits(Info.getDataSize()) << "\n";
   OS << "  Alignment:" << toBits(Info.getAlignment()) << "\n";
   if (Target->defaultsToAIXPowerAlignment())

clang/lib/AST/TypePrinter.cpp

@@ -1861,6 +1861,17 @@ void TypePrinter::printAttributedBefore(const AttributedType *T,
   if (T->getAttrKind() == attr::ObjCKindOf)
     OS << "__kindof ";
 
+  if (T->getAttrKind() == attr::PreserveNone) {
+    OS << "__attribute__((preserve_none)) ";
+    spaceBeforePlaceHolder(OS);
+  } else if (T->getAttrKind() == attr::PreserveMost) {
+    OS << "__attribute__((preserve_most)) ";
+    spaceBeforePlaceHolder(OS);
+  } else if (T->getAttrKind() == attr::PreserveAll) {
+    OS << "__attribute__((preserve_all)) ";
+    spaceBeforePlaceHolder(OS);
+  }
+
   if (T->getAttrKind() == attr::AddressSpace)
     printBefore(T->getEquivalentType(), OS);
   else
@@ -1972,6 +1983,13 @@ void TypePrinter::printAttributedAfter(const AttributedType *T,
     return;
   }
 
+  if (T->getAttrKind() == attr::PreserveAll ||
+      T->getAttrKind() == attr::PreserveMost ||
+      T->getAttrKind() == attr::PreserveNone) {
+    // This has to be printed before the type.
+    return;
+  }
+
   OS << " __attribute__((";
   switch (T->getAttrKind()) {
 #define TYPE_ATTR(NAME)
@@ -2036,6 +2054,9 @@ void TypePrinter::printAttributedAfter(const AttributedType *T,
   case attr::Blocking:
   case attr::Allocating:
   case attr::SwiftAttr:
+  case attr::PreserveAll:
+  case attr::PreserveMost:
+  case attr::PreserveNone:
     llvm_unreachable("This attribute should have been handled already");
 
   case attr::NSReturnsRetained:
@@ -2071,20 +2092,12 @@ void TypePrinter::printAttributedAfter(const AttributedType *T,
   case attr::DeviceKernel:
     OS << T->getAttr()->getSpelling();
     break;
-  case attr::IntelOclBicc: OS << "inteloclbicc"; break;
-  case attr::PreserveMost:
-    OS << "preserve_most";
-    break;
-
-  case attr::PreserveAll:
-    OS << "preserve_all";
+  case attr::IntelOclBicc:
+    OS << "inteloclbicc";
     break;
   case attr::M68kRTD:
     OS << "m68k_rtd";
     break;
-  case attr::PreserveNone:
-    OS << "preserve_none";
-    break;
   case attr::RISCVVectorCC:
     OS << "riscv_vector_cc";
     break;

clang/lib/Basic/TargetInfo.cpp

@@ -176,6 +176,8 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : Triple(T) {
                     ? TargetCXXABI::Microsoft
                     : TargetCXXABI::GenericItanium);
 
+  HasMicrosoftRecordLayout = TheCXXABI.isMicrosoft();
+
   // Default to an empty address space map.
   AddrSpaceMap = &DefaultAddrSpaceMap;
   UseAddrSpaceMapMangling = false;
@@ -410,7 +412,8 @@ bool TargetInfo::isTypeSigned(IntType T) {
 /// Apply changes to the target information with respect to certain
 /// language options which change the target configuration and adjust
 /// the language based on the target options where applicable.
-void TargetInfo::adjust(DiagnosticsEngine &Diags, LangOptions &Opts) {
+void TargetInfo::adjust(DiagnosticsEngine &Diags, LangOptions &Opts,
+                        const TargetInfo *Aux) {
   if (Opts.NoBitFieldTypeAlign)
     UseBitFieldTypeAlignment = false;
 
@@ -550,6 +553,10 @@ void TargetInfo::adjust(DiagnosticsEngine &Diags, LangOptions &Opts) {
 
   if (Opts.FakeAddressSpaceMap)
     AddrSpaceMap = &FakeAddrSpaceMap;
+
+  // Check if it's CUDA device compilation; ensure layout consistency with host.
+  if (Opts.CUDA && Opts.CUDAIsDevice && Aux && !HasMicrosoftRecordLayout)
+    HasMicrosoftRecordLayout = Aux->getCXXABI().isMicrosoft();
 }
 
 bool TargetInfo::initFeatureMap(