LinxISA: package the PTO C++ frontend ABI

clang/include/clang/Basic/BuiltinsLinxISA.td

@@ -20,6 +20,9 @@ class LinxISABuiltin<string prototype, string features = ""> : TargetBuiltin {
 
 // PTO/Linx 0.58 tile bridge: 4 KiB frontend values as _Vector<1024, int>.
 let Attributes = [NoThrow] in {
+  // Read the current processing-element ID from the read-only PEID SSR.
+  def get_thread_idx : LinxISABuiltin<"unsigned int()">;
+
   // TLSU transfers and Local-to-Local movement.
   def tile_tload : LinxISABuiltin<"_Vector<1024, int>(void const *, _Constant unsigned int, _Constant unsigned int, _Constant long long int, long long int, long long int, long long int, long long int)">;
   def tile_tstore : LinxISABuiltin<"void(void *, _Vector<1024, int>, _Constant unsigned int, _Constant unsigned int, _Constant long long int, long long int, long long int, long long int, long long int)">;

clang/lib/AST/ASTContext.cpp

@@ -4632,10 +4632,12 @@ QualType ASTContext::getDependentVectorType(QualType VecType, Expr *SizeExpr,
 }
 
 /// getExtVectorType - Return the unique reference to an extended vector type of
-/// the specified element type and size. VectorType must be a built-in type.
+/// the specified element type and size. VectorType must be a built-in type,
+/// or a target-approved fixed-underlying enum used as an opaque storage type.
 QualType ASTContext::getExtVectorType(QualType vecType,
                                       unsigned NumElts) const {
-  assert(vecType->isBuiltinType() || vecType->isDependentType() ||
+  assert(vecType->isBuiltinType() || vecType->isEnumeralType() ||
+         vecType->isDependentType() ||
          (vecType->isBitIntType() &&
           // Only support _BitInt elements with byte-sized power of 2 NumBits.
           llvm::isPowerOf2_32(vecType->castAs<BitIntType>()->getNumBits())));

clang/lib/Basic/Targets/LinxISA.cpp

@@ -74,13 +74,29 @@ bool LinxISATargetInfo::validateAsmConstraint(
   // Z - first special register (ra/r10)
 
   switch (Name[0]) {
+  case 'T': {
+    if (Name[1] != 'r')
+      return false;
+    ++Name;
+    Info.setAllowsRegister();
+    return true;
+  }
+  case 'v': {
+    if (Name[1] != 'r')
+      return false;
+    ++Name;
+    Info.setAllowsRegister();
+    return true;
+  }
   case 'r': {
     // General purpose register
     Info.setAllowsRegister();
     return true;
   }
   case 'S': {
     // Compiler-allocated core-private Shared tile register S0..S255.
+    if (Name[1] == 'r')
+      ++Name;
     Info.setAllowsRegister();
     return true;
   }
@@ -145,6 +161,17 @@ bool LinxISATargetInfo::validateAsmConstraint(
   }
 }
 
+std::string
+LinxISATargetInfo::convertConstraint(const char *&Constraint) const {
+  if ((Constraint[0] == 'T' || Constraint[0] == 'S' || Constraint[0] == 'v') &&
+      Constraint[1] == 'r') {
+    std::string Result = std::string("^") + std::string(Constraint, 2);
+    ++Constraint;
+    return Result;
+  }
+  return TargetInfo::convertConstraint(Constraint);
+}
+
 bool LinxISATargetInfo::hasFeature(StringRef Feature) const {
   const bool Is64Bit = getTriple().isArch64Bit();
   return llvm::StringSwitch<bool>(Feature)

clang/lib/Basic/Targets/LinxISA.h

@@ -74,6 +74,7 @@ class LLVM_LIBRARY_VISIBILITY LinxISATargetInfo : public TargetInfo {
 
   bool validateAsmConstraint(const char *&Name,
                              TargetInfo::ConstraintInfo &Info) const override;
+  std::string convertConstraint(const char *&Constraint) const override;
 
   bool hasFeature(StringRef Feature) const override;
   bool handleTargetFeatures(std::vector<std::string> &Features,

clang/lib/CodeGen/CMakeLists.txt

@@ -141,6 +141,7 @@ add_clang_library(clangCodeGen
   Targets/DirectX.cpp
   Targets/Hexagon.cpp
   Targets/Lanai.cpp
+  Targets/LinxISA.cpp
   Targets/LoongArch.cpp
   Targets/M68k.cpp
   Targets/MSP430.cpp

clang/lib/CodeGen/CodeGenModule.cpp

@@ -117,6 +117,10 @@ createTargetCodeGenInfo(CodeGenModule &CGM) {
   default:
     return createDefaultTargetCodeGenInfo(CGM);
 
+  case llvm::Triple::linx32:
+  case llvm::Triple::linx64:
+    return createLinxISATargetCodeGenInfo(CGM);
+
   case llvm::Triple::m68k:
     return createM68kTargetCodeGenInfo(CGM);
   case llvm::Triple::mips:

clang/lib/CodeGen/TargetBuiltins/LinxISA.cpp

@@ -8,6 +8,7 @@
 
 #include "CGBuiltin.h"
 #include "clang/Basic/TargetBuiltins.h"
+#include "llvm/IR/InlineAsm.h"
 #include "llvm/IR/Intrinsics.h"
 #include "llvm/IR/IntrinsicsLinx.h"
 
@@ -25,6 +26,13 @@ llvm::Value *CodeGenFunction::EmitLinxISABuiltinExpr(unsigned BuiltinID,
   };
 
   switch (BuiltinID) {
+  case LinxISA::BI__builtin_linx_get_thread_idx: {
+    llvm::FunctionType *FTy =
+        llvm::FunctionType::get(Builder.getInt32Ty(), /*Variadic=*/false);
+    llvm::InlineAsm *ReadPEID = llvm::InlineAsm::get(
+        FTy, "ssrget 0x0802, ->$0", "=r", /*hasSideEffects=*/false);
+    return Builder.CreateCall(ReadPEID, {}, "linx.thread.idx");
+  }
   case LinxISA::BI__builtin_linx_tile_tload: {
     llvm::Value *Base = EmitScalarExpr(E->getArg(0));
     llvm::Value *Size = castToI32(EmitScalarExpr(E->getArg(1)));

clang/lib/CodeGen/TargetInfo.h

@@ -488,6 +488,9 @@ class TargetCodeGenInfo {
 std::unique_ptr<TargetCodeGenInfo>
 createDefaultTargetCodeGenInfo(CodeGenModule &CGM);
 
+std::unique_ptr<TargetCodeGenInfo>
+createLinxISATargetCodeGenInfo(CodeGenModule &CGM);
+
 enum class AArch64ABIKind {
   AAPCS = 0,
   DarwinPCS,

clang/lib/CodeGen/Targets/LinxISA.cpp

@@ -0,0 +1,44 @@
+//===-- LinxISA.cpp - LinxISA-specific CodeGen hooks ---------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "ABIInfoImpl.h"
+#include "TargetInfo.h"
+#include "llvm/IR/DerivedTypes.h"
+
+using namespace clang;
+using namespace clang::CodeGen;
+
+namespace {
+
+class LinxISATargetCodeGenInfo final : public TargetCodeGenInfo {
+public:
+  explicit LinxISATargetCodeGenInfo(CodeGenTypes &CGT)
+      : TargetCodeGenInfo(std::make_unique<DefaultABIInfo>(CGT)) {}
+
+  llvm::Type *adjustInlineAsmType(CodeGenFunction &CGF, StringRef Constraint,
+                                  llvm::Type *Ty) const override {
+    Constraint.consume_front("&");
+    if (Constraint == "Sr" || Constraint == "^Sr")
+      return Ty->isIntegerTy(64) ? Ty : nullptr;
+    if (Constraint != "Tr" && Constraint != "^Tr")
+      return Ty;
+
+    auto *VT = dyn_cast<llvm::FixedVectorType>(Ty);
+    if (!VT || VT->getPrimitiveSizeInBits() != 32768)
+      return nullptr;
+
+    return llvm::FixedVectorType::get(CGF.Int32Ty, 1024);
+  }
+};
+
+} // namespace
+
+std::unique_ptr<TargetCodeGenInfo>
+CodeGen::createLinxISATargetCodeGenInfo(CodeGenModule &CGM) {
+  return std::make_unique<LinxISATargetCodeGenInfo>(CGM.getTypes());
+}

clang/lib/Driver/ToolChains/Clang.cpp

@@ -939,6 +939,16 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
   if (JA.isOffloading(Action::OFK_SYCL))
     getToolChain().addSYCLIncludeArgs(Args, CmdArgs);
 
+  const llvm::Triple::ArchType TargetArch =
+      getToolChain().getTriple().getArch();
+  if ((TargetArch == llvm::Triple::linx32 ||
+       TargetArch == llvm::Triple::linx64) &&
+      !Inputs.empty() && types::isCXX(Inputs[0].getType()) &&
+      !Args.hasArg(options::OPT_nostdinc, options::OPT_nobuiltininc)) {
+    CmdArgs.push_back("-include");
+    CmdArgs.push_back("linx_blkc.h");
+  }
+
   // If we are offloading to a target via OpenMP we need to include the
   // openmp_wrappers folder which contains alternative system headers.
   if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&

clang/lib/Headers/CMakeLists.txt

@@ -10,6 +10,7 @@ set(core_files
   inttypes.h
   iso646.h
   limits.h
+  linx_blkc.h
   module.modulemap
   stdalign.h
   stdarg.h

clang/lib/Headers/linx_blkc.h

@@ -0,0 +1,105 @@
+//===-- linx_blkc.h - Linx block C++ frontend ABI ------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __LINX_BLKC_H
+#define __LINX_BLKC_H
+
+#if !defined(__linx) && !defined(__LINX__)
+#error "linx_blkc.h is only supported for the Linx target"
+#endif
+
+#define PTO_LINX_COMPAT_TYPES_PROVIDED 1
+
+using __fp32 = float;
+using __half = _Float16;
+
+#define __LINX_STORAGE_TYPE(NAME, STORAGE)                                 \
+  enum class __attribute__((annotate("linx.storage_type"))) NAME : STORAGE {}
+
+__LINX_STORAGE_TYPE(__tf32, unsigned int);
+__LINX_STORAGE_TYPE(__hf32, unsigned int);
+__LINX_STORAGE_TYPE(__blkc_bf16, unsigned short);
+__LINX_STORAGE_TYPE(__hif8, unsigned char);
+__LINX_STORAGE_TYPE(__fp8_e4m3, unsigned char);
+__LINX_STORAGE_TYPE(__fp8_e5m2, unsigned char);
+__LINX_STORAGE_TYPE(__fp8_e6m2, unsigned char);
+__LINX_STORAGE_TYPE(__fp6_e3m2, unsigned char);
+__LINX_STORAGE_TYPE(__fp6_e2m3, unsigned char);
+__LINX_STORAGE_TYPE(__fp4_e2m1x2, unsigned char);
+__LINX_STORAGE_TYPE(__fp4_e1m2x2, unsigned char);
+__LINX_STORAGE_TYPE(__fp8_e8m0, unsigned char);
+__LINX_STORAGE_TYPE(__fp4_hif4x2, unsigned char);
+__LINX_STORAGE_TYPE(__int4x2, signed char);
+__LINX_STORAGE_TYPE(__uint4x2, unsigned char);
+
+#undef __LINX_STORAGE_TYPE
+
+#define __LINX_PACKED_STORAGE_TYPE(NAME, STORAGE)                            \
+  struct NAME {                                                              \
+    STORAGE data;                                                            \
+  }
+
+__LINX_PACKED_STORAGE_TYPE(__fp16x2, unsigned int);
+__LINX_PACKED_STORAGE_TYPE(__bf16x2, unsigned int);
+__LINX_PACKED_STORAGE_TYPE(__uint16x2, unsigned int);
+__LINX_PACKED_STORAGE_TYPE(__int16x2, unsigned int);
+__LINX_PACKED_STORAGE_TYPE(__fp8_e4m3x4, unsigned int);
+__LINX_PACKED_STORAGE_TYPE(__fp8_e5m2x4, unsigned int);
+__LINX_PACKED_STORAGE_TYPE(__uint8x4, unsigned int);
+__LINX_PACKED_STORAGE_TYPE(__int8x4, unsigned int);
+__LINX_PACKED_STORAGE_TYPE(__fp8_e6m2x2, unsigned short);
+__LINX_PACKED_STORAGE_TYPE(__fp8_e4m3x2, unsigned short);
+__LINX_PACKED_STORAGE_TYPE(__fp8_e5m2x2, unsigned short);
+
+#undef __LINX_PACKED_STORAGE_TYPE
+
+template <typename Storage, typename Format>
+static __inline__ Storage &__linx_storage_ref(Format &Value) {
+  static_assert(sizeof(Storage) == sizeof(Format));
+  return *reinterpret_cast<Storage *>(&Value);
+}
+
+template <typename Storage, typename Format>
+static __inline__ const Storage &__linx_storage_ref(const Format &Value) {
+  static_assert(sizeof(Storage) == sizeof(Format));
+  return *reinterpret_cast<const Storage *>(&Value);
+}
+
+#define __tf32_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __hf32_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __blkc_bf16_STORAGE(value) __linx_storage_ref<unsigned short>(value)
+#define __hif8_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp8_e4m3_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp8_e5m2_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp8_e6m2_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp6_e3m2_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp6_e2m3_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp4_e2m1x2_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp4_e1m2x2_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp8_e8m0_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp4_hif4x2_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __int4x2_STORAGE(value) __linx_storage_ref<signed char>(value)
+#define __uint4x2_STORAGE(value) __linx_storage_ref<unsigned char>(value)
+#define __fp16x2_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __bf16x2_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __blkc_bf16x2_STORAGE(value) __bf16x2_STORAGE(value)
+#define __uint16x2_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __int16x2_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __fp8_e4m3x4_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __fp8_e5m2x4_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __uint8x4_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __int8x4_STORAGE(value) __linx_storage_ref<unsigned int>(value)
+#define __fp8_e6m2x2_STORAGE(value) __linx_storage_ref<unsigned short>(value)
+#define __fp8_e4m3x2_STORAGE(value) __linx_storage_ref<unsigned short>(value)
+#define __fp8_e5m2x2_STORAGE(value) __linx_storage_ref<unsigned short>(value)
+
+// TileOP uses this type modifier after the element type. The element count is
+// physical storage count, so packed public formats still use byte carriers.
+#define tile_size(elements) __attribute__((ext_vector_type(elements)))
+
+#endif // __LINX_BLKC_H

clang/lib/Sema/SemaType.cpp

@@ -2413,8 +2413,19 @@ QualType Sema::BuildExtVectorType(QualType T, Expr *SizeExpr,
   //
   // We explicitly allow bool elements in ext_vector_type for C/C++.
   bool IsNoBoolVecLang = getLangOpts().OpenCL || getLangOpts().OpenCLCPlusPlus;
+  const llvm::Triple::ArchType Arch =
+      Context.getTargetInfo().getTriple().getArch();
+  const auto *ET = T->getAs<EnumType>();
+  const bool HasLinxStorageAnnotation =
+      ET && llvm::any_of(ET->getDecl()->specific_attrs<AnnotateAttr>(),
+                         [](const AnnotateAttr *A) {
+                           return A->getAnnotation() == "linx.storage_type";
+                         });
+  const bool IsLinxStorageEnum =
+      HasLinxStorageAnnotation &&
+      (Arch == llvm::Triple::linx32 || Arch == llvm::Triple::linx64);
   if ((!T->isDependentType() && !T->isIntegerType() &&
-       !T->isRealFloatingType()) ||
+       !T->isRealFloatingType() && !IsLinxStorageEnum) ||
       (IsNoBoolVecLang && T->isBooleanType())) {
     Diag(AttrLoc, diag::err_attribute_invalid_vector_type) << T;
     return QualType();