merge main into amd-staging (llvm#1468)

Author: ronlieb

clang/docs/ReleaseNotes.rst

@@ -167,6 +167,7 @@ Non-comprehensive list of changes in this release
 
 - Support parsing the `cc` operand modifier and alias it to the `c` modifier (#GH127719).
 - Added `__builtin_elementwise_exp10`.
+- For AMDPGU targets, added `__builtin_v_cvt_off_f32_i4` that maps to the `v_cvt_off_f32_i4` instruction.
 
 New Compiler Flags
 ------------------

clang/include/clang/AST/StmtOpenACC.h

@@ -829,24 +829,42 @@ class OpenACCUpdateConstruct final
 
 // This class represents the 'atomic' construct, which has an associated
 // statement, but no clauses.
-class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
+class OpenACCAtomicConstruct final
+    : public OpenACCAssociatedStmtConstruct,
+      private llvm::TrailingObjects<OpenACCAtomicConstruct,
+                                    const OpenACCClause *> {
 
   friend class ASTStmtReader;
+  friend TrailingObjects;
   OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None;
 
-  OpenACCAtomicConstruct(EmptyShell)
+  OpenACCAtomicConstruct(unsigned NumClauses)
       : OpenACCAssociatedStmtConstruct(
             OpenACCAtomicConstructClass, OpenACCDirectiveKind::Atomic,
             SourceLocation{}, SourceLocation{}, SourceLocation{},
-            /*AssociatedStmt=*/nullptr) {}
+            /*AssociatedStmt=*/nullptr) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
 
   OpenACCAtomicConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
                          OpenACCAtomicKind AtKind, SourceLocation End,
+                         ArrayRef<const OpenACCClause *> Clauses,
                          Stmt *AssociatedStmt)
       : OpenACCAssociatedStmtConstruct(OpenACCAtomicConstructClass,
                                        OpenACCDirectiveKind::Atomic, Start,
                                        DirectiveLoc, End, AssociatedStmt),
-        AtomicKind(AtKind) {}
+        AtomicKind(AtKind) {
+    // Initialize the trailing storage.
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
 
   void setAssociatedStmt(Stmt *S) {
     OpenACCAssociatedStmtConstruct::setAssociatedStmt(S);
@@ -857,10 +875,12 @@ class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
     return T->getStmtClass() == OpenACCAtomicConstructClass;
   }
 
-  static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C);
+  static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C,
+                                             unsigned NumClauses);
   static OpenACCAtomicConstruct *
   Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
-         OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt);
+         OpenACCAtomicKind AtKind, SourceLocation End,
+         ArrayRef<const OpenACCClause *> Clauses, Stmt *AssociatedStmt);
 
   OpenACCAtomicKind getAtomicKind() const { return AtomicKind; }
   const Stmt *getAssociatedStmt() const {

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -140,6 +140,7 @@ BUILTIN(__builtin_amdgcn_cvt_pknorm_u16, "E2Usff", "nc")
 BUILTIN(__builtin_amdgcn_cvt_pk_i16, "E2sii", "nc")
 BUILTIN(__builtin_amdgcn_cvt_pk_u16, "E2UsUiUi", "nc")
 BUILTIN(__builtin_amdgcn_cvt_pk_u8_f32, "UifUiUi", "nc")
+BUILTIN(__builtin_amdgcn_cvt_off_f32_i4, "fi", "nc")
 BUILTIN(__builtin_amdgcn_sad_u8, "UiUiUiUi", "nc")
 BUILTIN(__builtin_amdgcn_msad_u8, "UiUiUiUi", "nc")
 BUILTIN(__builtin_amdgcn_sad_hi_u8, "UiUiUiUi", "nc")
@@ -254,7 +255,7 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s", "t", "at
 TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "atomic-global-pk-add-bf16-inst")
 TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atomic-ds-pk-add-16-insts")
 TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t", "atomic-ds-pk-add-16-insts")
-TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "gfx940-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "vmem-to-lds-load-insts")
 
 //===----------------------------------------------------------------------===//
 // Deep learning builtins.

clang/include/clang/Basic/DiagnosticGroups.td

