Initial Commit

thirdparty/capstone/suite/synctools/tablegen/include/llvm/Frontend/OpenMP/OMPAssume.h

@@ -0,0 +1,55 @@
+//===- OpenMP/OMPAssume.h --- OpenMP assumption helper functions  - C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+/// \file
+///
+/// This file provides helper functions and classes to deal with OpenMP
+/// assumptions, e.g., as used by `[begin/end] assumes` and `assume`.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_FRONTEND_OPENMP_OMPASSUME_H
+#define LLVM_FRONTEND_OPENMP_OMPASSUME_H
+
+#include "llvm/ADT/StringRef.h"
+
+namespace llvm {
+
+namespace omp {
+
+/// Helper to describe assume clauses.
+struct AssumptionClauseMappingInfo {
+  /// The identifier describing the (beginning of the) clause.
+  llvm::StringLiteral Identifier;
+  /// Flag to determine if the identifier is a full name or the start of a name.
+  bool StartsWith;
+  /// Flag to determine if a directive lists follows.
+  bool HasDirectiveList;
+  /// Flag to determine if an expression follows.
+  bool HasExpression;
+};
+
+/// All known assume clauses.
+static constexpr AssumptionClauseMappingInfo AssumptionClauseMappings[] = {
+#define OMP_ASSUME_CLAUSE(Identifier, StartsWith, HasDirectiveList,            \
+                          HasExpression)                                       \
+  {Identifier, StartsWith, HasDirectiveList, HasExpression},
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+inline std::string getAllAssumeClauseOptions() {
+  std::string S;
+  for (const AssumptionClauseMappingInfo &ACMI : AssumptionClauseMappings)
+    S += (S.empty() ? "'" : "', '") + ACMI.Identifier.str();
+  return S + "'";
+}
+
+} // namespace omp
+
+} // namespace llvm
+
+#endif // LLVM_FRONTEND_OPENMP_OMPASSUME_H

thirdparty/capstone/suite/synctools/tablegen/include/llvm/Frontend/OpenMP/OMPConstants.h

