Merge branch 'main' into vector_gather_doc_improvement

Author: newling

clang/include/clang/APINotes/Types.h

@@ -141,6 +141,9 @@ class CommonTypeInfo : public CommonEntityInfo {
   /// The NS error domain for this type.
   std::optional<std::string> NSErrorDomain;
 
+  /// The Swift protocol that this type should be automatically conformed to.
+  std::optional<std::string> SwiftConformance;
+
 public:
   CommonTypeInfo() {}
 
@@ -165,6 +168,14 @@ class CommonTypeInfo : public CommonEntityInfo {
                            : std::nullopt;
   }
 
+  std::optional<std::string> getSwiftConformance() const {
+    return SwiftConformance;
+  }
+
+  void setSwiftConformance(std::optional<std::string> conformance) {
+    SwiftConformance = conformance;
+  }
+
   friend bool operator==(const CommonTypeInfo &, const CommonTypeInfo &);
 
   CommonTypeInfo &operator|=(const CommonTypeInfo &RHS) {
@@ -175,6 +186,8 @@ class CommonTypeInfo : public CommonEntityInfo {
       setSwiftBridge(RHS.getSwiftBridge());
     if (!NSErrorDomain)
       setNSErrorDomain(RHS.getNSErrorDomain());
+    if (SwiftConformance)
+      setSwiftConformance(RHS.getSwiftConformance());
 
     return *this;
   }
@@ -185,7 +198,8 @@ class CommonTypeInfo : public CommonEntityInfo {
 inline bool operator==(const CommonTypeInfo &LHS, const CommonTypeInfo &RHS) {
   return static_cast<const CommonEntityInfo &>(LHS) == RHS &&
          LHS.SwiftBridge == RHS.SwiftBridge &&
-         LHS.NSErrorDomain == RHS.NSErrorDomain;
+         LHS.NSErrorDomain == RHS.NSErrorDomain &&
+         LHS.SwiftConformance == RHS.SwiftConformance;
 }
 
 inline bool operator!=(const CommonTypeInfo &LHS, const CommonTypeInfo &RHS) {
@@ -739,9 +753,6 @@ class TagInfo : public CommonTypeInfo {
   std::optional<std::string> SwiftReleaseOp;
   std::optional<std::string> SwiftDefaultOwnership;
 
-  /// The Swift protocol that this type should be automatically conformed to.
-  std::optional<std::string> SwiftConformance;
-
   std::optional<EnumExtensibilityKind> EnumExtensibility;
 
   TagInfo()
@@ -790,9 +801,6 @@ class TagInfo : public CommonTypeInfo {
     if (!SwiftDefaultOwnership)
       SwiftDefaultOwnership = RHS.SwiftDefaultOwnership;
 
-    if (!SwiftConformance)
-      SwiftConformance = RHS.SwiftConformance;
-
     if (!HasFlagEnum)
       setFlagEnum(RHS.isFlagEnum());
 
@@ -819,7 +827,6 @@ inline bool operator==(const TagInfo &LHS, const TagInfo &RHS) {
          LHS.SwiftRetainOp == RHS.SwiftRetainOp &&
          LHS.SwiftReleaseOp == RHS.SwiftReleaseOp &&
          LHS.SwiftDefaultOwnership == RHS.SwiftDefaultOwnership &&
-         LHS.SwiftConformance == RHS.SwiftConformance &&
          LHS.isFlagEnum() == RHS.isFlagEnum() &&
          LHS.isSwiftCopyable() == RHS.isSwiftCopyable() &&
          LHS.isSwiftEscapable() == RHS.isSwiftEscapable() &&

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -503,6 +503,9 @@ TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vv*i", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vv*i", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vv*", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "vIs", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_get_named_barrier_state, "Uiv*", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_prefetch_data, "vvC*Ui", "nc", "gfx12-insts")

clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h

@@ -212,6 +212,13 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
     return create<cir::AllocaOp>(loc, addrType, type, name, alignment);
   }
 
+  /// Get constant address of a global variable as an MLIR attribute.
+  cir::GlobalViewAttr getGlobalViewAttr(cir::PointerType type,
+                                        cir::GlobalOp globalOp) {
+    auto symbol = mlir::FlatSymbolRefAttr::get(globalOp.getSymNameAttr());
+    return cir::GlobalViewAttr::get(type, symbol);
+  }
+
   mlir::Value createGetGlobal(mlir::Location loc, cir::GlobalOp global) {
     assert(!cir::MissingFeatures::addressSpace());
     return create<cir::GetGlobalOp>(loc, getPointerTo(global.getSymType()),

clang/include/clang/CIR/Dialect/IR/CIRAttrs.td

@@ -370,6 +370,53 @@ def CIR_ConstPtrAttr : CIR_Attr<"ConstPtr", "ptr", [TypedAttrInterface]> {
   }];
 }
 
+//===----------------------------------------------------------------------===//
+// GlobalViewAttr
+//===----------------------------------------------------------------------===//
+
+def CIR_GlobalViewAttr : CIR_Attr<"GlobalView", "global_view", [
+  TypedAttrInterface
+]> {
+  let summary = "Provides constant access to a global address";
+  let description = [{
+    Get constant address of global `symbol`. It provides a way to access globals
+    from other global and always produces a pointer.
+
+    The type of the input symbol can be different from `#cir.global_view`
+    output type, since a given view of the global might require a static
+    cast for initializing other globals.
+
+    The result type of this attribute may be an integer type. In such a case,
+    the pointer to the referenced global is casted to an integer and this
+    attribute represents the casted result.
+
+    Example:
+
+    ```
+      cir.global external @s = @".str2": !cir.ptr<i8>
+      cir.global external @x = #cir.global_view<@s> : !cir.ptr<i8>
+      cir.global external @s_addr = #cir.global_view<@s> : !s64i
+    ```
+  }];
+
+  let parameters = (ins AttributeSelfTypeParameter<"">:$type,
+                        "mlir::FlatSymbolRefAttr":$symbol);
+
+  let builders = [
+    AttrBuilderWithInferredContext<(ins "mlir::Type":$type,
+                                        "mlir::FlatSymbolRefAttr":$symbol), [{
+      return $_get(type.getContext(), type, symbol);
+    }]>
+  ];
+
+  // let genVerifyDecl = 1;
+  let assemblyFormat = [{
+    `<`
+      $symbol
+    `>`
+  }];
+}
+
 //===----------------------------------------------------------------------===//
 // ConstComplexAttr
 //===----------------------------------------------------------------------===//

clang/include/clang/CIR/MissingFeatures.h

@@ -192,6 +192,7 @@ struct MissingFeatures {
   static bool constantFoldSwitchStatement() { return false; }
   static bool constructABIArgDirectExtend() { return false; }
   static bool coverageMapping() { return false; }
+  static bool createInvariantGroup() { return false; }
   static bool createProfileWeightsForLoop() { return false; }
   static bool ctorMemcpyizer() { return false; }
   static bool cudaSupport() { return false; }
@@ -200,6 +201,8 @@ struct MissingFeatures {
   static bool deferredCXXGlobalInit() { return false; }
   static bool ehCleanupFlags() { return false; }
   static bool ehCleanupScope() { return false; }
+  static bool ehCleanupScopeRequiresEHCleanup() { return false; }
+  static bool ehCleanupBranchFixups() { return false; }
   static bool ehstackBranches() { return false; }
   static bool emitCheckedInBoundsGEP() { return false; }
   static bool emitCondLikelihoodViaExpectIntrinsic() { return false; }
@@ -210,9 +213,12 @@ struct MissingFeatures {
   static bool fastMathFlags() { return false; }
   static bool fpConstraints() { return false; }
   static bool generateDebugInfo() { return false; }
+  static bool globalViewIndices() { return false; }
+  static bool globalViewIntLowering() { return false; }
   static bool hip() { return false; }
   static bool implicitConstructorArgs() { return false; }
   static bool incrementProfileCounter() { return false; }
+  static bool innermostEHScope() { return false; }
   static bool insertBuiltinUnpredictable() { return false; }
   static bool instrumentation() { return false; }
   static bool intrinsics() { return false; }
@@ -233,7 +239,6 @@ struct MissingFeatures {
   static bool objCGC() { return false; }
   static bool objCLifetime() { return false; }
   static bool openMP() { return false; }
-  static bool opGlobalViewAttr() { return false; }
   static bool opTBAA() { return false; }
   static bool peepholeProtection() { return false; }
   static bool pgoUse() { return false; }
@@ -260,6 +265,7 @@ struct MissingFeatures {
   static bool appleKext() { return false; }
   static bool dtorCleanups() { return false; }
   static bool vtableInitialization() { return false; }
+  static bool vtableRelativeLayout() { return false; }
   static bool msvcBuiltins() { return false; }
   static bool vlas() { return false; }
 

clang/include/clang/Interpreter/RemoteJITUtils.h

@@ -26,8 +26,7 @@
 
 llvm::Expected<std::unique_ptr<llvm::orc::SimpleRemoteEPC>>
 launchExecutor(llvm::StringRef ExecutablePath, bool UseSharedMemory,
-               llvm::StringRef SlabAllocateSizeString,
-               std::function<void()> CustomizeFork = nullptr);
+               llvm::StringRef SlabAllocateSizeString);
 
 /// Create a JITLinkExecutor that connects to the given network address
 /// through a TCP socket. A valid NetworkAddress provides hostname and port,
@@ -36,13 +35,4 @@ llvm::Expected<std::unique_ptr<llvm::orc::SimpleRemoteEPC>>
 connectTCPSocket(llvm::StringRef NetworkAddress, bool UseSharedMemory,
                  llvm::StringRef SlabAllocateSizeString);
 
-#ifdef LLVM_ON_UNIX
-/// Returns PID of last launched executor.
-pid_t getLastLaunchedExecutorPID();
-
-/// Returns PID of nth launched executor.
-/// 1-based indexing.
-pid_t getNthLaunchedExecutorPID(int n);
-#endif
-
 #endif // LLVM_CLANG_INTERPRETER_REMOTEJITUTILS_H

clang/lib/APINotes/APINotesFormat.h

@@ -24,7 +24,7 @@ const uint16_t VERSION_MAJOR = 0;
 /// API notes file minor version number.
 ///
 /// When the format changes IN ANY WAY, this number should be incremented.
-const uint16_t VERSION_MINOR = 35; // SwiftDefaultOwnership
+const uint16_t VERSION_MINOR = 36; // Typedef SwiftConformsTo
 
 const uint8_t kSwiftConforms = 1;
 const uint8_t kSwiftDoesNotConform = 2;

clang/lib/APINotes/APINotesReader.cpp

@@ -134,6 +134,13 @@ void ReadCommonTypeInfo(const uint8_t *&Data, CommonTypeInfo &Info) {
         reinterpret_cast<const char *>(Data), ErrorDomainLength - 1)));
     Data += ErrorDomainLength - 1;
   }
+
+  if (unsigned ConformanceLength =
+          endian::readNext<uint16_t, llvm::endianness::little>(Data)) {
+    Info.setSwiftConformance(std::string(reinterpret_cast<const char *>(Data),
+                                         ConformanceLength - 1));
+    Data += ConformanceLength - 1;
+  }
 }
 
 /// Used to deserialize the on-disk identifier table.
@@ -629,12 +636,6 @@ class TagTableInfo
           reinterpret_cast<const char *>(Data), DefaultOwnershipLength - 1);
       Data += DefaultOwnershipLength - 1;
     }
-    if (unsigned ConformanceLength =
-            endian::readNext<uint16_t, llvm::endianness::little>(Data)) {
-      Info.SwiftConformance = std::string(reinterpret_cast<const char *>(Data),
-                                          ConformanceLength - 1);
-      Data += ConformanceLength - 1;
-    }
 
     ReadCommonTypeInfo(Data, Info);
     return Info;

clang/lib/APINotes/APINotesWriter.cpp

@@ -536,7 +536,8 @@ unsigned getCommonEntityInfoSize(const CommonEntityInfo &CEI) {
 // in on-disk hash tables.
 unsigned getCommonTypeInfoSize(const CommonTypeInfo &CTI) {
   return 2 + (CTI.getSwiftBridge() ? CTI.getSwiftBridge()->size() : 0) + 2 +
-         (CTI.getNSErrorDomain() ? CTI.getNSErrorDomain()->size() : 0) +
+         (CTI.getNSErrorDomain() ? CTI.getNSErrorDomain()->size() : 0) + 2 +
+         (CTI.getSwiftConformance() ? CTI.getSwiftConformance()->size() : 0) +
          getCommonEntityInfoSize(CTI);
 }
 
@@ -557,6 +558,12 @@ void emitCommonTypeInfo(raw_ostream &OS, const CommonTypeInfo &CTI) {
   } else {
     writer.write<uint16_t>(0);
   }
+  if (auto conformance = CTI.getSwiftConformance()) {
+    writer.write<uint16_t>(conformance->size() + 1);
+    OS.write(conformance->c_str(), conformance->size());
+  } else {
+    writer.write<uint16_t>(0);
+  }
 }
 
 /// Used to serialize the on-disk Objective-C property table.
@@ -1274,7 +1281,6 @@ class TagTableInfo : public CommonTypeTableInfo<TagTableInfo, TagInfo> {
            2 + (TI.SwiftRetainOp ? TI.SwiftRetainOp->size() : 0) +
            2 + (TI.SwiftReleaseOp ? TI.SwiftReleaseOp->size() : 0) +
            2 + (TI.SwiftDefaultOwnership ? TI.SwiftDefaultOwnership->size() : 0) +
-           2 + (TI.SwiftConformance ? TI.SwiftConformance->size() : 0) +
            3 + getCommonTypeInfoSize(TI);
     // clang-format on
   }
@@ -1328,12 +1334,6 @@ class TagTableInfo : public CommonTypeTableInfo<TagTableInfo, TagInfo> {
     } else {
       writer.write<uint16_t>(0);
     }
-    if (auto Conformance = TI.SwiftConformance) {
-      writer.write<uint16_t>(Conformance->size() + 1);
-      OS.write(Conformance->c_str(), Conformance->size());
-    } else {
-      writer.write<uint16_t>(0);
-    }
 
     emitCommonTypeInfo(OS, TI);
   }

clang/lib/APINotes/APINotesYAMLCompiler.cpp

@@ -251,6 +251,7 @@ struct Class {
   std::optional<StringRef> NSErrorDomain;
   std::optional<bool> SwiftImportAsNonGeneric;
   std::optional<bool> SwiftObjCMembers;
+  std::optional<std::string> SwiftConformance;
   MethodsSeq Methods;
   PropertiesSeq Properties;
 };
@@ -275,6 +276,7 @@ template <> struct MappingTraits<Class> {
     IO.mapOptional("NSErrorDomain", C.NSErrorDomain);
     IO.mapOptional("SwiftImportAsNonGeneric", C.SwiftImportAsNonGeneric);
     IO.mapOptional("SwiftObjCMembers", C.SwiftObjCMembers);
+    IO.mapOptional("SwiftConformsTo", C.SwiftConformance);
     IO.mapOptional("Methods", C.Methods);
     IO.mapOptional("Properties", C.Properties);
   }
@@ -525,6 +527,7 @@ struct Typedef {
   std::optional<StringRef> SwiftBridge;
   std::optional<StringRef> NSErrorDomain;
   std::optional<SwiftNewTypeKind> SwiftType;
+  std::optional<std::string> SwiftConformance;
 };
 
 typedef std::vector<Typedef> TypedefsSeq;
@@ -553,6 +556,7 @@ template <> struct MappingTraits<Typedef> {
     IO.mapOptional("SwiftBridge", T.SwiftBridge);
     IO.mapOptional("NSErrorDomain", T.NSErrorDomain);
     IO.mapOptional("SwiftWrapper", T.SwiftType);
+    IO.mapOptional("SwiftConformsTo", T.SwiftConformance);
   }
 };
 } // namespace yaml
@@ -802,6 +806,8 @@ class YAMLConverter {
     if (Common.SwiftBridge)
       Info.setSwiftBridge(std::string(*Common.SwiftBridge));
     Info.setNSErrorDomain(Common.NSErrorDomain);
+    if (auto conformance = Common.SwiftConformance)
+      Info.setSwiftConformance(conformance);
   }
 
   // Translate from Method into ObjCMethodInfo and write it out.
@@ -990,8 +996,6 @@ class YAMLConverter {
       TI.SwiftRetainOp = T.SwiftRetainOp;
     if (T.SwiftReleaseOp)
       TI.SwiftReleaseOp = T.SwiftReleaseOp;
-    if (T.SwiftConformance)
-      TI.SwiftConformance = T.SwiftConformance;
     if (T.SwiftDefaultOwnership)
       TI.SwiftDefaultOwnership = T.SwiftDefaultOwnership;