Merged master:3b4a0adec27 into amd-gfx:d08bdca3b40

Local branch amd-gfx d08bdca Merged master:0f03b2bfda9 into amd-gfx:145e568c385
Remote branch master 3b4a0ad [DWARFYAML][test] Use --ignore-case to suppress errors.

clang/include/clang/Frontend/CompilerInstance.h

@@ -764,10 +764,7 @@ class CompilerInstance : public ModuleLoader {
   static bool InitializeSourceManager(const FrontendInputFile &Input,
                                       DiagnosticsEngine &Diags,
                                       FileManager &FileMgr,
-                                      SourceManager &SourceMgr,
-                                      HeaderSearch *HS,
-                                      DependencyOutputOptions &DepOpts,
-                                      const FrontendOptions &Opts);
+                                      SourceManager &SourceMgr);
 
   /// }
 

clang/lib/Basic/Targets/AMDGPU.h

@@ -130,8 +130,26 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
         "exec_hi", "tma_lo", "tma_hi", "tba_lo", "tba_hi",
     });
 
+    switch (*Name) {
+    case 'I':
+      Info.setRequiresImmediate(-16, 64);
+      return true;
+    case 'J':
+      Info.setRequiresImmediate(-32768, 32767);
+      return true;
+    case 'A':
+    case 'B':
+    case 'C':
+      Info.setRequiresImmediate();
+      return true;
+    default:
+      break;
+    }
+
     StringRef S(Name);