@@ -377,12 +377,13 @@ def CXX11WarnSuggestOverride : DiagGroup<"suggest-override">;
 def WarnUnnecessaryVirtualSpecifier : DiagGroup<"unnecessary-virtual-specifier"> {
   code Documentation = [{
 Warns when a ``final`` class contains a virtual method (including virtual
-destructors) that does not override anything. Since ``final`` classes cannot be
-subclassed, their methods cannot be overridden, so there is no point to
-introducing new ``virtual`` methods.
+destructors). Since ``final`` classes cannot be subclassed, their methods
+cannot be overridden, and hence the ``virtual`` specifier is useless.
 
 The warning also detects virtual methods in classes whose destructor is
 ``final``, for the same reason.
+
+The warning does not fire on virtual methods which are also marked ``override``.
   }];
 }
 

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -2733,7 +2733,7 @@ def note_final_dtor_non_final_class_silence : Note<
   "mark %0 as '%select{final|sealed}1' to silence this warning">;
 def warn_unnecessary_virtual_specifier : Warning<
   "virtual method %0 is inside a 'final' class and can never be overridden">,
-  InGroup<WarnUnnecessaryVirtualSpecifier>;
+  InGroup<WarnUnnecessaryVirtualSpecifier>, DefaultIgnore;
 
 // C++11 attributes
 def err_repeat_attribute : Error<"%0 attribute cannot be repeated">;

clang/lib/AST/StmtOpenACC.cpp

@@ -307,20 +307,26 @@ OpenACCUpdateConstruct::Create(const ASTContext &C, SourceLocation Start,
 }
 
 OpenACCAtomicConstruct *
-OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C) {
-  void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
-  auto *Inst = new (Mem) OpenACCAtomicConstruct(EmptyShell{});
+OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C, unsigned NumClauses) {
+  void *Mem = C.Allocate(
+      OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCAtomicConstruct(NumClauses);
   return Inst;
 }
 
 OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create(
     const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
-    OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt) {
-  void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
-  auto *Inst = new (Mem)
-      OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind, End, AssociatedStmt);
+    OpenACCAtomicKind AtKind, SourceLocation End,
+    ArrayRef<const OpenACCClause *> Clauses, Stmt *AssociatedStmt) {
+  void *Mem = C.Allocate(
+      OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst = new (Mem) OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind,
+                                                End, Clauses, AssociatedStmt);
   return Inst;
 }
+
 OpenACCCacheConstruct *OpenACCCacheConstruct::CreateEmpty(const ASTContext &C,
                                                           unsigned NumVars) {
   void *Mem =

clang/lib/AST/StmtPrinter.cpp

@@ -1258,6 +1258,7 @@ void StmtPrinter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
   if (S->getAtomicKind() != OpenACCAtomicKind::None)
     OS << " " << S->getAtomicKind();
 
+  PrintOpenACCClauseList(S);
   OS << '\n';
   PrintStmt(S->getAssociatedStmt());
 }

clang/lib/AST/StmtProfile.cpp

@@ -2839,6 +2839,8 @@ void StmtProfiler::VisitOpenACCUpdateConstruct(
 void StmtProfiler::VisitOpenACCAtomicConstruct(
     const OpenACCAtomicConstruct *S) {
   VisitStmt(S);
+  OpenACCClauseProfiler P{*this};
+  P.VisitOpenACCClauseList(S->clauses());
 }
 
 void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {

clang/lib/Basic/Targets/AMDGPU.cpp

@@ -260,7 +260,7 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple,
 
   MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
   CUMode = !(GPUFeatures & llvm::AMDGPU::FEATURE_WGP);
-  for (auto F : {"image-insts", "gws"})
+  for (auto F : {"image-insts", "gws", "vmem-to-lds-load-insts"})
     ReadOnlyFeatures.insert(F);
   HalfArgsAndReturns = true;
 }

clang/lib/Sema/SemaOpenACC.cpp

@@ -1832,9 +1832,8 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
                                           EndLoc, Clauses);
   }
   case OpenACCDirectiveKind::Atomic: {
-    assert(Clauses.empty() && "Atomic doesn't allow clauses");
     return OpenACCAtomicConstruct::Create(
-        getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc,
+        getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc, Clauses,
         AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
   }
   case OpenACCDirectiveKind::Cache: {