diff --git a/clang/include/clang/Basic/BuiltinsLinxISA.td b/clang/include/clang/Basic/BuiltinsLinxISA.td index c897cb3ef0625..2f41d27542a0b 100644 --- a/clang/include/clang/Basic/BuiltinsLinxISA.td +++ b/clang/include/clang/Basic/BuiltinsLinxISA.td @@ -20,6 +20,9 @@ class LinxISABuiltin : 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)">; diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index f52470a4d7458..92eebf6da3c96 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/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()->getNumBits()))); diff --git a/clang/lib/Basic/Targets/LinxISA.cpp b/clang/lib/Basic/Targets/LinxISA.cpp index a7086810f1f36..8423b2cb4da57 100644 --- a/clang/lib/Basic/Targets/LinxISA.cpp +++ b/clang/lib/Basic/Targets/LinxISA.cpp @@ -74,6 +74,20 @@ 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(); @@ -81,6 +95,8 @@ bool LinxISATargetInfo::validateAsmConstraint( } 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(Feature) diff --git a/clang/lib/Basic/Targets/LinxISA.h b/clang/lib/Basic/Targets/LinxISA.h index ac24cbb8aef49..e181ec39edf19 100644 --- a/clang/lib/Basic/Targets/LinxISA.h +++ b/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 &Features, diff --git a/clang/lib/CodeGen/CMakeLists.txt b/clang/lib/CodeGen/CMakeLists.txt index ac10858817097..7ab3db89b7474 100644 --- a/clang/lib/CodeGen/CMakeLists.txt +++ b/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 diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index d50c9605a30b3..c86d7c265e41f 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/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: diff --git a/clang/lib/CodeGen/TargetBuiltins/LinxISA.cpp b/clang/lib/CodeGen/TargetBuiltins/LinxISA.cpp index 9324519293c48..45b0a3e177775 100644 --- a/clang/lib/CodeGen/TargetBuiltins/LinxISA.cpp +++ b/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))); diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h index db06584d766bf..46ba666f1b4d4 100644 --- a/clang/lib/CodeGen/TargetInfo.h +++ b/clang/lib/CodeGen/TargetInfo.h @@ -488,6 +488,9 @@ class TargetCodeGenInfo { std::unique_ptr createDefaultTargetCodeGenInfo(CodeGenModule &CGM); +std::unique_ptr +createLinxISATargetCodeGenInfo(CodeGenModule &CGM); + enum class AArch64ABIKind { AAPCS = 0, DarwinPCS, diff --git a/clang/lib/CodeGen/Targets/LinxISA.cpp b/clang/lib/CodeGen/Targets/LinxISA.cpp new file mode 100644 index 0000000000000..35a313673791f --- /dev/null +++ b/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(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(Ty); + if (!VT || VT->getPrimitiveSizeInBits() != 32768) + return nullptr; + + return llvm::FixedVectorType::get(CGF.Int32Ty, 1024); + } +}; + +} // namespace + +std::unique_ptr +CodeGen::createLinxISATargetCodeGenInfo(CodeGenModule &CGM) { + return std::make_unique(CGM.getTypes()); +} diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 5147d95506a71..0f66034cee885 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/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) && diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index c92b370b88d2d..fa3cbffe3507a 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/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 diff --git a/clang/lib/Headers/linx_blkc.h b/clang/lib/Headers/linx_blkc.h new file mode 100644 index 0000000000000..3181943435290 --- /dev/null +++ b/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 +static __inline__ Storage &__linx_storage_ref(Format &Value) { + static_assert(sizeof(Storage) == sizeof(Format)); + return *reinterpret_cast(&Value); +} + +template +static __inline__ const Storage &__linx_storage_ref(const Format &Value) { + static_assert(sizeof(Storage) == sizeof(Format)); + return *reinterpret_cast(&Value); +} + +#define __tf32_STORAGE(value) __linx_storage_ref(value) +#define __hf32_STORAGE(value) __linx_storage_ref(value) +#define __blkc_bf16_STORAGE(value) __linx_storage_ref(value) +#define __hif8_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e4m3_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e5m2_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e6m2_STORAGE(value) __linx_storage_ref(value) +#define __fp6_e3m2_STORAGE(value) __linx_storage_ref(value) +#define __fp6_e2m3_STORAGE(value) __linx_storage_ref(value) +#define __fp4_e2m1x2_STORAGE(value) __linx_storage_ref(value) +#define __fp4_e1m2x2_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e8m0_STORAGE(value) __linx_storage_ref(value) +#define __fp4_hif4x2_STORAGE(value) __linx_storage_ref(value) +#define __int4x2_STORAGE(value) __linx_storage_ref(value) +#define __uint4x2_STORAGE(value) __linx_storage_ref(value) +#define __fp16x2_STORAGE(value) __linx_storage_ref(value) +#define __bf16x2_STORAGE(value) __linx_storage_ref(value) +#define __blkc_bf16x2_STORAGE(value) __bf16x2_STORAGE(value) +#define __uint16x2_STORAGE(value) __linx_storage_ref(value) +#define __int16x2_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e4m3x4_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e5m2x4_STORAGE(value) __linx_storage_ref(value) +#define __uint8x4_STORAGE(value) __linx_storage_ref(value) +#define __int8x4_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e6m2x2_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e4m3x2_STORAGE(value) __linx_storage_ref(value) +#define __fp8_e5m2x2_STORAGE(value) __linx_storage_ref(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 diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 91d36f0502d85..8a355ac69bf29 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/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(); + const bool HasLinxStorageAnnotation = + ET && llvm::any_of(ET->getDecl()->specific_attrs(), + [](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(); diff --git a/clang/test/CodeGen/LinxISA/cxx-frontend.cpp b/clang/test/CodeGen/LinxISA/cxx-frontend.cpp new file mode 100644 index 0000000000000..d0e83584f11d5 --- /dev/null +++ b/clang/test/CodeGen/LinxISA/cxx-frontend.cpp @@ -0,0 +1,90 @@ +// RUN: %clang -target linx64-unknown-linux-musl -std=c++20 -S -emit-llvm %s -o - | FileCheck %s --check-prefix=IR +// RUN: %clang -target linx64-unknown-linux-musl -std=c++20 -c -O2 %s -o %t +// RUN: llvm-objdump -d %t | FileCheck %s --check-prefix=OBJ + +using E8M0Tile = __fp8_e8m0 tile_size(4096); +using FP16Tile = __half tile_size(2048); +using FP32Tile = __fp32 tile_size(1024); +using IntTile = int tile_size(1024); + +unsigned thread_index() { + return __builtin_linx_get_thread_idx(); +} + +// IR-LABEL: define{{.*}} i32 @_Z12thread_indexv() +// IR: call i32 asm "ssrget 0x0802, ->$0", "=r"() +// OBJ-LABEL: <_Z12thread_indexv>: +// OBJ: ssrget{{[[:space:]]+}}0x802,{{[[:space:]]+}}->a0 + +void tile_constraint(E8M0Tile &Dst, const E8M0Tile &Src) { + asm volatile("" : "=Tr"(Dst) : "Tr"(Src)); +} + +// IR-LABEL: define{{.*}} void @_Z15tile_constraint +// IR: bitcast <4096 x i8> {{.*}} to <1024 x i32> +// IR: call <1024 x i32> asm sideeffect "", "=^Tr,^Tr" +// IR: bitcast <1024 x i32> {{.*}} to <4096 x i8> + +void shared_chain() { + unsigned long Handle; + asm volatile("B.IOS mask=1111, ->%S0<4KB>" : "=Sr"(Handle)); + asm volatile("B.IOS %S0, mask=1111" : : "Sr"(Handle)); +} + +// IR-LABEL: define{{.*}} void @_Z12shared_chainv() +// IR: call i64 asm sideeffect "B.IOS mask=1111, ->${0:S}<4KB>", "=^Sr"() +// IR: call void asm sideeffect "B.IOS ${0:S}, mask=1111", "^Sr"(i64 {{.*}}) +// OBJ-LABEL: <_Z12shared_chainv>: +// OBJ-NOT: addi +// OBJ-NOT: LDI +// OBJ-NOT: SDI +// OBJ: B.IOS{{[[:space:]]+}}mask=1111, ->S0<4KB> +// OBJ: B.IOS{{[[:space:]]+}}S0, mask=1111 +// OBJ-NOT: addi +// OBJ-NOT: LDI +// OBJ-NOT: SDI + +#define TR_SIZE_CODE_TEST(Name, Type, SizeCode) \ + void Name() { \ + Type Value; \ + asm volatile("B.IOT mask=1111, last, ->%0<%Z1>" \ + : "=Tr"(Value) \ + : "i"(SizeCode)); \ + } + +TR_SIZE_CODE_TEST(fp32_128b, FP32Tile, 1) +TR_SIZE_CODE_TEST(fp16_4kb, FP16Tile, 6) +TR_SIZE_CODE_TEST(int_64kb, IntTile, 10) +TR_SIZE_CODE_TEST(mx_8kb, E8M0Tile, 7) + +// OBJ: B.IOT{{[[:space:]]+}}mask=1111, last, ->t<128B> +// OBJ: B.IOT{{[[:space:]]+}}mask=1111, last, ->t<4KB> +// OBJ: B.IOT{{[[:space:]]+}}mask=1111, last, ->t<64KB> +// OBJ: B.IOT{{[[:space:]]+}}mask=1111, last, ->t<8KB> + +#define SR_SIZE_CODE_TEST(Name, SizeCode) \ + void Name() { \ + unsigned long Handle; \ + asm volatile("B.IOS mask=1111, ->%S0<%Z1>" \ + : "=Sr"(Handle) \ + : "i"(SizeCode)); \ + } + +SR_SIZE_CODE_TEST(shared_128kb, 11) +SR_SIZE_CODE_TEST(shared_256kb, 12) + +// OBJ: B.IOS{{[[:space:]]+}}mask=1111, ->S0<128KB> +// OBJ: B.IOS{{[[:space:]]+}}mask=1111, ->S0<256KB> + +void tile_add(FP32Tile &Dst, const FP32Tile &A, const FP32Tile &B) { + asm volatile("BSTART.TEPL 0, 0, S32\n" + "B.IOT %1, %2, mask=1111, last, ->%0<4KB>" + : "=&Tr"(Dst) + : "Tr"(A), "Tr"(B)); +} + +// OBJ: BSTART.TLOAD{{[[:space:]]+}}S32 +// OBJ: BSTART.TLOAD{{[[:space:]]+}}S32 +// OBJ: BSTART.VEC{{[[:space:]]+}}TADD, S32 +// OBJ: B.IOT{{[[:space:]]+}}t#2, t#1, mask=1111, last, ->t<4KB> +// OBJ: BSTART.TSTORE{{[[:space:]]+}}S32 diff --git a/clang/test/Driver/linx-blkc-header.cpp b/clang/test/Driver/linx-blkc-header.cpp new file mode 100644 index 0000000000000..efb3c5725464d --- /dev/null +++ b/clang/test/Driver/linx-blkc-header.cpp @@ -0,0 +1,5 @@ +// RUN: %clang -target linx64-unknown-linux-musl -std=c++20 -### -fsyntax-only %s 2>&1 | FileCheck %s --check-prefix=AUTO +// RUN: %clang -target linx64-unknown-linux-musl -std=c++20 -### -fsyntax-only -nobuiltininc %s 2>&1 | FileCheck %s --check-prefix=NOAUTO + +// AUTO: "-include" "linx_blkc.h" +// NOAUTO-NOT: "-include" "linx_blkc.h" diff --git a/clang/test/Headers/linx-blkc-negative.cpp b/clang/test/Headers/linx-blkc-negative.cpp new file mode 100644 index 0000000000000..c991937c058ac --- /dev/null +++ b/clang/test/Headers/linx-blkc-negative.cpp @@ -0,0 +1,72 @@ +// RUN: not %clang -target x86_64-unknown-linux-gnu -std=c++20 -fsyntax-only \ +// RUN: -include %S/../../lib/Headers/linx_blkc.h %s 2>&1 | \ +// RUN: FileCheck %s --check-prefix=HOST +// RUN: not %clang -target linx64-unknown-linux-musl -std=c++20 -S -o /dev/null \ +// RUN: -DBAD_TR %s 2>&1 | FileCheck %s --check-prefix=TR +// RUN: not %clang -target linx64-unknown-linux-musl -std=c++20 -S -o /dev/null \ +// RUN: -DBAD_TR_SIZE %s 2>&1 | FileCheck %s --check-prefix=TR-SIZE +// RUN: not %clang -target linx64-unknown-linux-musl -std=c++20 -c -O2 -o /dev/null \ +// RUN: -DBAD_TR_TIED %s 2>&1 | FileCheck %s --check-prefix=TR-TIED +// RUN: not %clang -target linx64-unknown-linux-musl -std=c++20 -c -O2 -o /dev/null \ +// RUN: -DBAD_SIZE_CODE %s 2>&1 | FileCheck %s --check-prefix=SIZE-CODE +// RUN: not %clang -target linx64-unknown-linux-musl -std=c++20 -S -O2 -o /dev/null \ +// RUN: -DBAD_SR_ESCAPE %s 2>&1 | FileCheck %s --check-prefix=SR-ESCAPE +// RUN: not %clang -target linx64-unknown-linux-musl -std=c++20 -S -o /dev/null \ +// RUN: -DBAD_SR_TYPE %s 2>&1 | FileCheck %s --check-prefix=SR-TYPE +// RUN: not %clang -target linx64-unknown-linux-musl -std=c++20 -fsyntax-only \ +// RUN: -DBAD_VECTOR %s 2>&1 | FileCheck %s --check-prefix=VECTOR + +// HOST: error: "linx_blkc.h is only supported for the Linx target" + +#if defined(BAD_TR) +void bad_tr(unsigned &Value) { + asm volatile("" : "+Tr"(Value)); +} +// TR: error: invalid type 'unsigned int' in asm input for constraint '^Tr' +#endif + +#if defined(BAD_TR_SIZE) +using SmallTile = float tile_size(32); +void bad_tr_size(SmallTile &Value) { + asm volatile("" : "+Tr"(Value)); +} +// TR-SIZE: error: invalid type 'SmallTile' (vector of 32 'float' values) in asm input for constraint '^Tr' +#endif + +#if defined(BAD_TR_TIED) +using TiedTile = float tile_size(1024); +void bad_tr_tied(TiedTile &Value) { + asm volatile("B.IOT %0, mask=1111, last, ->%0<4KB>" : "+Tr"(Value)); +} +// TR-TIED: error: invalid operand in inline asm +#endif + +#if defined(BAD_SIZE_CODE) +using SizeTile = float tile_size(1024); +void bad_size_code(SizeTile &Value) { + asm volatile("B.IOT mask=1111, last, ->%0<%Z1>" + : "=Tr"(Value) + : "i"(0)); +} +// SIZE-CODE: error: invalid operand in inline asm +#endif + +#if defined(BAD_SR_ESCAPE) +void bad_sr_escape(unsigned long &Dst, unsigned long Src) { + asm volatile("" : "=Sr"(Dst) : "Sr"(Src)); +} +// SR-ESCAPE: error: {{.*}}Sr values are compiler-local Shared handles and cannot be copied to or from ordinary registers +#endif + +#if defined(BAD_SR_TYPE) +void bad_sr_type(float &Value) { + asm volatile("" : "=Sr"(Value)); +} +// SR-TYPE: error: invalid type 'float' in asm input for constraint '^Sr' +#endif + +#if defined(BAD_VECTOR) +enum class UserEnum : unsigned char {}; +using BadVector = UserEnum tile_size(16); +// VECTOR: error: invalid vector element type 'UserEnum' +#endif diff --git a/clang/test/Headers/linx-blkc.cpp b/clang/test/Headers/linx-blkc.cpp new file mode 100644 index 0000000000000..975a9ca2c9804 --- /dev/null +++ b/clang/test/Headers/linx-blkc.cpp @@ -0,0 +1,60 @@ +// RUN: %clang -target linx64-unknown-linux-musl -std=c++20 -fsyntax-only %s +// RUN: %clang -target linx32-unknown-linux-musl -std=c++20 -fsyntax-only %s + +#ifndef PTO_LINX_COMPAT_TYPES_PROVIDED +#error "linx_blkc.h was not included by the Linx C++ driver" +#endif + +static_assert(sizeof(__fp32) == 4); +static_assert(sizeof(__half) == 2); +static_assert(sizeof(__tf32) == 4); +static_assert(sizeof(__hf32) == 4); +static_assert(sizeof(__blkc_bf16) == 2); +static_assert(sizeof(__hif8) == 1); +static_assert(sizeof(__fp8_e4m3) == 1); +static_assert(sizeof(__fp8_e5m2) == 1); +static_assert(sizeof(__fp8_e6m2) == 1); +static_assert(sizeof(__fp6_e3m2) == 1); +static_assert(sizeof(__fp6_e2m3) == 1); +static_assert(sizeof(__fp4_e2m1x2) == 1); +static_assert(sizeof(__fp4_e1m2x2) == 1); +static_assert(sizeof(__fp8_e8m0) == 1); +static_assert(sizeof(__fp4_hif4x2) == 1); +static_assert(sizeof(__int4x2) == 1); +static_assert(sizeof(__uint4x2) == 1); + +static_assert(!__is_same(__fp8_e4m3, __fp8_e5m2)); +static_assert(!__is_same(__fp8_e8m0, __hif8)); +static_assert(!__is_same(__int4x2, __uint4x2)); + +static_assert(sizeof(__fp16x2) == 4); +static_assert(sizeof(__bf16x2) == 4); +static_assert(sizeof(__uint16x2) == 4); +static_assert(sizeof(__int16x2) == 4); +static_assert(sizeof(__fp8_e4m3x4) == 4); +static_assert(sizeof(__fp8_e5m2x4) == 4); +static_assert(sizeof(__uint8x4) == 4); +static_assert(sizeof(__int8x4) == 4); +static_assert(sizeof(__fp8_e6m2x2) == 2); +static_assert(sizeof(__fp8_e4m3x2) == 2); +static_assert(sizeof(__fp8_e5m2x2) == 2); + +using E8M0Tile = __fp8_e8m0 tile_size(4096); +static_assert(sizeof(E8M0Tile) == 4096); + +void tile_constraint(E8M0Tile &Dst, const E8M0Tile &Src) { + asm volatile("" : "=Tr"(Dst) : "Tr"(Src)); +} + +void vector_register_constraint(unsigned &Dst, unsigned Src) { + asm volatile("" : "=vr"(Dst) : "vr"(Src)); +} + +void shared_constraint(unsigned long &Dst, unsigned long Src) { + asm volatile("" : "=Sr"(Dst) : "Sr"(Src)); +} + +void storage_lvalues(__tf32 &TF32, __bf16x2 &BF16x2) { + __tf32_STORAGE(TF32) = 0; + __blkc_bf16x2_STORAGE(BF16x2) = 0; +} diff --git a/llvm/lib/Target/LinxISA/AsmParser/LinxISAAsmParser.cpp b/llvm/lib/Target/LinxISA/AsmParser/LinxISAAsmParser.cpp index fb259599564e7..e34b29fcea22c 100644 --- a/llvm/lib/Target/LinxISA/AsmParser/LinxISAAsmParser.cpp +++ b/llvm/lib/Target/LinxISA/AsmParser/LinxISAAsmParser.cpp @@ -1587,6 +1587,7 @@ bool LinxISAAsmParser::parseInstruction(ParseInstructionInfo &Info, bool IsTileIODesc = false; bool IsSharedIODesc = false; bool IsBlockAttr = false; + bool HasDataTypeField = false; unsigned MaxArrowDests = 0; { std::string Key = toUpperStr(Name); @@ -1643,6 +1644,7 @@ bool LinxISAAsmParser::parseInstruction(ParseInstructionInfo &Info, } for (unsigned FormIndex : It->second) { const linxisa_inst_form &F = linxisa_inst_forms[FormIndex]; + HasDataTypeField |= hasField(F, "DataType"); if (hasField(F, "RegDst1") || hasField(F, "RegDst0")) MaxArrowDests = std::max(MaxArrowDests, 2u); else if (hasField(F, "RegDst")) @@ -2034,6 +2036,15 @@ bool LinxISAAsmParser::parseInstruction(ParseInstructionInfo &Info, } } + if (HasDataTypeField && parseDataTypeKeyword(getTok().getString())) { + const MCExpr *Expr = nullptr; + SMLoc Start, End; + if (parseImmOperand(Expr, Start, End)) + return true; + Operands.push_back(LinxOperand::createImm(Expr, Start, End)); + continue; + } + // Registers (including suffixed forms like a0.sw). StringRef Tok = getTok().getString(); StringRef BaseTok = Tok; diff --git a/llvm/lib/Target/LinxISA/LinxISAAsmPrinter.cpp b/llvm/lib/Target/LinxISA/LinxISAAsmPrinter.cpp index 01b369cd08874..3401b705c622e 100644 --- a/llvm/lib/Target/LinxISA/LinxISAAsmPrinter.cpp +++ b/llvm/lib/Target/LinxISA/LinxISAAsmPrinter.cpp @@ -388,19 +388,29 @@ static StringRef linxTileDataTypeName(int64_t Value) { } } +static StringRef linxTileSizeName(int64_t Value) { + static constexpr StringLiteral Names[] = { + "", "128B", "256B", "512B", "1KB", "2KB", "4KB", + "8KB", "16KB", "32KB", "64KB", "128KB", "256KB"}; + if (Value < 1 || Value >= static_cast(std::size(Names))) + return {}; + return Names[Value]; +} + bool LinxISAAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, const char *ExtraCode, raw_ostream &OS) { // Clang/GCC use %cN for an immediate without target punctuation. Linx block // templates rely on that conventional modifier for selector and dimension - // fields. %DN prints a tile dtype keyword for attribute commands. %SN + // fields. %DN prints a tile dtype keyword for attribute commands. %ZN + // prints a SizeCode as its canonical byte size. %SN // prints a compiler-allocated Shared register as S0..S255. %qN // prints only a tile register's destination queue bank because // canonical B.IOT destinations are written as ->t/u/m/n, while source // operands name a concrete queue slot such as t#1. if (ExtraCode && ExtraCode[0] != 0 && !((ExtraCode[0] == 'c' || ExtraCode[0] == 'D' || ExtraCode[0] == 'S' || - ExtraCode[0] == 'q') && + ExtraCode[0] == 'Z' || ExtraCode[0] == 'q') && ExtraCode[1] == 0)) return true; @@ -423,6 +433,34 @@ bool LinxISAAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, return false; } if (LinxISA::TILERegClass.contains(MO.getReg())) { + if (MO.isDef()) { + // A plain pure output denotes the destination queue bank, matching + // the explicit %q modifier. A read/write operand is ambiguous because + // the same placeholder would need both bank and ranked-source syntax. + for (unsigned FlagNo = InlineAsm::MIOp_FirstOperand; + FlagNo < MI->getNumOperands();) { + const MachineOperand &FlagMO = MI->getOperand(FlagNo); + if (!FlagMO.isImm()) + break; + const InlineAsm::Flag Flag(FlagMO.getImm()); + const unsigned NumRegs = Flag.getNumOperandRegisters(); + if (FlagNo + 1u + NumRegs > MI->getNumOperands()) + return true; + unsigned MatchedOp = 0; + if (Flag.isRegUseKind() && Flag.isUseOperandTiedToDef(MatchedOp)) { + (void)MatchedOp; + for (unsigned I = 0; I != NumRegs; ++I) { + const MachineOperand &UseMO = MI->getOperand(FlagNo + 1u + I); + if (UseMO.isReg() && UseMO.getReg() == MO.getReg()) + return true; + } + } + FlagNo += 1u + NumRegs; + } + static constexpr char Banks[4] = {'t', 'u', 'm', 'n'}; + OS << Banks[(Enc >> 3) & 0x3]; + return false; + } auto It = InlineAsmTileRanks.find(MI); if (It == InlineAsmTileRanks.end() || OpNo >= It->second.size() || It->second[OpNo] == 0u) @@ -437,8 +475,9 @@ bool LinxISAAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, } if (MO.isImm()) { - if (ExtraCode && ExtraCode[0] == 'D') { - StringRef Name = linxTileDataTypeName(MO.getImm()); + if (ExtraCode && (ExtraCode[0] == 'D' || ExtraCode[0] == 'Z')) { + StringRef Name = ExtraCode[0] == 'D' ? linxTileDataTypeName(MO.getImm()) + : linxTileSizeName(MO.getImm()); if (Name.empty()) return true; OS << Name; @@ -449,8 +488,10 @@ bool LinxISAAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, } if (MO.isCImm()) { - if (ExtraCode && ExtraCode[0] == 'D') { - StringRef Name = linxTileDataTypeName(MO.getCImm()->getSExtValue()); + if (ExtraCode && (ExtraCode[0] == 'D' || ExtraCode[0] == 'Z')) { + StringRef Name = ExtraCode[0] == 'D' + ? linxTileDataTypeName(MO.getCImm()->getSExtValue()) + : linxTileSizeName(MO.getCImm()->getSExtValue()); if (Name.empty()) return true; OS << Name; diff --git a/llvm/lib/Target/LinxISA/LinxISAISelLowering.cpp b/llvm/lib/Target/LinxISA/LinxISAISelLowering.cpp index 4106d68fd87dc..13a6af2b0c263 100644 --- a/llvm/lib/Target/LinxISA/LinxISAISelLowering.cpp +++ b/llvm/lib/Target/LinxISA/LinxISAISelLowering.cpp @@ -1879,7 +1879,8 @@ const char *LinxISATargetLowering::getTargetNodeName(unsigned Opcode) const { LinxISATargetLowering::ConstraintType LinxISATargetLowering::getConstraintType(StringRef Constraint) const { - if (Constraint == "S") + if (Constraint == "S" || Constraint == "Sr" || Constraint == "Tr" || + Constraint == "vr") return C_RegisterClass; return TargetLowering::getConstraintType(Constraint); } @@ -1890,11 +1891,13 @@ LinxISATargetLowering::getRegForInlineAsmConstraint( if (Constraint.size() == 1) { switch (Constraint[0]) { case 'r': - if (VT.isVector() || VT == MVT::linxtile) - // PTO 0.58 uses the ordinary register constraint for typed tile - // carriers; the value type, rather than a retired multi-letter - // constraint, selects the architectural TILE register file. + if (VT == MVT::v1024i32 || VT == MVT::linxtile) + // The ordinary register constraint remains valid for typed tile + // carriers; the value type selects the architectural TILE register + // file. PTO C++ wrappers may spell this class explicitly as "Tr". return std::make_pair(0u, &LinxISA::TILERegClass); + if (VT.isVector()) + return std::make_pair(0u, nullptr); // The Linx GPR encoding space includes queue pseudo-registers // (t#k/u#k and the special RegDst encodings for ->t/->u). These are not // general-purpose registers and are not safe for C inline-asm operands @@ -1908,6 +1911,18 @@ LinxISATargetLowering::getRegForInlineAsmConstraint( } } + if (Constraint == "Tr") { + if (VT == MVT::v1024i32 || VT == MVT::linxtile) + return std::make_pair(0u, &LinxISA::TILERegClass); + return std::make_pair(0u, nullptr); + } + if (Constraint == "Sr") + return VT == MVT::i64 + ? std::make_pair(0u, &LinxISA::SHAREDRegClass) + : std::make_pair(0u, nullptr); + if (Constraint == "vr") + return std::make_pair(0u, &LinxISA::GPR_ArchRegClass); + // Explicit register constraints: "{a0}", "{a1}", ... if (Constraint.size() > 2 && Constraint.front() == '{' && Constraint.back() == '}') { diff --git a/llvm/lib/Target/LinxISA/LinxISAInstrInfo.cpp b/llvm/lib/Target/LinxISA/LinxISAInstrInfo.cpp index 632a2a56453e0..aa39c4b1d0d28 100644 --- a/llvm/lib/Target/LinxISA/LinxISAInstrInfo.cpp +++ b/llvm/lib/Target/LinxISA/LinxISAInstrInfo.cpp @@ -10,6 +10,7 @@ #include "LinxISASubtarget.h" #include "MCTargetDesc/LinxISAMCTargetDesc.h" #include "llvm/CodeGen/MachineInstrBuilder.h" +#include "llvm/IR/Function.h" #include "llvm/Support/ErrorHandling.h" using namespace llvm; @@ -28,6 +29,17 @@ void LinxISAInstrInfo::copyPhysReg(MachineBasicBlock &MBB, Register SrcReg, bool KillSrc, bool RenamableDest, bool RenamableSrc) const { + const bool DestShared = LinxISA::SHAREDRegClass.contains(DestReg); + const bool SrcShared = LinxISA::SHAREDRegClass.contains(SrcReg); + if (DestShared || SrcShared) { + Function &F = MBB.getParent()->getFunction(); + F.getContext().emitError( + F.getName() + + ": Sr values are compiler-local Shared handles and cannot be copied " + "to or from ordinary registers"); + return; + } + if (!LinxISA::GPRRegClass.contains(DestReg, SrcReg)) report_fatal_error("Linx: unsupported reg-to-reg copy"); diff --git a/llvm/test/MC/LinxISA/v0583-s8-dtype-context.s b/llvm/test/MC/LinxISA/v0583-s8-dtype-context.s new file mode 100644 index 0000000000000..5b983a92a72dc --- /dev/null +++ b/llvm/test/MC/LinxISA/v0583-s8-dtype-context.s @@ -0,0 +1,17 @@ +# RUN: llvm-mc -triple=linx64 -show-encoding %s | FileCheck %s --check-prefix=ENC +# RUN: llvm-mc -triple=linx64 -filetype=obj %s -o %t +# RUN: llvm-objdump -d %t | FileCheck %s --check-prefix=DIS +# RUN: llvm-mc -triple=linx64 -show-encoding %s | sed -E 's/[[:space:]]*#.*$//' | llvm-mc -triple=linx64 -filetype=obj -o %t.re +# RUN: cmp %t %t.re + +BSTART.TLOAD S8 +BSTART.TMATMUL S8 +add s8, zero, ->a0 + +# ENC: BSTART.TLOAD S8{{.*}}encoding: +# ENC: BSTART.TMATMUL S8{{.*}}encoding: +# ENC: add s8, zero, ->a0{{.*}}encoding: + +# DIS: BSTART.TLOAD{{[[:space:]]+}}S8 +# DIS: BSTART.TMATMUL{{[[:space:]]+}}S8 +# DIS: add{{[[:space:]]+}}s8, zero, ->a0