@@ -0,0 +1,123 @@
+//===- OMPConstants.h - OpenMP related constants and helpers ------ C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+/// \file
+///
+/// This file defines constants and helpers used when dealing with OpenMP.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_FRONTEND_OPENMP_OMPCONSTANTS_H
+#define LLVM_FRONTEND_OPENMP_OMPCONSTANTS_H
+
+#include "llvm/ADT/BitmaskEnum.h"
+
+#include "llvm/ADT/StringRef.h"
+#include "llvm/Frontend/OpenMP/OMP.h.inc"
+
+namespace llvm {
+namespace omp {
+LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
+
+/// IDs for all Internal Control Variables (ICVs).
+enum class InternalControlVar {
+#define ICV_DATA_ENV(Enum, ...) Enum,
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+#define ICV_DATA_ENV(Enum, ...)                                                \
+  constexpr auto Enum = omp::InternalControlVar::Enum;
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+
+enum class ICVInitValue {
+#define ICV_INIT_VALUE(Enum, Name) Enum,
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+#define ICV_INIT_VALUE(Enum, Name)                                             \
+  constexpr auto Enum = omp::ICVInitValue::Enum;
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+
+/// IDs for all omp runtime library (RTL) functions.
+enum class RuntimeFunction {
+#define OMP_RTL(Enum, ...) Enum,
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+#define OMP_RTL(Enum, ...) constexpr auto Enum = omp::RuntimeFunction::Enum;
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+
+/// IDs for the different default kinds.
+enum class DefaultKind {
+#define OMP_DEFAULT_KIND(Enum, Str) Enum,
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+#define OMP_DEFAULT_KIND(Enum, ...)                                            \
+  constexpr auto Enum = omp::DefaultKind::Enum;
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+
+/// IDs for all omp runtime library ident_t flag encodings (see
+/// their definition in openmp/runtime/src/kmp.h).
+enum class IdentFlag {
+#define OMP_IDENT_FLAG(Enum, Str, Value) Enum = Value,
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+  LLVM_MARK_AS_BITMASK_ENUM(0x7FFFFFFF)
+};
+
+#define OMP_IDENT_FLAG(Enum, ...) constexpr auto Enum = omp::IdentFlag::Enum;
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+
+/// \note This needs to be kept in sync with kmp.h enum sched_type.
+/// Todo: Update kmp.h to include this file, and remove the enums in kmp.h
+///       To complete this, more enum values will need to be moved here.
+enum class OMPScheduleType {
+  StaticChunked = 33,
+  Static = 34, // static unspecialized
+  DistributeChunked = 91,
+  Distribute = 92,
+  DynamicChunked = 35,
+  GuidedChunked = 36, // guided unspecialized
+  Runtime = 37,
+  Auto = 38, // auto
+
+  StaticBalancedChunked = 45, // static with chunk adjustment (e.g., simd)
+  GuidedSimd = 46,            // guided with chunk adjustment
+  RuntimeSimd = 47,           // runtime with chunk adjustment
+
+  ModifierMonotonic =
+      (1 << 29), // Set if the monotonic schedule modifier was present
+  ModifierNonmonotonic =
+      (1 << 30), // Set if the nonmonotonic schedule modifier was present
+  ModifierMask = ModifierMonotonic | ModifierNonmonotonic,
+  LLVM_MARK_AS_BITMASK_ENUM(/* LargestValue */ ModifierMask)
+};
+
+enum OMPTgtExecModeFlags : int8_t {
+  OMP_TGT_EXEC_MODE_GENERIC = 1 << 0,
+  OMP_TGT_EXEC_MODE_SPMD = 1 << 1,
+  OMP_TGT_EXEC_MODE_GENERIC_SPMD =
+      OMP_TGT_EXEC_MODE_GENERIC | OMP_TGT_EXEC_MODE_SPMD,
+  LLVM_MARK_AS_BITMASK_ENUM(/* LargestValue */ OMP_TGT_EXEC_MODE_GENERIC_SPMD)
+};
+
+enum class AddressSpace : unsigned {
+  Generic = 0,
+  Global = 1,
+  Shared = 3,
+  Constant = 4,
+  Local = 5,
+};
+
+/// \note This needs to be kept in sync with interop.h enum kmp_interop_type_t.:
+enum class OMPInteropType { Unknown, Target, TargetSync };
+
+} // end namespace omp
+
+} // end namespace llvm
+
+#endif // LLVM_FRONTEND_OPENMP_OMPCONSTANTS_H

thirdparty/capstone/suite/synctools/tablegen/include/llvm/Frontend/OpenMP/OMPContext.h

@@ -0,0 +1,210 @@
+//===- OpenMP/OMPContext.h ----- OpenMP context helper functions  - C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+/// \file
+///
+/// This file provides helper functions and classes to deal with OpenMP
+/// contexts as used by `[begin/end] declare variant` and `metadirective`.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_FRONTEND_OPENMP_OMPCONTEXT_H
+#define LLVM_FRONTEND_OPENMP_OMPCONTEXT_H
+
+#include "llvm/ADT/APSInt.h"
+#include "llvm/ADT/BitVector.h"
+#include "llvm/ADT/SetVector.h"
+#include "llvm/ADT/SmallSet.h"
+#include "llvm/ADT/Triple.h"
+#include "llvm/Frontend/OpenMP/OMPConstants.h"
+
+namespace llvm {
+namespace omp {
+
+/// OpenMP Context related IDs and helpers
+///
+///{
+
+/// IDs for all OpenMP context selector trait sets (construct/device/...).
+enum class TraitSet {
+#define OMP_TRAIT_SET(Enum, ...) Enum,
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+/// IDs for all OpenMP context selector trait (device={kind/isa...}/...).
+enum class TraitSelector {
+#define OMP_TRAIT_SELECTOR(Enum, ...) Enum,
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+/// IDs for all OpenMP context trait properties (host/gpu/bsc/llvm/...)
+enum class TraitProperty {
+#define OMP_TRAIT_PROPERTY(Enum, ...) Enum,
+#define OMP_LAST_TRAIT_PROPERTY(Enum) Last = Enum
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+/// Parse \p Str and return the trait set it matches or TraitSet::invalid.
+TraitSet getOpenMPContextTraitSetKind(StringRef Str);
+
+/// Return the trait set for which \p Selector is a selector.
+TraitSet getOpenMPContextTraitSetForSelector(TraitSelector Selector);
+
+/// Return the trait set for which \p Property is a property.
+TraitSet getOpenMPContextTraitSetForProperty(TraitProperty Property);
+
+/// Return a textual representation of the trait set \p Kind.
+StringRef getOpenMPContextTraitSetName(TraitSet Kind);
+
+/// Parse \p Str and return the trait set it matches or
+/// TraitSelector::invalid.
+TraitSelector getOpenMPContextTraitSelectorKind(StringRef Str);
+
+/// Return the trait selector for which \p Property is a property.
+TraitSelector getOpenMPContextTraitSelectorForProperty(TraitProperty Property);
+
+/// Return a textual representation of the trait selector \p Kind.
+StringRef getOpenMPContextTraitSelectorName(TraitSelector Kind);
+
+/// Parse \p Str and return the trait property it matches in the set \p Set and
+/// selector \p Selector or TraitProperty::invalid.
+TraitProperty getOpenMPContextTraitPropertyKind(TraitSet Set,
+                                                TraitSelector Selector,
+                                                StringRef Str);
+
+/// Return the trait property for a singleton selector \p Selector.
+TraitProperty getOpenMPContextTraitPropertyForSelector(TraitSelector Selector);
+
+/// Return a textual representation of the trait property \p Kind, which might
+/// be the raw string we parsed (\p RawString) if we do not translate the
+/// property into a (distinct) enum.
+StringRef getOpenMPContextTraitPropertyName(TraitProperty Kind,
+                                            StringRef RawString);
+
+/// Return a textual representation of the trait property \p Kind with selector
+/// and set name included.
+StringRef getOpenMPContextTraitPropertyFullName(TraitProperty Kind);
+
+/// Return a string listing all trait sets.
+std::string listOpenMPContextTraitSets();
+
+/// Return a string listing all trait selectors for \p Set.
+std::string listOpenMPContextTraitSelectors(TraitSet Set);
+
+/// Return a string listing all trait properties for \p Set and \p Selector.
+std::string listOpenMPContextTraitProperties(TraitSet Set,
+                                             TraitSelector Selector);
+///}
+
+/// Return true if \p Selector can be nested in \p Set. Also sets
+/// \p AllowsTraitScore and \p RequiresProperty to true/false if the user can
+/// specify a score for properties in \p Selector and if the \p Selector
+/// requires at least one property.
+bool isValidTraitSelectorForTraitSet(TraitSelector Selector, TraitSet Set,
+                                     bool &AllowsTraitScore,
+                                     bool &RequiresProperty);
+
+/// Return true if \p Property can be nested in \p Selector and \p Set.
+bool isValidTraitPropertyForTraitSetAndSelector(TraitProperty Property,
+                                                TraitSelector Selector,
+                                                TraitSet Set);
+
+/// Variant match information describes the required traits and how they are
+/// scored (via the ScoresMap). In addition, the required construct nesting is
+/// described as well.
+struct VariantMatchInfo {
+  /// Add the trait \p Property to the required trait set. \p RawString is the
+  /// string we parsed and derived \p Property from. If \p Score is not null, it
+  /// recorded as well. If \p Property is in the `construct` set it is recorded
+  /// in-order in the ConstructTraits as well.
+  void addTrait(TraitProperty Property, StringRef RawString,
+                APInt *Score = nullptr) {
+    addTrait(getOpenMPContextTraitSetForProperty(Property), Property, RawString,
+             Score);
+  }
+  /// Add the trait \p Property which is in set \p Set to the required trait
+  /// set. \p RawString is the string we parsed and derived \p Property from. If
+  /// \p Score is not null, it recorded as well. If \p Set is the `construct`
+  /// set it is recorded in-order in the ConstructTraits as well.
+  void addTrait(TraitSet Set, TraitProperty Property, StringRef RawString,
+                APInt *Score = nullptr) {
+    if (Score)
+      ScoreMap[Property] = *Score;
+
+    // Special handling for `device={isa(...)}` as we do not match the enum but
+    // the raw string.
+    if (Property == TraitProperty::device_isa___ANY)
+      ISATraits.push_back(RawString);
+
+    RequiredTraits.set(unsigned(Property));
+    if (Set == TraitSet::construct)
+      ConstructTraits.push_back(Property);
+  }
+
+  BitVector RequiredTraits = BitVector(unsigned(TraitProperty::Last) + 1);
+  SmallVector<StringRef, 8> ISATraits;
+  SmallVector<TraitProperty, 8> ConstructTraits;
+  SmallDenseMap<TraitProperty, APInt> ScoreMap;
+};
+
+/// The context for a source location is made up of active property traits,
+/// e.g., device={kind(host)}, and constructs traits which describe the nesting
+/// in OpenMP constructs at the location.
+struct OMPContext {
+  OMPContext(bool IsDeviceCompilation, Triple TargetTriple);
+  virtual ~OMPContext() = default;
+
+  void addTrait(TraitProperty Property) {
+    addTrait(getOpenMPContextTraitSetForProperty(Property), Property);
+  }
+  void addTrait(TraitSet Set, TraitProperty Property) {
+    ActiveTraits.set(unsigned(Property));
+    if (Set == TraitSet::construct)
+      ConstructTraits.push_back(Property);
+  }
+
+  /// Hook for users to check if an ISA trait matches. The trait is described as
+  /// the string that got parsed and it depends on the target and context if
+  /// this matches or not.
+  virtual bool matchesISATrait(StringRef) const { return false; }
+
+  BitVector ActiveTraits = BitVector(unsigned(TraitProperty::Last) + 1);
+  SmallVector<TraitProperty, 8> ConstructTraits;
+};
+
+/// Return true if \p VMI is applicable in \p Ctx, that is, all traits required
+/// by \p VMI are available in the OpenMP context \p Ctx. If \p DeviceSetOnly is
+/// true, only the device selector set, if present, are checked. Note that we
+/// still honor extension traits provided by the user.
+bool isVariantApplicableInContext(const VariantMatchInfo &VMI,
+                                  const OMPContext &Ctx,
+                                  bool DeviceSetOnly = false);
+
+/// Return the index (into \p VMIs) of the variant with the highest score
+/// from the ones applicble in \p Ctx. See llvm::isVariantApplicableInContext.
+int getBestVariantMatchForContext(const SmallVectorImpl<VariantMatchInfo> &VMIs,
+                                  const OMPContext &Ctx);
+
+} // namespace omp
+
+template <> struct DenseMapInfo<omp::TraitProperty> {
+  static inline omp::TraitProperty getEmptyKey() {
+    return omp::TraitProperty(-1);
+  }
+  static inline omp::TraitProperty getTombstoneKey() {
+    return omp::TraitProperty(-2);
+  }
+  static unsigned getHashValue(omp::TraitProperty val) {
+    return std::hash<unsigned>{}(unsigned(val));
+  }
+  static bool isEqual(omp::TraitProperty LHS, omp::TraitProperty RHS) {
+    return LHS == RHS;
+  }
+};
+
+} // end namespace llvm
+#endif // LLVM_FRONTEND_OPENMP_OMPCONTEXT_H

thirdparty/capstone/suite/synctools/tablegen/include/llvm/Frontend/OpenMP/OMPGridValues.h

@@ -0,0 +1,120 @@
+//====--- OMPGridValues.h - Language-specific address spaces --*- C++ -*-====//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// \brief Provides definitions for Target specific Grid Values
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_FRONTEND_OPENMP_OMPGRIDVALUES_H
+#define LLVM_FRONTEND_OPENMP_OMPGRIDVALUES_H
+
+namespace llvm {
+
+namespace omp {
+
+/// \brief Defines various target-specific GPU grid values that must be
+///        consistent between host RTL (plugin), device RTL, and clang.
+///        We can change grid values for a "fat" binary so that different
+///        passes get the correct values when generating code for a
+///        multi-target binary. Both amdgcn and nvptx values are stored in
+///        this file. In the future, should there be differences between GPUs
+///        of the same architecture, then simply make a different array and
+///        use the new array name.
+///
+/// Example usage in clang:
+///   const unsigned slot_size =
+///   ctx.GetTargetInfo().getGridValue().GV_Warp_Size;
+///
+/// Example usage in libomptarget/deviceRTLs:
+///   #include "llvm/Frontend/OpenMP/OMPGridValues.h"
+///   #ifdef __AMDGPU__
+///     #define GRIDVAL AMDGPUGridValues
+///   #else
+///     #define GRIDVAL NVPTXGridValues
+///   #endif
+///   ... Then use this reference for GV_Warp_Size in the deviceRTL source.
+///   llvm::omp::GRIDVAL().GV_Warp_Size
+///
+/// Example usage in libomptarget hsa plugin:
+///   #include "llvm/Frontend/OpenMP/OMPGridValues.h"
+///   #define GRIDVAL AMDGPUGridValues
+///   ... Then use this reference to access GV_Warp_Size in the hsa plugin.
+///   llvm::omp::GRIDVAL().GV_Warp_Size
+///
+/// Example usage in libomptarget cuda plugin:
+///    #include "llvm/Frontend/OpenMP/OMPGridValues.h"
+///    #define GRIDVAL NVPTXGridValues
+///   ... Then use this reference to access GV_Warp_Size in the cuda plugin.
+///    llvm::omp::GRIDVAL().GV_Warp_Size
+///
+
+struct GV {
+  /// The size reserved for data in a shared memory slot.
+  const unsigned GV_Slot_Size;
+  /// The default value of maximum number of threads in a worker warp.
+  const unsigned GV_Warp_Size;
+
+  constexpr unsigned warpSlotSize() const {
+    return GV_Warp_Size * GV_Slot_Size;
+  }
+
+  /// the maximum number of teams.
+  const unsigned GV_Max_Teams;
+  // An alternative to the heavy data sharing infrastructure that uses global
+  // memory is one that uses device __shared__ memory.  The amount of such space
+  // (in bytes) reserved by the OpenMP runtime is noted here.
+  const unsigned GV_SimpleBufferSize;
+  // The absolute maximum team size for a working group
+  const unsigned GV_Max_WG_Size;
+  // The default maximum team size for a working group
+  const unsigned GV_Default_WG_Size;
+
+  constexpr unsigned maxWarpNumber() const {
+    return GV_Max_WG_Size / GV_Warp_Size;
+  }
+};
+
+/// For AMDGPU GPUs
+static constexpr GV AMDGPUGridValues64 = {
+    256,  // GV_Slot_Size
+    64,   // GV_Warp_Size
+    128,  // GV_Max_Teams
+    896,  // GV_SimpleBufferSize
+    1024, // GV_Max_WG_Size,
+    256,  // GV_Default_WG_Size
+};
+
+static constexpr GV AMDGPUGridValues32 = {
+    256,  // GV_Slot_Size
+    32,   // GV_Warp_Size
+    128,  // GV_Max_Teams
+    896,  // GV_SimpleBufferSize
+    1024, // GV_Max_WG_Size,
+    256,  // GV_Default_WG_Size
+};
+
+template <unsigned wavesize> constexpr const GV &getAMDGPUGridValues() {
+  static_assert(wavesize == 32 || wavesize == 64, "");
+  return wavesize == 32 ? AMDGPUGridValues32 : AMDGPUGridValues64;
+}
+
+/// For Nvidia GPUs
+static constexpr GV NVPTXGridValues = {
+    256,  // GV_Slot_Size
+    32,   // GV_Warp_Size
+    1024, // GV_Max_Teams
+    896,  // GV_SimpleBufferSize
+    1024, // GV_Max_WG_Size
+    128,  // GV_Default_WG_Size
+};
+
+} // namespace omp
+} // namespace llvm
+
+#endif // LLVM_FRONTEND_OPENMP_OMPGRIDVALUES_H