-    if (S == "A") {
+
+    if (S == "DA" || S == "DB") {
+      Name++;
       Info.setRequiresImmediate();
       return true;
     }
@@ -203,6 +221,12 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
   // the constraint.  In practice, it won't be changed unless the
   // constraint is longer than one character.
   std::string convertConstraint(const char *&Constraint) const override {
+
+    StringRef S(Constraint);
+    if (S == "DA" || S == "DB") {
+      return std::string("^") + std::string(Constraint++, 2);
+    }
+
     const char *Begin = Constraint;
     TargetInfo::ConstraintInfo Info("", "");
     if (validateAsmConstraint(Constraint, Info))

clang/lib/Frontend/CompilerInstance.cpp

@@ -815,17 +815,15 @@ std::unique_ptr<llvm::raw_pwrite_stream> CompilerInstance::createOutputFile(
 // Initialization Utilities
 
 bool CompilerInstance::InitializeSourceManager(const FrontendInputFile &Input){
-  return InitializeSourceManager(
-      Input, getDiagnostics(), getFileManager(), getSourceManager(),
-      hasPreprocessor() ? &getPreprocessor().getHeaderSearchInfo() : nullptr,
-      getDependencyOutputOpts(), getFrontendOpts());
+  return InitializeSourceManager(Input, getDiagnostics(), getFileManager(),
+                                 getSourceManager());
 }
 
 // static
-bool CompilerInstance::InitializeSourceManager(
-    const FrontendInputFile &Input, DiagnosticsEngine &Diags,
-    FileManager &FileMgr, SourceManager &SourceMgr, HeaderSearch *HS,
-    DependencyOutputOptions &DepOpts, const FrontendOptions &Opts) {
+bool CompilerInstance::InitializeSourceManager(const FrontendInputFile &Input,
+                                               DiagnosticsEngine &Diags,
+                                               FileManager &FileMgr,
+                                               SourceManager &SourceMgr) {
   SrcMgr::CharacteristicKind Kind =
       Input.getKind().getFormat() == InputKind::ModuleMap
           ? Input.isSystem() ? SrcMgr::C_System_ModuleMap

clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl

@@ -33,3 +33,17 @@ kernel void test_agpr() {
          : "={a1}"(reg_a)
          : "{a1}"(reg_b));
 }
+
+kernel void test_constraint_DA() {
+  const long x = 0x200000001;
+  int res;
+  // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1 & 0xFFFFFFFF", "=v,^DA"(i64 8589934593)
+  __asm volatile("v_mov_b32 %0, %1 & 0xFFFFFFFF" : "=v"(res) : "DA"(x));
+}
+
+kernel void test_constraint_DB() {
+  const long x = 0x200000001;
+  int res;
+  // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1 & 0xFFFFFFFF", "=v,^DB"(i64 8589934593)
+  __asm volatile("v_mov_b32 %0, %1 & 0xFFFFFFFF" : "=v"(res) : "DB"(x));
+}

clang/test/Sema/inline-asm-validate-amdgpu.cl

@@ -18,9 +18,35 @@ kernel void test () {
   // vgpr constraints
   __asm__ ("v_mov_b32 %0, %1" : "=v" (vgpr) : "v" (imm) : );
 
-  // 'A' constraint
+  // 'I' constraint (an immediate integer in the range -16 to 64)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (imm) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (-16) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (64) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (-17) : ); // expected-error {{value '-17' out of range for constraint 'I'}}
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (65) : ); // expected-error {{value '65' out of range for constraint 'I'}}
+
+  // 'J' constraint (an immediate 16-bit signed integer)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (imm) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (-32768) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (32767) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (-32769) : ); // expected-error {{value '-32769' out of range for constraint 'J'}}
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (32768) : ); // expected-error {{value '32768' out of range for constraint 'J'}}
+
+  // 'A' constraint (an immediate constant that can be inlined)
   __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "A" (imm) : );
 
+  // 'B' constraint (an immediate 32-bit signed integer)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "B" (imm) : );
+
+  // 'C' constraint (an immediate 32-bit unsigned integer or 'A' constraint)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "C" (imm) : );
+
+  // 'DA' constraint (an immediate 64-bit constant that can be split into two 'A' constants)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "DA" (imm) : );
+
+  // 'DB' constraint (an immediate 64-bit constant that can be split into two 'B' constants)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "DB" (imm) : );
+
 }
 
 __kernel void

flang/runtime/file.cpp

@@ -352,7 +352,7 @@ bool OpenFile::RawSeekToEnd() {
 
 int OpenFile::PendingResult(const Terminator &terminator, int iostat) {
   int id{nextId_++};
-  pending_.reset(&New<Pending>{}(terminator, id, iostat, std::move(pending_)));
+  pending_ = New<Pending>{terminator}(id, iostat, std::move(pending_));
   return id;
 }
 } // namespace Fortran::runtime::io

flang/runtime/io-api.cpp

@@ -28,9 +28,10 @@ Cookie BeginInternalArrayListIO(const Descriptor &descriptor,
     void ** /*scratchArea*/, std::size_t /*scratchBytes*/,
     const char *sourceFile, int sourceLine) {
   Terminator oom{sourceFile, sourceLine};
-  return &New<InternalListIoStatementState<DIR>>{}(
-      oom, descriptor, sourceFile, sourceLine)
-              .ioStatementState();
+  return &New<InternalListIoStatementState<DIR>>{oom}(
+      descriptor, sourceFile, sourceLine)
+              .release()
+              ->ioStatementState();
 }
 
 Cookie IONAME(BeginInternalArrayListOutput)(const Descriptor &descriptor,
@@ -52,9 +53,10 @@ Cookie BeginInternalArrayFormattedIO(const Descriptor &descriptor,
     const char *format, std::size_t formatLength, void ** /*scratchArea*/,
     std::size_t /*scratchBytes*/, const char *sourceFile, int sourceLine) {
   Terminator oom{sourceFile, sourceLine};
-  return &New<InternalFormattedIoStatementState<DIR>>{}(
-      oom, descriptor, format, formatLength, sourceFile, sourceLine)
-              .ioStatementState();
+  return &New<InternalFormattedIoStatementState<DIR>>{oom}(
+      descriptor, format, formatLength, sourceFile, sourceLine)
+              .release()
+              ->ioStatementState();
 }
 
 Cookie IONAME(BeginInternalArrayFormattedOutput)(const Descriptor &descriptor,
@@ -78,9 +80,10 @@ Cookie BeginInternalFormattedIO(
     void ** /*scratchArea*/, std::size_t /*scratchBytes*/,
     const char *sourceFile, int sourceLine) {
   Terminator oom{sourceFile, sourceLine};
-  return &New<InternalFormattedIoStatementState<DIR>>{}(oom, internal,
-      internalLength, format, formatLength, sourceFile, sourceLine)
-              .ioStatementState();
+  return &New<InternalFormattedIoStatementState<DIR>>{oom}(
+      internal, internalLength, format, formatLength, sourceFile, sourceLine)
+              .release()
+              ->ioStatementState();
 }
 
 Cookie IONAME(BeginInternalFormattedOutput)(char *internal,
@@ -234,8 +237,9 @@ Cookie IONAME(BeginClose)(
   } else {
     // CLOSE(UNIT=bad unit) is just a no-op
     Terminator oom{sourceFile, sourceLine};
-    return &New<NoopCloseStatementState>{}(oom, sourceFile, sourceLine)
-                .ioStatementState();
+    return &New<NoopCloseStatementState>{oom}(sourceFile, sourceLine)
+                .release()
+                ->ioStatementState();
   }
 }
 

flang/runtime/memory.h

@@ -32,20 +32,32 @@ template <typename A> void FreeMemoryAndNullify(A *&p) {
   p = nullptr;
 }
 
-template <typename A> struct New {
-  template <typename... X>
-  [[nodiscard]] A &operator()(const Terminator &terminator, X &&... x) {
-    return *new (AllocateMemoryOrCrash(terminator, sizeof(A)))
-        A{std::forward<X>(x)...};
-  }
-};
-
 template <typename A> struct OwningPtrDeleter {
   void operator()(A *p) { FreeMemory(p); }
 };
 
 template <typename A> using OwningPtr = std::unique_ptr<A, OwningPtrDeleter<A>>;
 
+template <typename A> class SizedNew {
+public:
+  explicit SizedNew(const Terminator &terminator) : terminator_{terminator} {}
+  template <typename... X>
+  [[nodiscard]] OwningPtr<A> operator()(std::size_t bytes, X &&... x) {
+    return OwningPtr<A>{new (AllocateMemoryOrCrash(terminator_, bytes))
+            A{std::forward<X>(x)...}};
+  }
+
+private:
+  const Terminator &terminator_;
+};
+
+template <typename A> struct New : public SizedNew<A> {
+  using SizedNew<A>::SizedNew;
+  template <typename... X> [[nodiscard]] OwningPtr<A> operator()(X &&... x) {
+    return SizedNew<A>::operator()(sizeof(A), std::forward<X>(x)...);
+  }
+};
+
 template <typename A> struct Allocator {
   using value_type = A;
   explicit Allocator(const Terminator &t) : terminator{t} {}

flang/runtime/unit-map.cpp

@@ -64,7 +64,7 @@ void UnitMap::CloseAll(IoErrorHandler &handler) {
 }
 
 ExternalFileUnit &UnitMap::Create(int n, const Terminator &terminator) {
-  Chain &chain{New<Chain>{}(terminator, n)};
+  Chain &chain{*New<Chain>{terminator}(n).release()};
   chain.next.reset(&chain);
   bucket_[Hash(n)].swap(chain.next); // pushes new node as list head
   return chain.unit;

flang/runtime/unit.cpp

@@ -95,7 +95,7 @@ UnitMap &ExternalFileUnit::GetUnitMap() {
     return *unitMap;
   }
   Terminator terminator{__FILE__, __LINE__};
-  unitMap = &New<UnitMap>{}(terminator);
+  unitMap = New<UnitMap>{terminator}().release();
   ExternalFileUnit &out{ExternalFileUnit::LookUpOrCreate(6, terminator)};
   out.Predefine(1);
   out.set_mayRead(false);

lldb/source/Plugins/SymbolFile/DWARF/DWARFASTParserClang.cpp

@@ -2741,31 +2741,21 @@ void DWARFASTParserClang::ParseSingleMember(
 
           if (TypeSystemClang::IsCXXClassType(member_clang_type) &&
               !member_clang_type.GetCompleteType()) {
-            if (die.GetCU()->GetProducer() == eProducerClang)
-              module_sp->ReportError(
-                  "DWARF DIE at 0x%8.8x (class %s) has a member variable "
-                  "0x%8.8x (%s) whose type is a forward declaration, not a "
-                  "complete definition.\nTry compiling the source file "
-                  "with -fstandalone-debug",
-                  parent_die.GetOffset(), parent_die.GetName(), die.GetOffset(),
-                  name);
-            else
-              module_sp->ReportError(
-                  "DWARF DIE at 0x%8.8x (class %s) has a member variable "
-                  "0x%8.8x (%s) whose type is a forward declaration, not a "
-                  "complete definition.\nPlease file a bug against the "
-                  "compiler and include the preprocessed output for %s",
-                  parent_die.GetOffset(), parent_die.GetName(), die.GetOffset(),
-                  name, GetUnitName(parent_die).c_str());
-            // We have no choice other than to pretend that the member
-            // class is complete. If we don't do this, clang will crash
-            // when trying to layout the class. Since we provide layout
-            // assistance, all ivars in this class and other classes will
-            // be fine, this is the best we can do short of crashing.
+            // Mark the class as complete, ut we make a note of the fact that
+            // this class is not _really_ complete so we can later search for a
+            // definition in a different module.
+            // Since we provide layout assistance, all ivars in this class and
+            // other classes will be fine even if we are not able to find the
+            // definition elsewhere.
             if (TypeSystemClang::StartTagDeclarationDefinition(
                     member_clang_type)) {
               TypeSystemClang::CompleteTagDeclarationDefinition(
                   member_clang_type);
+              const auto *td = TypeSystemClang::GetQualType(
+                                   member_clang_type.GetOpaqueQualType())
+                                   .getTypePtr()
+                                   ->getAsTagDecl();
+              m_ast.GetMetadata(td)->SetIsForcefullyCompleted();
             } else {
               module_sp->ReportError(
                   "DWARF DIE at 0x%8.8x (class %s) has a member variable "