diff --git a/include/PTO/IR/CMakeLists.txt b/include/PTO/IR/CMakeLists.txt index c0dcef2881..00098d0b63 100644 --- a/include/PTO/IR/CMakeLists.txt +++ b/include/PTO/IR/CMakeLists.txt @@ -31,6 +31,12 @@ mlir_tablegen(PTOAttrs.cpp.inc -gen-attrdef-defs -attrdefs-dialect=pto) mlir_tablegen(PTOEnums.h.inc -gen-enum-decls) mlir_tablegen(PTOEnums.cpp.inc -gen-enum-defs) +# 生成 Attribute Interfaces 定义 +set(LLVM_TARGET_DEFINITIONS PTOAttrInterfaces.td) +mlir_tablegen(PTOAttrInterfaces.h.inc -gen-attr-interface-decls) +mlir_tablegen(PTOAttrInterfaces.cpp.inc -gen-attr-interface-defs) +set(LLVM_TARGET_DEFINITIONS PTOOps.td) + # ============================================================ # 2. 处理 Interfaces (防止下个报错) # ============================================================ diff --git a/include/PTO/IR/PTO.h b/include/PTO/IR/PTO.h index 695072cd80..5ff7e50fc4 100644 --- a/include/PTO/IR/PTO.h +++ b/include/PTO/IR/PTO.h @@ -57,6 +57,8 @@ // PTO Attributes //===----------------------------------------------------------------------===// +#include "PTO/IR/PTOAttrInterfaces.h.inc" + #define GET_ATTRDEF_CLASSES #include "PTO/IR/PTOAttrs.h.inc" diff --git a/include/PTO/IR/PTOAttrInterfaces.td b/include/PTO/IR/PTOAttrInterfaces.td new file mode 100644 index 0000000000..80b290fc21 --- /dev/null +++ b/include/PTO/IR/PTOAttrInterfaces.td @@ -0,0 +1,51 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- PTOAttrInterfaces.td - PTO attribute interfaces ------*- tablegen -*-===// +//===----------------------------------------------------------------------===// +// +// This is the definition file for PTO dialect attribute interfaces. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_PTO_IR_PTOATTRINTERFACES +#define MLIR_DIALECT_PTO_IR_PTOATTRINTERFACES + +include "mlir/IR/AttrTypeBase.td" +include "mlir/IR/OpBase.td" + +//===----------------------------------------------------------------------===// +// Enum token reflection +//===----------------------------------------------------------------------===// + +// Implemented by PTO enum attributes that can feed a VPTO bridge template +// token. This MLIR build has no generic C++ runtime base for TableGen enum +// attributes, so the declarative bridge lowering reflects the enum case +// through this interface instead of knowing each attribute class: it +// assembles `EnumType::Case` spellings from the whitelist `enum_type` +// prefix and the enumerator symbol exposed here. +def PTO_EnumTokenAttr : AttrInterface<"EnumTokenAttr"> { + let cppNamespace = "::mlir::pto"; + let description = [{ + An attribute whose value is one case of a PTO enum and can be rendered + as a C++ template token by the declarative VPTO bridge lowering. + }]; + let methods = [ + InterfaceMethod< + /*desc=*/[{ + Returns the C++ enumerator symbol of the attribute's enum case + (e.g. "Partial" for `pto::AccPhase::Partial`). + }], + /*retTy=*/"llvm::StringRef", + /*methodName=*/"getEnumCaseSymbol", + /*args=*/(ins) + >, + ]; +} + +#endif // MLIR_DIALECT_PTO_IR_PTOATTRINTERFACES diff --git a/include/PTO/IR/PTOAttrs.td b/include/PTO/IR/PTOAttrs.td index 283d9ab105..608fab1816 100644 --- a/include/PTO/IR/PTOAttrs.td +++ b/include/PTO/IR/PTOAttrs.td @@ -17,6 +17,7 @@ #define MLIR_DIALECT_PTO_IR_PTOATTRS include "PTO/IR/PTODialect.td" +include "PTO/IR/PTOAttrInterfaces.td" include "mlir/Dialect/LLVMIR/LLVMOpBase.td" include "mlir/IR/AttrTypeBase.td" @@ -867,7 +868,8 @@ def PTO_AccPhaseEnum : PTO_I32Enum< I32EnumAttrCase<"Final", 3, "final"> ]>; -def PTO_AccPhaseAttr : EnumAttr { +def PTO_AccPhaseAttr : EnumAttr]> { let summary = "TMATMUL / TMATMUL_ACC accumulation phase attribute"; } diff --git a/include/PTO/IR/PTOOps.td b/include/PTO/IR/PTOOps.td index 48b8498ced..9ade14ad73 100644 --- a/include/PTO/IR/PTOOps.td +++ b/include/PTO/IR/PTOOps.td @@ -2677,6 +2677,51 @@ def TFreeOp : PTO_TOp<"tfree", [ }]; } +//===----------------------------------------------------------------------===// +// C++ Interface Bridge Ops +//===----------------------------------------------------------------------===// +// +// These internal ops form the generic VPTO C++ interface bridge. A +// family-specific pass +// (e.g. the TPipe family pass) rewrites its internal ops into bridge ops +// carrying only the wrapper callee name and ABI values. The generic bridge +// lowering pass then mechanically converts them into calls to wrapper +// functions whose implementations come from externally compiled device +// bitcode linked into the VPTO device module. + +def BridgeCallOp : PTO_Op<"bridge_call", [ + DeclareOpInterfaceMethods +]> { + let summary = "Generic call into an externally linked C++ wrapper entry"; + + let arguments = (ins + StrAttr:$callee, + OptionalAttr:$storage_size_callee, + Variadic:$args + ); + + let results = (outs Variadic:$results); + + let assemblyFormat = [{ + $callee + (`{` `storage_size_callee` `=` $storage_size_callee^ `}`)? + (`(` $args^ `:` type($args) `)`)? + attr-dict + (`->` type($results)^)? + }]; +} + +def BridgeIntToPtrOp : PTO_Op<"bridge_inttoptr"> { + let summary = "Generic bridge address materialization (i64 to PTO pointer)"; + + let arguments = (ins SignlessIntegerLike:$addr); + let results = (outs AnyType:$result); + + let assemblyFormat = [{ + $addr attr-dict `:` type($addr) `->` type($result) + }]; +} + //===----------------------------------------------------------------------===// // Synchronization Ops //===----------------------------------------------------------------------===// diff --git a/include/PTO/Transforms/PTOCppTokens.h b/include/PTO/Transforms/PTOCppTokens.h new file mode 100644 index 0000000000..549c0b1263 --- /dev/null +++ b/include/PTO/Transforms/PTOCppTokens.h @@ -0,0 +1,82 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// http://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- PTOCppTokens.h - shared PTO-ISA C++ token mappings -------*- C++ -*-===// +//===----------------------------------------------------------------------===// +// +// Single source of truth for the pure mappings from IR facts (element +// types, enums, integer attribute values) to PTO-ISA C++ spellings. Both +// the EmitC backend and the VPTO C++ interface bridge render the same +// pto-isa template tokens; they share these mapping functions and keep +// their own assembly logic (which template arguments to emit, how the +// tokens are consumed). +// +// Every builder takes a `qualifier` prefix that is prepended to the +// pto-isa constant/type spelling: the bridge passes "pto::" (the wrapper +// is a standalone translation unit and always spells fully qualified +// names), the EmitC backend passes an empty string (its output relies on +// the surrounding namespace context). +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_PTO_TRANSFORMS_PTOCPPTOKENS_H +#define MLIR_DIALECT_PTO_TRANSFORMS_PTOCPPTOKENS_H + +#include "PTO/IR/PTO.h" +#include "mlir/IR/Types.h" +#include "mlir/Support/LLVM.h" +#include "mlir/Support/LogicalResult.h" +#include "llvm/ADT/StringRef.h" +#include + +namespace mlir { +namespace pto { + +/// Builds the C++ element type token (e.g. "float", "half", "int8_t") for +/// an MLIR element type. Falls back to "float" for unrecognized types. +std::string getPTOCppElementTypeToken(Type elementType); + +/// Builds the `TileSplitAxis::TILE_*` token for a split value (0..4). +/// Fails for values outside that range. +FailureOr getPTOCppTileSplitToken(int64_t split, + llvm::StringRef qualifier); + +/// Builds the `Direction::DIR_*` token for a local pipe dir_mask (1=C2V, +/// 2=V2C, 3=BOTH). The L2G2L "_GM" variants are an EmitC-side extension +/// and are not part of this core mapping. Fails for other masks. +FailureOr getPTOCppDirectionToken(int8_t dirMask, + llvm::StringRef qualifier); + +/// Builds the `TileType::*` token for a local tile address space. Fails +/// for address spaces with no TileType mapping (e.g. global memory); +/// callers apply their own fallback policy for those. +FailureOr getPTOCppTileTypeToken(AddressSpace addressSpace, + llvm::StringRef qualifier); + +/// Builds the `BLayout::*` token. Fails for values outside the closed set. +FailureOr getPTOCppBLayoutToken(BLayout bLayout, + llvm::StringRef qualifier); + +/// Builds the `SLayout::*` token. Fails for values outside the closed set. +FailureOr getPTOCppSLayoutToken(SLayout sLayout, + llvm::StringRef qualifier); + +/// Renders the `TPipe` spelling; `dirTok` is an already rendered direction token. +std::string renderTPipeSpelling(int32_t flagBase, llvm::StringRef dirTok, + int32_t slotSize, int32_t slotNum, + int32_t localSlotNum, bool nosplit, + llvm::StringRef qualifier); + +} // namespace pto +} // namespace mlir + +#endif // MLIR_DIALECT_PTO_TRANSFORMS_PTOCPPTOKENS_H diff --git a/include/PTO/Transforms/Passes.h b/include/PTO/Transforms/Passes.h index 1a8dc14307..f0e838492c 100644 --- a/include/PTO/Transforms/Passes.h +++ b/include/PTO/Transforms/Passes.h @@ -45,6 +45,10 @@ std::unique_ptr createPTOValidatePhysicalSectionBoundariesPass(); std::unique_ptr createPTOMaterializeTileOpSectionsPass(); std::unique_ptr createVPTOSplitCVModulePass(); std::unique_ptr createVPTONormalizeContainerPass(); +std::unique_ptr createPTOLowerPipeFamilyOpsPass(); +std::unique_ptr createPTOLowerDeclarativeBridgeOpsPass(); +std::unique_ptr createVPTOBridgeLoweringPass(); +std::unique_ptr createVPTOBridgeWrapperGenPass(); std::unique_ptr createPTOVerifyTFreePass(); // Creates a pass for ... diff --git a/include/PTO/Transforms/Passes.td b/include/PTO/Transforms/Passes.td index 13a084cbc4..489d7e16de 100644 --- a/include/PTO/Transforms/Passes.td +++ b/include/PTO/Transforms/Passes.td @@ -426,6 +426,109 @@ def PTOMaterializeTileOpSections ]; } +def PTOLowerPipeFamilyOps : Pass<"pto-lower-pipe-family-ops", "mlir::func::FuncOp"> { + let summary = "Lower internal TPipe ops into generic VPTO bridge ops"; + let description = [{ + TPipe family pass of the VPTO C++ interface bridge. It rewrites the + internal pipe ops (pto.initialize_l2l_pipe / pto.tpush / pto.tpop / + pto.tfree) and the tile handles they consume (pto.alloc_tile / + pto.declare_tile / pto.tile_buf_addr) into generic pto.bridge_call / + pto.bridge_inttoptr ops that carry only wrapper callee names and ABI + values. All pipe family semantics (config validation, storage handle + flow, and the runtime rebinding of a declared tile to the FIFO slot + returned by TPOP) are resolved here so that the generic bridge lowering + stays family-agnostic. + }]; + let constructor = "mlir::pto::createPTOLowerPipeFamilyOpsPass()"; + let options = [ + Option<"whitelistPath", "whitelist-path", "std::string", + /*default=*/"\"\"", + "Path to the VPTO bridge whitelist YAML; falls back to the " + "PTOAS_VPTO_BRIDGE_WHITELIST environment variable, " + "then to the built-in default whitelist"> + ]; + let dependentDialects = [ + "mlir::pto::PTODialect", + "mlir::arith::ArithDialect" + ]; +} + +def PTOLowerDeclarativeBridgeOps : Pass<"pto-lower-declarative-bridge-ops", "mlir::func::FuncOp"> { + let summary = "Lower whitelist-routed ops through the declarative VPTO " + "bridge channel"; + let description = [{ + Declarative lowering channel of the VPTO C++ interface bridge, and the + default channel of every whitelist entry. It rewrites each entry into a + void pto.bridge_call using only the whitelist description: each abi row + binds a wrapper argument to an IR operand position whose planned + alloc_tile address becomes the i64 call argument, and the wrapper + specialization is collected from the operand tile types (keyed by the + abi role) plus optional enum attributes. Ops needing family semantics + (storage lifecycle, address rebinding) opt out with `lowering: custom` + and stay on their family pass. Routing is whitelist driven: unrouted + ops keep flowing through their regular non-bridge lowering path. + }]; + let constructor = "mlir::pto::createPTOLowerDeclarativeBridgeOpsPass()"; + let options = [ + Option<"whitelistPath", "whitelist-path", "std::string", + /*default=*/"\"\"", + "Path to the VPTO bridge whitelist YAML; falls back to the " + "PTOAS_VPTO_BRIDGE_WHITELIST environment variable, " + "then to the built-in default whitelist"> + ]; + let dependentDialects = [ + "mlir::pto::PTODialect" + ]; +} + +def VPTOBridgeLowering : Pass<"vpto-bridge-lowering", "ModuleOp"> { + let summary = "Lower generic bridge ops into C++ wrapper calls"; + let description = [{ + Generic bridge lowering pass of the VPTO C++ interface bridge. It knows + nothing about individual PTO-ISA interface families: it validates each + pto.bridge_call against the bridge whitelist and mechanically lowers it + into a call to the wrapper entry, materializing the wrapper declaration + at module level. Entries carrying storage_size_callee additionally + synthesize the stateful-object pattern (size query + stack storage). + pto.bridge_inttoptr is lowered to llvm.inttoptr. + }]; + let constructor = "mlir::pto::createVPTOBridgeLoweringPass()"; + let options = [ + Option<"whitelistPath", "whitelist-path", "std::string", + /*default=*/"\"\"", + "Path to the VPTO bridge whitelist YAML; falls back to the " + "PTOAS_VPTO_BRIDGE_WHITELIST environment variable, " + "then to the built-in default whitelist"> + ]; + let dependentDialects = [ + "mlir::pto::PTODialect", + "mlir::func::FuncDialect", + "mlir::LLVM::LLVMDialect" + ]; +} + +def VPTOBridgeWrapperGen : Pass<"pto-emit-vpto-bridge-wrapper", "ModuleOp"> { + let summary = "Render the VPTO bridge wrapper source from the collected spec"; + let description = [{ + Wrapper generation pass of the VPTO C++ interface bridge. It reads the + pipe bridge specialization collected by the family passes in the + `pto.vpto.bridge.spec` module attribute (TPipe/Tile/split C++ template + tokens plus the whitelist wrapper entry names), renders the complete + bridge wrapper C++ source, and stores it in the + `pto.vpto.bridge.wrapper_source` module attribute for object emission to + compile and link. The spec attribute is removed once rendered. Modules + without a bridge spec are left untouched. + }]; + let constructor = "mlir::pto::createVPTOBridgeWrapperGenPass()"; + let options = [ + Option<"whitelistPath", "whitelist-path", "std::string", + /*default=*/"\"\"", + "Path to the VPTO bridge whitelist YAML; falls back to the " + "PTOAS_VPTO_BRIDGE_WHITELIST environment variable, " + "then to the built-in default whitelist"> + ]; +} + def VPTOSplitCVModule : Pass<"vpto-split-cv-module", "ModuleOp"> { let summary = "Split a VPTO module with cube/vector sections into kernel modules"; let description = [{ diff --git a/include/PTO/Transforms/TileOpExpansionUtils.h b/include/PTO/Transforms/TileOpExpansionUtils.h index f2aa3960c6..1501712905 100644 --- a/include/PTO/Transforms/TileOpExpansionUtils.h +++ b/include/PTO/Transforms/TileOpExpansionUtils.h @@ -22,7 +22,8 @@ inline bool isTileLibExpandableOp(Operation *op) { } return !isa(op); + TFreeFromAicOp, TFreeFromAivOp, TAllocOp, TPushOp, TPopOp, + TFreeOp>(op); } } // namespace mlir::pto diff --git a/include/PTO/Transforms/VPTOBridgeSpecCollector.h b/include/PTO/Transforms/VPTOBridgeSpecCollector.h new file mode 100644 index 0000000000..1574fee79d --- /dev/null +++ b/include/PTO/Transforms/VPTOBridgeSpecCollector.h @@ -0,0 +1,116 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// http://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- VPTOBridgeSpecCollector.h - bridge spec collection -------*- C++ -*-===// +//===----------------------------------------------------------------------===// +// +// Per-function collection of the VPTO bridge specialization fields that the +// wrapper generation pass merges into the module spec. Both the declarative +// channel and the pipe family pass write through this collector so the +// same-key policies and the kBridgeFuncSpecAttrName attribute shape live in +// one place. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGESPECCOLLECTOR_H +#define MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGESPECCOLLECTOR_H + +#include "PTO/Transforms/VPTOBridgeTokens.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/IR/BuiltinAttributes.h" +#include "mlir/IR/Operation.h" +#include "llvm/ADT/SmallVector.h" +#include "llvm/ADT/StringMap.h" +#include "llvm/ADT/StringRef.h" +#include +#include + +namespace mlir { +namespace pto { + +/// Collects the per-function bridge specialization fields written as the +/// kBridgeFuncSpecAttrName function attribute. +class BridgeSpecCollector { +public: + /// Adds a spec field with same-value dedup: several ops sharing a tile + /// shape or a wrapper entry write the same token under one key, which is + /// harmless; a different token for a key already written is diagnosed as + /// a conflict on `op` (the wrapper renders one token per spec field). + void addField(Operation *op, llvm::StringRef key, llvm::StringRef token); + + /// Adds a spec field that may be written at most once per function (the + /// pipe family's single producer/consumer pair); any repeat is diagnosed + /// on `op`. + void addUniqueField(Operation *op, llvm::StringRef key, + llvm::StringRef token); + + /// Returns whether any conflict was diagnosed. + bool hadError() const { return hadError_; } + + /// Stores the collected fields as the kBridgeFuncSpecAttrName function + /// attribute. No-op when nothing was collected. + void store(func::FuncOp func) const; + +private: + llvm::SmallVector> fields; + llvm::StringMap written; + bool hadError_ = false; +}; + +inline void BridgeSpecCollector::addField(Operation *op, llvm::StringRef key, + llvm::StringRef token) { + auto inserted = written.try_emplace(key, token); + if (inserted.second) { + fields.emplace_back(key.str(), token.str()); + return; + } + if (inserted.first->second != token) { + op->emitError() << "VPTO bridge spec field '" << key + << "' was already collected as '" + << inserted.first->second + << "'; the wrapper renders one token per spec field"; + hadError_ = true; + } +} + +inline void BridgeSpecCollector::addUniqueField(Operation *op, + llvm::StringRef key, + llvm::StringRef token) { + if (written.count(key)) { + op->emitError() + << "VPTO bridge spec field '" << key + << "' was already collected; only one bridged producer/consumer pair " + "per function is supported"; + hadError_ = true; + return; + } + written.try_emplace(key, token); + fields.emplace_back(key.str(), token.str()); +} + +inline void BridgeSpecCollector::store(func::FuncOp func) const { + if (fields.empty()) { + return; + } + SmallVector specAttrs; + specAttrs.reserve(fields.size()); + for (const auto &field : fields) { + specAttrs.push_back({StringAttr::get(func.getContext(), field.first), + StringAttr::get(func.getContext(), field.second)}); + } + func->setAttr(kBridgeFuncSpecAttrName, + DictionaryAttr::get(func.getContext(), specAttrs)); +} + +} // namespace pto +} // namespace mlir + +#endif // MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGESPECCOLLECTOR_H diff --git a/include/PTO/Transforms/VPTOBridgeTokens.h b/include/PTO/Transforms/VPTOBridgeTokens.h new file mode 100644 index 0000000000..a506ffbc33 --- /dev/null +++ b/include/PTO/Transforms/VPTOBridgeTokens.h @@ -0,0 +1,109 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- VPTOBridgeTokens.h - C++ template token building ---------*- C++ -*-===// +//===----------------------------------------------------------------------===// +// +// Bridge-side construction of the PTO-ISA C++ template tokens used by the +// generated VPTO bridge wrapper. These utilities are the bridge analogue of +// EmitC's token builders: the IR-fact -> C++ spelling mapping rules are +// shared through PTOCppTokens, while the bridge assembly rules (fully +// qualified spellings, NoneBox trailing-argument omission) live here. +// +// The tokens are fully qualified C++ type/constant spellings (e.g. +// "pto::TPipe<0, pto::Direction::DIR_C2V, 1024, 8, 2, false>") suitable for +// direct substitution into the generated wrapper source. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGETOKENS_H +#define MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGETOKENS_H + +#include "mlir/IR/Types.h" +#include "mlir/Support/LLVM.h" +#include "mlir/Support/LogicalResult.h" +#include "llvm/ADT/StringRef.h" +#include + +namespace mlir { +namespace pto { + +class InitializeL2LPipeOp; +class TileBufType; + +/// Module attribute carrying the collected pipe bridge specialization (a +/// DictionaryAttr of StringAttr token fields). Written by the pipe family +/// pass, consumed by the bridge wrapper generation pass. +constexpr llvm::StringLiteral kBridgeSpecAttrName = "pto.vpto.bridge.spec"; + +/// Function attribute carrying one function's pipe bridge specialization +/// (a DictionaryAttr with the same keys as the module spec). Written by the +/// pipe family pass; the wrapper generation pass merges the per-function +/// specs into the module spec and removes them. The family pass instances +/// may run concurrently, so the shared module attribute is only written by +/// the single-threaded module-level pass. +constexpr llvm::StringLiteral kBridgeFuncSpecAttrName = + "pto.vpto.bridge.func_spec"; + +/// Module attribute carrying the rendered bridge wrapper C++ source (a +/// StringAttr). Written by the bridge wrapper generation pass, consumed by +/// object emission. +constexpr llvm::StringLiteral kBridgeWrapperSourceAttrName = + "pto.vpto.bridge.wrapper_source"; + +/// Spec DictionaryAttr keys for the pipe bridge specialization. +constexpr llvm::StringLiteral kBridgeSpecPipeKey = "pipe"; +constexpr llvm::StringLiteral kBridgeSpecProducerTileKey = "producer_tile"; +constexpr llvm::StringLiteral kBridgeSpecConsumerTileKey = "consumer_tile"; +constexpr llvm::StringLiteral kBridgeSpecSplitKey = "split"; +constexpr llvm::StringLiteral kBridgeSpecEntryInitKey = "entry.init"; +constexpr llvm::StringLiteral kBridgeSpecEntrySizeKey = "entry.size"; +constexpr llvm::StringLiteral kBridgeSpecEntryPushKey = "entry.push"; +constexpr llvm::StringLiteral kBridgeSpecEntryPopKey = "entry.pop"; +constexpr llvm::StringLiteral kBridgeSpecEntryFreeKey = "entry.free"; + +/// Spec DictionaryAttr keys of declarative-wrapper specializations. The +/// tile tokens are collected under the whitelist abi role of each operand +/// (these constants document the roles the built-in whitelist uses); the +/// acc phase token is only collected for a non-Unspecified phase. +constexpr llvm::StringLiteral kBridgeSpecLeftTileKey = "left_tile"; +constexpr llvm::StringLiteral kBridgeSpecRightTileKey = "right_tile"; +constexpr llvm::StringLiteral kBridgeSpecResultTileKey = "result_tile"; +constexpr llvm::StringLiteral kBridgeSpecAccInTileKey = "acc_in_tile"; +constexpr llvm::StringLiteral kBridgeSpecBiasTileKey = "bias_tile"; +constexpr llvm::StringLiteral kBridgeSpecAScaleTileKey = "a_scale_tile"; +constexpr llvm::StringLiteral kBridgeSpecBScaleTileKey = "b_scale_tile"; +constexpr llvm::StringLiteral kBridgeSpecAccPhaseKey = "acc_phase"; + +/// Builds the fully qualified `pto::TPipe` token from a local-to-local pipe init op. +/// Fails when the op lacks a flag_base attribute or carries an unsupported +/// dir_mask. The L2L pipe uses a fixed localSlotNum of 2. +FailureOr buildBridgePipeToken(InitializeL2LPipeOp init); + +/// Builds the fully qualified `pto::TileSplitAxis::TILE_*` token for a split +/// value (0..4). Fails for values outside that range. +FailureOr buildBridgeTileSplitToken(int64_t split); + +/// Builds the fully qualified `pto::Tile` token from a tile +/// buffer type. The SLayout/SFractalSize template arguments are emitted only +/// for boxed (non-NoneBox) storage layouts, matching the wrapper's Tile +/// specializations. Fails when the type lacks a resolvable address space or +/// element type. +FailureOr buildBridgeTileToken(TileBufType tile); + +/// Builds the C++ element type token (e.g. "float", "half", "int8_t") for an +/// MLIR element type. Falls back to "float" for unrecognized types, mirroring +/// the EmitC element token behavior. +std::string buildBridgeElementTypeToken(Type elementType); + +} // namespace pto +} // namespace mlir + +#endif // MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGETOKENS_H diff --git a/include/PTO/Transforms/VPTOBridgeWhitelist.h b/include/PTO/Transforms/VPTOBridgeWhitelist.h new file mode 100644 index 0000000000..6647c7a34e --- /dev/null +++ b/include/PTO/Transforms/VPTOBridgeWhitelist.h @@ -0,0 +1,293 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- VPTOBridgeWhitelist.h - C++ bridge whitelist --------------*- C++ -*-===// +//===----------------------------------------------------------------------===// +// +// Declarative description of which IR ops are routed to the VPTO C++ +// interface bridge and how their arguments map onto wrapper ABI values. +// The generic bridge lowering pass consumes this table to validate bridge +// calls; wrapper generation consumes it to synthesize wrapper sources. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGEWHITELIST_H +#define MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGEWHITELIST_H + +#include "mlir/IR/Types.h" +#include "mlir/Support/LLVM.h" +#include "llvm/ADT/StringRef.h" +#include +#include + +namespace mlir { +namespace pto { + +/// ABI argument of a wrapper entry. `type` is one of the supported carrier +/// tokens: "ptr", "i64", or "i32" (declarative entries may omit it, the +/// parser fills in the "i64" tile-address default). Declarative entries +/// additionally bind each argument to an IR operand position (`operand`, +/// positional because MLIR exposes no generic ODS operand-name +/// reflection) and the template role the operand's tile token is collected +/// under (`role`, which is also the spec key and the source of the entry's +/// tile typedef name). `arg` is the diagnostic label and rendered parameter +/// name (the ODS operand name); it defaults to the lowerCamelCase of the +/// role (see bridgeRoleParamName), so a binding registers with +/// operand/role only. +struct BridgeAbiArg { + std::string type; + int64_t operand = -1; + std::string arg; + std::string role; +}; + +/// Declarative template-argument mapping row: an IR field (`source` + +/// `field`) feeds a C++ template slot (`target`). Consumed by wrapper +/// generation to validate that the collected specialization covers the +/// declared slots; the authoritative token construction lives in +/// VPTOBridgeTokens. Declarative entries carry only `source: attr` rows +/// (an enum attribute mapped to a template slot, with `enumType` providing +/// the qualified C++ enum spelling and `omitValue` the case that renders +/// no template argument); their tile typedefs derive from the abi roles, +/// so tile rows are rejected at parse time. +struct BridgeTmplMapField { + std::string source; + std::string field; + std::string target; + std::string enumType; + std::string omitValue; +}; + +/// One whitelist row: an IR op routed to a wrapper entry of a bridged +/// PTO-ISA C++ interface. +struct BridgeWhitelistEntry { + /// IR op name, e.g. "pto.tpush". The whitelist is the routing table the + /// family passes consult; "internal" marks wrapper-internal helpers + /// (e.g. the size query entry) that are never routed from an IR op. + std::string op; + /// Wrapper source this entry is rendered into, e.g. "pipe". Entries + /// sharing a wrapper share one C++ translation unit, one template + /// specialization and one typedef set, so this is the unit of wrapper + /// generation rather than a taxonomy of ops: the N IR ops of the pipe + /// protocol (init/push/pop/free plus the internal size query) all name + /// wrapper "pipe" and render into a single source. Wrapper generation + /// dispatches on it to pick the renderer, and a module may use entries of + /// exactly one wrapper. It also scopes the tmpl_map source validation of + /// custom-channel entries. + std::string wrapper; + /// Lowering channel. "declarative" (the default) routes the op through the + /// generic declarative bridge lowering: a mechanical operand-adapter + /// mapping driven entirely by this table, needing no pass code. "custom" + /// opts the entry out of that channel because it carries real family + /// semantics (storage lifecycle, address rebinding) that only a dedicated + /// pass can express. + /// + /// The default is deliberately the generic channel: adding a mechanically + /// mapped interface family must cost zero pass code *and* zero ceremony, + /// while the exception stays explicit. It also makes the common mistake + /// self-detecting -- an entry that needs a custom pass but forgets the tag + /// lacks the operand/arg/role bindings the declarative channel requires, + /// so it is rejected at parse time with the missing field named, instead + /// of surviving until the post-lowering leftover check. + std::string lowering = "declarative"; + /// Wrapper entry name, e.g. "pto_vpto_pipe_push". This is the callee + /// the generic bridge lowering emits. Optional for declarative routed + /// entries, which default to a name derived from the op name (see + /// deriveDefaultBridgeEntry); custom entries must declare it. + std::string entry; + /// C++ call spelling the generic declarative renderer emits for this + /// entry, e.g. "pto::TMATMUL" or "TADD". Declarative routed entries + /// only; the call arguments are the abi-bound tiles in declaration + /// order. Optional: it defaults to a spelling derived from the op name + /// (see deriveDefaultBridgeCall); a wrong derivation fails loudly when + /// the generated wrapper source is compiled. + std::string call; + /// Template arguments rendered between the call spelling and its + /// argument list. Each item is either the `field` of one of this + /// entry's attr tmpl_map rows (rendered as the spec token collected + /// for that field; the whole template argument list is omitted when + /// the spec carries no token) or a literal qualified C++ spelling + /// (contains "::"). Declarative routed entries only. + std::vector tmplArgs; + /// Call-side ABI of the wrapper entry, including any synthesized + /// arguments such as the storage pointer of stateful entries. + std::vector abi; + /// Wrapper entry returning the size of the stateful object owned by this + /// entry. Declared on stateful entries (e.g. the pipe init) and consumed + /// by the family pass as the bridge call storage_size_callee. + std::string storageSizeEntry; + /// Declarative IR-field -> C++ template-slot mappings for wrapper + /// generation. Optional; empty when the entry needs no template mapping. + std::vector tmplMap; + + /// Returns whether the op lowers through the generic declarative channel. + bool isDeclarative() const { return lowering == kLoweringDeclarative; } + + /// `lowering` value (the default) routing the op through the generic + /// declarative bridge lowering. + static constexpr llvm::StringLiteral kLoweringDeclarative = "declarative"; + /// `lowering` value opting the entry out of the declarative channel: a + /// dedicated family pass owns the rewrite. + static constexpr llvm::StringLiteral kLoweringCustom = "custom"; +}; + +/// Declaration of a wrapper rendered by the generic declarative renderer: +/// the family knowledge no whitelist entry carries. `includes` lists the +/// PTO-ISA headers the wrapper translation unit includes; `core` selects +/// the core guard the entries render under ("cube" -> __DAV_CUBE__, +/// "vec" -> __DAV_VEC__, "both" -> no guard). `core` is optional: it +/// defaults to the kind of the tile address spaces the declarative +/// lowering collects for the wrapper (VEC tiles -> "vec", any cube-family +/// tile -> "cube"), so a single-core wrapper registers without it; a +/// wrapper whose used entries collected no tile declares it explicitly. +/// Wrappers whose entries carry `lowering: custom` own a dedicated +/// renderer and must not be declared here. +struct BridgeWrapperDecl { + std::string name; + std::vector includes; + std::string core; +}; + +/// Parsed whitelist document. +struct BridgeWhitelist { + std::vector bridgeOps; + std::vector wrappers; + + /// Returns the wrapper declaration named `name`, or nullptr. + const BridgeWrapperDecl *findWrapper(llvm::StringRef name) const { + for (const BridgeWrapperDecl &decl : wrappers) { + if (decl.name == name) { + return &decl; + } + } + return nullptr; + } + + /// Returns whether any entry routing into the wrapper `name` carries + /// `lowering: custom` (and thus needs a dedicated renderer). + bool wrapperHasCustomEntry(llvm::StringRef name) const { + for (const BridgeWhitelistEntry &entry : bridgeOps) { + if (entry.wrapper == name && !entry.isDeclarative()) { + return true; + } + } + return false; + } + + /// Returns the entry whose wrapper name is `entryName`, or nullptr. + const BridgeWhitelistEntry *findEntry(llvm::StringRef entryName) const { + for (const BridgeWhitelistEntry &entry : bridgeOps) { + if (entry.entry == entryName) { + return &entry; + } + } + return nullptr; + } + + /// Returns the entry routing the IR op `opName` (e.g. "pto.tpush"), or + /// nullptr. Wrapper-internal helpers (op == "internal") are never routed. + const BridgeWhitelistEntry *findOp(llvm::StringRef opName) const { + for (const BridgeWhitelistEntry &entry : bridgeOps) { + if (entry.op == opName && entry.op != kInternalOp) { + return &entry; + } + } + return nullptr; + } + + /// Marker `op` value of wrapper-internal helper entries that no IR op + /// routes to (e.g. the stateful-object size query). + static constexpr llvm::StringLiteral kInternalOp = "internal"; +}; + +/// Parses a whitelist YAML file. Diagnostics are written to `diagOS`. +/// Rejects unreadable files, YAML syntax errors, empty fields, duplicate +/// wrapper entry names, duplicate routed op names, unsupported ABI type +/// tokens, and dangling storage_size_entry references. +FailureOr parseBridgeWhitelist(llvm::StringRef path, + llvm::raw_ostream &diagOS); + +/// Parses a whitelist YAML document already in memory; `sourceName` is used +/// in diagnostics (e.g. a file path or the built-in whitelist marker). +FailureOr +parseBridgeWhitelistFromBuffer(llvm::StringRef content, + llvm::StringRef sourceName, + llvm::raw_ostream &diagOS); + +/// Resolves the whitelist path from a pass `whitelist-path` option value, +/// falling back to the PTOAS_VPTO_BRIDGE_WHITELIST environment variable. +/// Returns an empty string when neither is configured. +std::string resolveBridgeWhitelistPath(llvm::StringRef optionValue); + +/// Source name used in diagnostics when the built-in default whitelist is +/// in effect. +constexpr llvm::StringLiteral kBuiltinBridgeWhitelistSource = + ""; + +/// Loads the bridge whitelist through the formal resolution chain: pass +/// `whitelist-path` option, then PTOAS_VPTO_BRIDGE_WHITELIST, then the +/// built-in default whitelist (pipe + matmul families) shipped with ptoas. +/// Always returns a parsed whitelist unless the explicitly configured file +/// fails to parse. When `sourceName` is non-null it receives the resolved +/// source name (file path or the built-in marker) for diagnostics. +FailureOr loadBridgeWhitelist(llvm::StringRef optionValue, + llvm::raw_ostream &diagOS, + std::string *sourceName = nullptr); + +/// Returns whether `token` is one of the ABI carrier tokens accepted by the +/// generic bridge lowering ("ptr", "i64", "i32"). The set stays closed so +/// whitelist parsing and lowering agree on the carriers. +bool isSupportedBridgeAbiType(llvm::StringRef token); + +/// Returns whether the ABI carrier token describes `type` after bridge +/// lowering conversion ("ptr" -> LLVM pointer, "i64"/"i32" -> the integer +/// widths). +bool bridgeAbiTypeMatches(llvm::StringRef token, Type type); + +/// tmpl_map `source` tokens naming IR producers of a template argument. +constexpr llvm::StringLiteral kPipeInitTmplMapSource = "pipe.init"; +constexpr llvm::StringLiteral kTileTmplMapSource = "tile"; +/// tmpl_map `source` token mapping an enum attribute of the routed op to a +/// template slot. +constexpr llvm::StringLiteral kAttrTmplMapSource = "attr"; + +/// `core` values accepted by a wrapper declaration. +constexpr llvm::StringLiteral kBridgeWrapperCoreCube = "cube"; +constexpr llvm::StringLiteral kBridgeWrapperCoreVec = "vec"; +constexpr llvm::StringLiteral kBridgeWrapperCoreBoth = "both"; + +/// Derives the default wrapper entry name of a declarative routed entry +/// from its IR op name: strip the `pto.` prefix, drop the tile-world `t` +/// mnemonic lead, replace dots with underscores and prepend `pto_vpto_` +/// (`pto.tmatmul.mx.acc` -> `pto_vpto_matmul_mx_acc`). +std::string deriveDefaultBridgeEntry(llvm::StringRef opName); + +/// Derives the default C++ call spelling of a declarative routed entry +/// from its IR op name: strip the `pto.` prefix (keeping the tile-world +/// `t` mnemonic lead the entry name drops), uppercase the remainder, +/// replace dots with underscores and qualify with `pto::` +/// (`pto.tadd` -> `pto::TADD`, `pto.tmatmul.mx` -> `pto::TMATMUL_MX`). +/// Variants whose interface call does not follow the convention declare +/// `call` explicitly. +std::string deriveDefaultBridgeCall(llvm::StringRef opName); + +/// Renders the lowerCamelCase parameter name of an abi role +/// (`left_tile` -> `leftTile`, `a_scale_tile` -> `aScaleTile`), the +/// default `arg` of declarative abi bindings. +std::string bridgeRoleParamName(llvm::StringRef role); + +/// Renders the CamelCase typedef target name of an abi role +/// (`left_tile` -> `LeftTile`, `a_scale_tile` -> `AScaleTile`). Tile +/// typedefs of declarative entries are role driven, so this is the single +/// source of the typedef names the wrapper bodies reference. +std::string bridgeRoleTypedefTarget(llvm::StringRef role); + +} // namespace pto +} // namespace mlir + +#endif // MLIR_DIALECT_PTO_TRANSFORMS_VPTOBRIDGEWHITELIST_H diff --git a/lib/PTO/IR/PTO.cpp b/lib/PTO/IR/PTO.cpp index 48602ff8ce..4347c4bd25 100644 --- a/lib/PTO/IR/PTO.cpp +++ b/lib/PTO/IR/PTO.cpp @@ -243,6 +243,8 @@ static ParseResult parseI32LiteralAttr(OpAsmParser &parser, IntegerAttr &attr); #define GET_ATTRDEF_CLASSES #include "PTO/IR/PTOAttrs.cpp.inc" +#include "PTO/IR/PTOAttrInterfaces.cpp.inc" + #include "PTO/IR/PTODialect.cpp.inc" [[maybe_unused]] static LogicalResult parseShapeAndElemStable(mlir::AsmParser &parser, @@ -18147,8 +18149,11 @@ getEnclosingFunctionKernelKind(Operation *op) { } static bool isInsideSectionOrAttributedKernel(Operation *op) { - return isInsideSectionCube(op) || isInsideSectionVector(op) || - isInsideTileOpHelper(op) || getEnclosingFunctionKernelKind(op).has_value(); + if (isInsideSectionCube(op) || isInsideSectionVector(op) || + isInsideTileOpHelper(op) || getEnclosingFunctionKernelKind(op).has_value()) + return true; + auto module = op->getParentOfType(); + return module && module->hasAttr(FunctionKernelKindAttr::name); } static LogicalResult verifySplitAttr(Operation *op, int64_t split) { @@ -21807,6 +21812,21 @@ void TFreeOp::getEffects( addEffect(effects, &getPipeHandleMutable(), MemoryEffects::Write::get()); } +void BridgeCallOp::getEffects( + SmallVectorImpl> + &effects) { + for (OpOperand &arg : getArgsMutable()) { + addEffect(effects, &arg, MemoryEffects::Read::get()); + } + // The wrapper may mutate state the bridge cannot model (e.g. a FIFO or + // the storage it is handed), so the call itself is conservatively marked + // as reading and writing the default resource. + effects.emplace_back(MemoryEffects::Read::get(), + SideEffects::DefaultResource::get()); + effects.emplace_back(MemoryEffects::Write::get(), + SideEffects::DefaultResource::get()); +} + void SetQuantScalarOp::getEffects( SmallVectorImpl> &effects) { diff --git a/lib/PTO/IR/PTOAttrs.cpp b/lib/PTO/IR/PTOAttrs.cpp index a6b4a60a7f..c0b20f7c49 100644 --- a/lib/PTO/IR/PTOAttrs.cpp +++ b/lib/PTO/IR/PTOAttrs.cpp @@ -35,6 +35,18 @@ constexpr int32_t kCompactModeRowPlusOne = } // namespace +llvm::StringRef AccPhaseAttr::getEnumCaseSymbol() const { + switch (getValue()) { + case AccPhase::Unspecified: + return "Unspecified"; + case AccPhase::Partial: + return "Partial"; + case AccPhase::Final: + return "Final"; + } + llvm_unreachable("unknown PTO AccPhase case"); +} + TileBufConfigAttr TileBufConfigAttr::getDefault(MLIRContext *ctx) { Builder b(ctx); BLayoutAttr bl = BLayoutAttr::get(ctx, BLayout::RowMajor); diff --git a/lib/PTO/Transforms/CMakeLists.txt b/lib/PTO/Transforms/CMakeLists.txt index eb92bf1edb..e788f6767e 100644 --- a/lib/PTO/Transforms/CMakeLists.txt +++ b/lib/PTO/Transforms/CMakeLists.txt @@ -95,6 +95,13 @@ add_mlir_dialect_library(PTOTransforms InsertTemplateAttributes.cpp ExpandTileOp.cpp FoldTileBufIntrinsics.cpp + PTOLowerPipeFamilyOps.cpp + PTOLowerDeclarativeBridgeOps.cpp + VPTOBridgeLowering.cpp + PTOCppTokens.cpp + VPTOBridgeTokens.cpp + VPTOBridgeWhitelist.cpp + VPTOBridgeWrapperGen.cpp LowerPTOToUBufOps.cpp PTOLowerToOpLibCalls.cpp PTOInstantiateAndInlineOpLib.cpp diff --git a/lib/PTO/Transforms/FoldTileBufIntrinsics.cpp b/lib/PTO/Transforms/FoldTileBufIntrinsics.cpp index db074fbdd6..78efcdad29 100644 --- a/lib/PTO/Transforms/FoldTileBufIntrinsics.cpp +++ b/lib/PTO/Transforms/FoldTileBufIntrinsics.cpp @@ -223,6 +223,14 @@ static bool isSCFTileCarrier(Value value) { isa_and_nonnull(blockArg.getOwner()->getParentOp()); } +/// Returns whether the declared tile is rebound by a TPOP. The VPTO pipe +/// bridge resolves such tiles at runtime from the FIFO slot address returned +/// by the pop, so folding their address here would break address propagation. +static bool isDeclareTileReboundByTPop(pto::DeclareTileOp decl) { + return llvm::any_of(decl.getResult().getUsers(), + [](Operation *user) { return isa(user); }); +} + static std::optional resolveTileHandle(Value tileBuf, Operation *user) { if (auto regionResult = dyn_cast(tileBuf)) { @@ -257,6 +265,17 @@ static std::optional resolveTileHandle(Value tileBuf, alloc.getValidCol(), tileTy.getConfigAttr()}; } + if (auto decl = tileBuf.getDefiningOp()) { + if (!isDeclareTileReboundByTPop(decl)) { + user->emitError("FoldTileBufIntrinsics: pto.declare_tile address requires " + "a matching pto.tpop rebinding"); + return std::nullopt; + } + auto tileTy = cast(decl.getResult().getType()); + return TileHandleInfo{decl.getResult(), Value{}, Value{}, + tileTy.getConfigAttr()}; + } + if (auto reshape = tileBuf.getDefiningOp()) { auto sourceInfo = resolveTileHandle(reshape.getSrc(), user); if (!sourceInfo) { @@ -279,7 +298,8 @@ static std::optional resolveTileHandle(Value tileBuf, user->emitError("FoldTileBufIntrinsics: expected tile_buf to be defined by " "the active materialized tile-handle bridge " - "(pto.alloc_tile or pto.treshape, " + "(pto.alloc_tile, pto.declare_tile rebound by pto.tpop, or " + "pto.treshape, " "or a pto.fusion_region result that yields one of them)"); return std::nullopt; } @@ -811,6 +831,13 @@ struct FoldTileBufIntrinsicsPass continue; } + // A declare_tile rebound by TPOP gets its address at runtime from the + // FIFO slot; folding it to the placeholder value here would break the + // VPTO pipe bridge address propagation. + if (auto decl = addrOp.getSrc().getDefiningOp(); + decl && isDeclareTileReboundByTPop(decl)) + continue; + auto handleInfo = resolveTileHandle(addrOp.getSrc(), addrOp); if (!handleInfo) { return signalPassFailure(); diff --git a/lib/PTO/Transforms/PTOCppTokens.cpp b/lib/PTO/Transforms/PTOCppTokens.cpp new file mode 100644 index 0000000000..18cbe7a8eb --- /dev/null +++ b/lib/PTO/Transforms/PTOCppTokens.cpp @@ -0,0 +1,156 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"); you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// http://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- PTOCppTokens.cpp - shared PTO-ISA C++ token mappings --------------===// +//===----------------------------------------------------------------------===// +// +// Implementation of the shared IR-fact -> C++ spelling mapping functions; +// see include/PTO/Transforms/PTOCppTokens.h. +// +//===----------------------------------------------------------------------===// + +#include "PTO/Transforms/PTOCppTokens.h" +#include "PTO/IR/PTOTypeUtils.h" + +using namespace mlir; +using namespace mlir::pto; + +std::string pto::getPTOCppElementTypeToken(Type elementType) { + if (pto::isPTOFloat8E4M3LikeType(elementType)) + return "float8_e4m3_t"; + if (pto::isPTOFloat8E5M2LikeType(elementType)) + return "float8_e5m2_t"; + if (pto::isPTOF8E8M0Type(elementType)) + return "float8_e8m0_t"; + if (isa(elementType)) + return "hifloat8_t"; + if (isa(elementType)) + return "float4_e1m2x2_t"; + if (isa(elementType)) + return "float4_e2m1x2_t"; + if (elementType.isF16()) + return "half"; + if (elementType.isBF16()) + return "bfloat16_t"; + if (elementType.isF32()) + return "float"; + if (elementType.isF64()) + return "double"; + if (elementType.isInteger(8)) + return (elementType.isSignlessInteger(8) || + elementType.isSignedInteger(8)) + ? "int8_t" + : "uint8_t"; + if (elementType.isInteger(16)) + return (elementType.isSignlessInteger(16) || + elementType.isSignedInteger(16)) + ? "int16_t" + : "uint16_t"; + if (elementType.isInteger(32)) + return (elementType.isSignlessInteger(32) || + elementType.isSignedInteger(32)) + ? "int32_t" + : "uint32_t"; + if (elementType.isInteger(64)) + return cast(elementType).isUnsigned() ? "uint64_t" + : "int64_t"; + return "float"; +} + +FailureOr pto::getPTOCppTileSplitToken(int64_t split, + StringRef qualifier) { + switch (split) { + case 0: + return (qualifier + "TileSplitAxis::TILE_NO_SPLIT").str(); + case 1: + return (qualifier + "TileSplitAxis::TILE_UP_DOWN").str(); + case 2: + return (qualifier + "TileSplitAxis::TILE_LEFT_RIGHT").str(); + case 3: + return (qualifier + "TileSplitAxis::TILE_UP_DOWN_ODD").str(); + case 4: + return (qualifier + "TileSplitAxis::TILE_LEFT_RIGHT_ODD").str(); + default: + return failure(); + } +} + +FailureOr pto::getPTOCppDirectionToken(int8_t dirMask, + StringRef qualifier) { + switch (dirMask) { + case 1: + return (qualifier + "Direction::DIR_C2V").str(); + case 2: + return (qualifier + "Direction::DIR_V2C").str(); + case 3: + return (qualifier + "Direction::DIR_BOTH").str(); + default: + return failure(); + } +} + +FailureOr pto::getPTOCppTileTypeToken(AddressSpace addressSpace, + StringRef qualifier) { + switch (addressSpace) { + case AddressSpace::MAT: + return (qualifier + "TileType::Mat").str(); + case AddressSpace::LEFT: + return (qualifier + "TileType::Left").str(); + case AddressSpace::RIGHT: + return (qualifier + "TileType::Right").str(); + case AddressSpace::ACC: + return (qualifier + "TileType::Acc").str(); + case AddressSpace::VEC: + return (qualifier + "TileType::Vec").str(); + case AddressSpace::BIAS: + return (qualifier + "TileType::Bias").str(); + case AddressSpace::SCALING: + return (qualifier + "TileType::Scaling").str(); + default: + return failure(); + } +} + +FailureOr pto::getPTOCppBLayoutToken(BLayout bLayout, + StringRef qualifier) { + switch (bLayout) { + case BLayout::RowMajor: + return (qualifier + "BLayout::RowMajor").str(); + case BLayout::ColMajor: + return (qualifier + "BLayout::ColMajor").str(); + } + return failure(); +} + +FailureOr pto::getPTOCppSLayoutToken(SLayout sLayout, + StringRef qualifier) { + switch (sLayout) { + case SLayout::NoneBox: + return (qualifier + "SLayout::NoneBox").str(); + case SLayout::RowMajor: + return (qualifier + "SLayout::RowMajor").str(); + case SLayout::ColMajor: + return (qualifier + "SLayout::ColMajor").str(); + } + return failure(); +} + +std::string pto::renderTPipeSpelling(int32_t flagBase, StringRef dirTok, + int32_t slotSize, int32_t slotNum, + int32_t localSlotNum, bool nosplit, + StringRef qualifier) { + std::string token = (qualifier + "TPipe<").str() + + std::to_string(flagBase) + ", " + dirTok.str() + ", " + + std::to_string(slotSize) + ", " + + std::to_string(slotNum) + ", " + + std::to_string(localSlotNum) + ", " + + (nosplit ? "true" : "false") + ">"; + return token; +} diff --git a/lib/PTO/Transforms/PTOLowerDeclarativeBridgeOps.cpp b/lib/PTO/Transforms/PTOLowerDeclarativeBridgeOps.cpp new file mode 100644 index 0000000000..3644259e47 --- /dev/null +++ b/lib/PTO/Transforms/PTOLowerDeclarativeBridgeOps.cpp @@ -0,0 +1,309 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- PTOLowerDeclarativeBridgeOps.cpp - declarative bridge lowering ----===// +//===----------------------------------------------------------------------===// +// +// Generic declarative lowering channel of the VPTO C++ interface bridge. +// It rewrites every whitelist entry marked `lowering: declarative` into a +// void pto.bridge_call using only the whitelist description: each abi row +// binds a wrapper argument to an IR operand position whose planned +// alloc_tile address becomes the i64 call argument, and the template +// specialization is collected from the operand tile types (keyed by the +// abi role) plus optional enum attributes (tmpl_map `source: attr` rows). +// No family semantics are understood here; ops needing storage lifecycle +// or address rebinding stay on their family pass (`lowering: family`). +// +// Routing is whitelist driven: an op the whitelist does not route (or +// routes to the family channel) is left untouched, so unrouted matmul ops +// keep flowing through the regular tile-op expansion path. +// +// The collected spec keys deliberately match the constants consumed by the +// wrapper generation pass: the role name is the tile spec key, the attr +// row `field` is the enum spec key, the entry spec key is derived +// from the op name (`pto.tmatmul.mx.acc` -> `entry.matmul_mx_acc`), and +// the reserved `core.` key carries the derived core guard of a +// wrapper declaration that omits `core`. +// +//===----------------------------------------------------------------------===// + +#include "PTO/IR/PTO.h" +#include "PTO/Transforms/VPTOBridgeSpecCollector.h" +#include "PTO/Transforms/VPTOBridgeTokens.h" +#include "PTO/Transforms/VPTOBridgeWhitelist.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/IR/Builders.h" +#include "mlir/Pass/Pass.h" +#include "llvm/ADT/SmallVector.h" +#include "llvm/ADT/StringMap.h" +#include +#include + +namespace mlir { +namespace pto { + +#define GEN_PASS_DECL_PTOLOWERDECLARATIVEBRIDGEOPS +#define GEN_PASS_DEF_PTOLOWERDECLARATIVEBRIDGEOPS +#include "PTO/Transforms/Passes.h.inc" + +namespace { + +/// Derives the wrapper entry spec key from the routed IR op name. +/// Tile-world ops carry the `pto.t` mnemonic prefix, which is not part of +/// the wrapper's entry naming (`pto.tmatmul.mx.acc` -> `entry.matmul_mx_acc`). +static std::string deriveEntrySpecKey(llvm::StringRef opName) { + constexpr llvm::StringLiteral kTileWorldOpPrefix = "pto.t"; + if (!opName.consume_front(kTileWorldOpPrefix)) { + opName.consume_front("pto."); + } + std::string key = ("entry." + opName).str(); + constexpr llvm::StringLiteral kEntryKeyPrefix = "entry."; + std::replace(key.begin() + kEntryKeyPrefix.size(), key.end(), '.', '_'); + return key; +} + +/// Derives the camelCase spelling of a snake_case whitelist field name. +/// ODS attribute names are camelCase ($accPhase) while the whitelist field +/// doubles as the spec key (acc_phase), so the attribute lookup tries both +/// spellings. +static std::string camelCaseFieldName(llvm::StringRef fieldName) { + std::string camel; + camel.reserve(fieldName.size()); + bool upperNext = false; + for (char c : fieldName) { + if (c == '_') { + upperNext = true; + continue; + } + camel.push_back(upperNext && c >= 'a' && c <= 'z' + ? static_cast(c - 'a' + 'A') + : c); + upperNext = false; + } + return camel; +} + +/// Maps a bridged tile to the core kind its wrapper renders under when the +/// wrapper declaration omits `core`: VEC tiles run on the vector core, the +/// cube-family tile spaces (mat/left/right/acc/bias/scaling) on the cube +/// core. Tiles without a supported address space fail earlier in +/// buildBridgeTileToken, so the cube default here is unreachable for a +/// successfully collected tile. +static llvm::StringLiteral bridgeCoreKindForTile(Value tile) { + auto tileTy = cast(tile.getType()); + auto spaceAttr = + dyn_cast_or_null(tileTy.getMemorySpace()); + if (spaceAttr && spaceAttr.getAddressSpace() == AddressSpace::VEC) + return kBridgeWrapperCoreVec; + return kBridgeWrapperCoreCube; +} + +struct PTOLowerDeclarativeBridgeOpsPass final + : public impl::PTOLowerDeclarativeBridgeOpsBase< + PTOLowerDeclarativeBridgeOpsPass> { + MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(PTOLowerDeclarativeBridgeOpsPass) + + void runOnOperation() override { + func::FuncOp func = getOperation(); + + // The whitelist always resolves through the formal chain (pass option, + // PTOAS_VPTO_BRIDGE_WHITELIST, built-in default); kernels that want the + // regular tile-op expansion route the op out of the whitelist with an + // explicit whitelist file. + FailureOr whitelistOr = + loadBridgeWhitelist(whitelistPath, llvm::errs()); + if (failed(whitelistOr)) { + signalPassFailure(); + return; + } + const BridgeWhitelist &whitelist = *whitelistOr; + + // Collect the ops routed to the declarative channel first; rewriting + // during the walk would invalidate the walker. Ops with no whitelist + // entry or with a family entry are left untouched: unrouted ops keep + // their non-bridge lowering (e.g. the matmul mad expansion), and family + // entries are rewritten by their family pass. + SmallVector> routed; + func.walk([&](Operation *op) { + const BridgeWhitelistEntry *entry = + whitelist.findOp(op->getName().getStringRef()); + if (entry && entry->isDeclarative()) { + routed.push_back({op, entry}); + } + }); + if (routed.empty()) { + return; + } + + bool hadError = false; + // Wrapper specialization fields collected while lowering this function; + // stored as a function attribute once lowering succeeds. The module-level + // wrapper generation pass merges the per-function specs deterministically + // (the pass instances may run concurrently). + BridgeSpecCollector spec; + // Tile handles consumed by bridged ops; erased once use-empty. + SmallVector bridgedAllocs; + + // Resolves a tile operand to its planned address. The bridge wrapper + // binds each tile to the address at runtime, so every operand must be an + // alloc_tile carrying a planned address. + auto resolvePlannedTile = [&](Operation *op, const BridgeAbiArg &abiArg, + Value tile) -> Value { + auto tileTy = dyn_cast(tile.getType()); + if (!tileTy) { + op->emitError() << "VPTO declarative bridge: operand #" << abiArg.operand + << " ('" << abiArg.arg << "', role " << abiArg.role + << ") must be a tile_buf"; + hadError = true; + return nullptr; + } + auto alloc = tile.getDefiningOp(); + if (!alloc || !alloc.getAddr()) { + op->emitError() << "VPTO declarative bridge: operand #" << abiArg.operand + << " ('" << abiArg.arg << "', role " << abiArg.role + << ") tile must come from an alloc_tile with a " + "planned address"; + hadError = true; + return nullptr; + } + return alloc.getAddr(); + }; + + // Collects the tile template token of one abi-bound operand into the + // spec under the operand's role, plus the core kind of the wrapper the + // entry routes into: a wrapper declaration may omit `core`, in which + // case the renderer picks the guard up from the reserved + // `core.` spec key collected here. + auto collectTileToken = [&](Operation *op, const BridgeAbiArg &abiArg, + Value tile, + const BridgeWhitelistEntry &entry) { + auto tileTokOr = buildBridgeTileToken(cast(tile.getType())); + if (failed(tileTokOr)) { + op->emitError() << "VPTO declarative bridge failed to build the " + << abiArg.role << " tile template token for operand '" + << abiArg.arg << "'"; + hadError = true; + return; + } + spec.addField(op, abiArg.role, *tileTokOr); + spec.addField(op, "core." + entry.wrapper, + bridgeCoreKindForTile(tile)); + if (auto alloc = tile.getDefiningOp()) { + bridgedAllocs.push_back(alloc); + } + }; + + // Renders a tmpl_map attr row into the spec. The attribute is reflected + // through the enum token interface; a missing attribute or the omit + // case renders no template argument (e.g. the Unspecified accumulation + // phase, and entries such as tmatmul.mx.bias that carry no phase). + auto collectAttrToken = [&](Operation *op, + const BridgeTmplMapField &field) { + Attribute attrValue = op->getAttr(field.field); + if (!attrValue) { + attrValue = op->getAttr(camelCaseFieldName(field.field)); + } + if (!attrValue) { + return; + } + auto enumTokenAttr = dyn_cast(attrValue); + if (!enumTokenAttr) { + op->emitError() << "VPTO declarative bridge: attribute '" + << field.field + << "' must be a PTO enum attribute to feed a " + "template slot"; + hadError = true; + return; + } + llvm::StringRef caseSymbol = enumTokenAttr.getEnumCaseSymbol(); + if (!field.omitValue.empty() && caseSymbol == field.omitValue) { + return; + } + spec.addField(op, field.field, field.enumType + "::" + caseSymbol.str()); + }; + + for (auto &[op, entry] : routed) { + // Per-op error flag: one broken op must not keep the other routed ops + // from lowering (matching the family pass continue semantics); the + // pass fails at the end when any error was recorded. + bool opFailed = false; + if (op->getNumResults() > 0) { + op->emitError("VPTO declarative bridge supports the buffer form " + "without a tensor result"); + hadError = true; + continue; + } + // Resolve the call arguments in abi order; every operand must bind to + // a planned tile address. + SmallVector callArgs; + for (const BridgeAbiArg &abiArg : entry->abi) { + if (abiArg.operand >= static_cast(op->getNumOperands())) { + op->emitError() << "VPTO declarative bridge: whitelist entry '" + << entry->entry << "' binds operand #" + << abiArg.operand << " but the op has only " + << op->getNumOperands() << " operands"; + hadError = true; + opFailed = true; + break; + } + Value tile = op->getOperand(abiArg.operand); + Value addr = resolvePlannedTile(op, abiArg, tile); + if (!addr) { + opFailed = true; + break; + } + callArgs.push_back(addr); + } + if (opFailed) { + continue; + } + // Template specialization: one tile token per abi role, then the + // enum attribute rows, then the wrapper entry name. + for (const BridgeAbiArg &abiArg : entry->abi) { + collectTileToken(op, abiArg, op->getOperand(abiArg.operand), *entry); + } + for (const BridgeTmplMapField &field : entry->tmplMap) { + if (field.source == kAttrTmplMapSource) { + collectAttrToken(op, field); + } + } + spec.addField(op, deriveEntrySpecKey(op->getName().getStringRef()), + entry->entry); + OpBuilder builder(op); + builder.create(op->getLoc(), /*results=*/TypeRange{}, + /*callee=*/entry->entry, + /*storage_size_callee=*/nullptr, + /*args=*/callArgs); + op->erase(); + } + + // Erase the tile handles consumed by the bridged ops. Handles with + // surviving users (e.g. a tile_buf_addr feeding a non-bridged op) stay + // on the regular lowering path. + for (AllocTileOp alloc : bridgedAllocs) { + if (alloc.use_empty()) { + alloc.erase(); + } + } + + if (hadError || spec.hadError()) { + signalPassFailure(); + return; + } + spec.store(func); + } +}; + +} // namespace + +std::unique_ptr createPTOLowerDeclarativeBridgeOpsPass() { + return std::make_unique(); +} + +} // namespace pto +} // namespace mlir diff --git a/lib/PTO/Transforms/PTOLowerPipeFamilyOps.cpp b/lib/PTO/Transforms/PTOLowerPipeFamilyOps.cpp new file mode 100644 index 0000000000..229a16919c --- /dev/null +++ b/lib/PTO/Transforms/PTOLowerPipeFamilyOps.cpp @@ -0,0 +1,406 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- PTOLowerPipeFamilyOps.cpp - TPipe family bridge lowering ----------===// +//===----------------------------------------------------------------------===// +// +// TPipe family pass of the VPTO C++ interface bridge. It understands the +// semantics of the internal pipe ops (initialize_l2l_pipe / tpush / tpop / +// tfree) plus the tile handles they consume (alloc_tile / declare_tile / +// tile_buf_addr) and rewrites them into generic pto.bridge_call / +// pto.bridge_inttoptr ops that carry only wrapper callee names and ABI +// values. All family semantics (config validation, storage handle flow, and +// the runtime rebinding of a declared tile to the FIFO slot returned by +// TPOP) are resolved here; the generic bridge lowering pass only sees the +// resulting bridge ops. +// +// Routing is whitelist driven: the wrapper callee of every converted op is +// looked up in the bridge whitelist by IR op name, so this pass holds no +// hardcoded wrapper entry names. Functions without pipe family ops are +// left untouched entirely (their tile handles keep flowing through the +// regular FoldTileBufIntrinsics path). +// +//===----------------------------------------------------------------------===// + +#include "PTO/IR/PTO.h" +#include "PTO/Transforms/VPTOBridgeSpecCollector.h" +#include "PTO/Transforms/VPTOBridgeTokens.h" +#include "PTO/Transforms/VPTOBridgeWhitelist.h" +#include "mlir/Dialect/Arith/IR/Arith.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/IR/Builders.h" +#include "mlir/Pass/Pass.h" +#include "llvm/ADT/DenseMap.h" +#include "llvm/ADT/SmallVector.h" + +namespace mlir { +namespace pto { + +#define GEN_PASS_DECL_PTOLOWERPIPEFAMILYOPS +#define GEN_PASS_DEF_PTOLOWERPIPEFAMILYOPS +#include "PTO/Transforms/Passes.h.inc" + +namespace { + +/// Emits a bridge call with no results and no synthesized storage. +static BridgeCallOp emitVoidBridgeCall(OpBuilder &builder, Location loc, + llvm::StringRef callee, + ValueRange args) { + return builder.create( + loc, /*results=*/TypeRange{}, /*callee=*/callee, + /*storage_size_callee=*/nullptr, /*args=*/args); +} + +/// Returns the address value a tile_buf_addr operand resolves to, or nullptr +/// when the source cannot be resolved (the caller emits the diagnostic). +/// alloc_tile carries the planned address as an i64 operand; a declare_tile +/// rebound by TPOP resolves to the FIFO slot address returned by the pop. +static Value resolveTileAddress(Value tile, OpBuilder &builder, + llvm::DenseMap &popAddresses) { + if (auto alloc = tile.getDefiningOp()) { + return alloc.getAddr(); + } + if (isa_and_nonnull(tile.getDefiningOp())) { + auto it = popAddresses.find(tile); + if (it == popAddresses.end()) { + return nullptr; + } + return it->second; + } + return nullptr; +} + +struct PTOLowerPipeFamilyOpsPass final + : public impl::PTOLowerPipeFamilyOpsBase { + MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(PTOLowerPipeFamilyOpsPass) + + void runOnOperation() override { + func::FuncOp func = getOperation(); + OpBuilder builder(func); + bool hadError = false; + // Wrapper specialization fields collected while lowering this function; + // merged into the module bridge spec attribute once lowering succeeds. + BridgeSpecCollector spec; + + // Collect first; rewriting during the walk would invalidate the walker. + SmallVector inits; + SmallVector pushes; + SmallVector pops; + SmallVector frees; + SmallVector addrs; + SmallVector allocs; + SmallVector decls; + func.walk([&](Operation *op) { + if (auto init = dyn_cast(op)) { + inits.push_back(init); + } else if (auto push = dyn_cast(op)) { + pushes.push_back(push); + } else if (auto pop = dyn_cast(op)) { + pops.push_back(pop); + } else if (auto free = dyn_cast(op)) { + frees.push_back(free); + } else if (auto addr = dyn_cast(op)) { + addrs.push_back(addr); + } else if (auto alloc = dyn_cast(op)) { + allocs.push_back(alloc); + } else if (auto decl = dyn_cast(op)) { + decls.push_back(decl); + } + }); + + // Whitelist-driven routing: the pass only acts on functions that carry + // pipe family ops. Tile handles of pipe-less functions keep flowing + // through the regular lowering (FoldTileBufIntrinsics), matching the + // pre-bridge behavior. + if (inits.empty() && pushes.empty() && pops.empty() && frees.empty()) { + return; + } + + // The whitelist always resolves through the formal chain (pass option, + // PTOAS_VPTO_BRIDGE_WHITELIST, built-in default), so routing is + // guaranteed; `whitelistName` is only for diagnostics. + std::string whitelistName; + FailureOr whitelistOr = + loadBridgeWhitelist(whitelistPath, llvm::errs(), &whitelistName); + if (failed(whitelistOr)) { + signalPassFailure(); + return; + } + const BridgeWhitelist &whitelist = *whitelistOr; + + // Resolves the whitelist entry routing `op`, or nullptr after emitting + // a diagnostic. Pipe ops have no non-bridge VPTO lowering, so a missing + // routing entry is a hard error rather than a silent fallback. + auto routeOp = [&](Operation *op) -> const BridgeWhitelistEntry * { + StringRef opName = op->getName().getStringRef(); + const BridgeWhitelistEntry *entry = whitelist.findOp(opName); + if (!entry) { + op->emitError() + << "VPTO pipe bridge: '" << opName + << "' is not routed in the bridge whitelist '" << whitelistName + << "'"; + hadError = true; + } + return entry; + }; + + // Phase 1: initialize_l2l_pipe -> storage-producing bridge init call. + // The SSA pipe value becomes the bridge call result (the storage handle); + // push/pop/free below consume that same value. + for (InitializeL2LPipeOp init : inits) { + const BridgeWhitelistEntry *entry = routeOp(init); + if (!entry) { + continue; + } + if (entry->storageSizeEntry.empty()) { + init.emitError() + << "VPTO pipe bridge: whitelist entry '" << entry->entry + << "' must declare a storage_size_entry for the stateful pipe " + "storage"; + hadError = true; + continue; + } + if (!isSupportedPipeCapability(init)) { + init.emitError( + "VPTO pipe bridge supports only a local pipe with dir_mask 1 " + "(C2V) or 2 (V2C), no acc_push_epilogue, and an i32 local buffer " + "address"); + hadError = true; + continue; + } + auto pipeTokOr = buildBridgePipeToken(init); + if (failed(pipeTokOr)) { + init.emitError("VPTO pipe bridge failed to build the TPipe template " + "token from the init attributes (flag_base is " + "required, dir_mask must be 1, 2 or 3)"); + hadError = true; + continue; + } + spec.addUniqueField(init, kBridgeSpecPipeKey, *pipeTokOr); + spec.addUniqueField(init, kBridgeSpecEntryInitKey, entry->entry); + spec.addUniqueField(init, kBridgeSpecEntrySizeKey, + entry->storageSizeEntry); + builder.setInsertionPoint(init); + BridgeCallOp call = builder.create( + init.getLoc(), /*results=*/TypeRange{init.getPipe().getType()}, + /*callee=*/entry->entry, + /*storage_size_callee=*/ + builder.getStringAttr(entry->storageSizeEntry), + /*args=*/ValueRange{init.getLocalAddr()}); + // The bridge call result becomes the storage handle: push/pop/free + // consume the same SSA value instead of the erased pipe op. + init.getPipe().replaceAllUsesWith(call.getResults().front()); + init.erase(); + } + + // Phase 2: tpop -> bridge pop call; record the returned FIFO slot + // address for the declared tile it rebinds. + llvm::DenseMap popAddresses; + // The wrapper renders one shared TileSplitAxis template argument for the + // push/pop/free entries, so every bridged op of the function must agree + // on the split value. The first bridged op fixes it; later ops check + // against it. Cross-function mismatches surface as a spec merge conflict + // in the wrapper generation pass. + std::optional bridgedSplit; + auto checkSplitConsistency = [&](Operation *op, int64_t split, + llvm::StringRef opName) { + if (bridgedSplit && *bridgedSplit != split) { + op->emitError() << "VPTO pipe bridge " << opName << " split " << split + << " does not match the split " << *bridgedSplit + << " already bridged in this function; the wrapper " + "renders one shared TileSplitAxis"; + hadError = true; + return false; + } + if (!bridgedSplit) { + auto splitTokOr = buildBridgeTileSplitToken(split); + if (failed(splitTokOr)) { + op->emitError() << "VPTO pipe bridge " << opName + << " carries an unsupported split value " << split; + hadError = true; + return false; + } + spec.addUniqueField(op, kBridgeSpecSplitKey, *splitTokOr); + bridgedSplit = split; + } + return true; + }; + for (TPopOp pop : pops) { + const BridgeWhitelistEntry *entry = routeOp(pop); + if (!entry) { + continue; + } + if (popAddresses.count(pop.getTile())) { + pop.emitError( + "VPTO pipe bridge supports at most one TPOP per declared tile; " + "sequential rebind consumption is not supported yet"); + hadError = true; + continue; + } + auto consumerTileTy = dyn_cast(pop.getTile().getType()); + if (!consumerTileTy) { + pop.emitError("VPTO pipe bridge TPOP tile must be a tile_buf"); + hadError = true; + continue; + } + auto consumerTokOr = buildBridgeTileToken(consumerTileTy); + if (failed(consumerTokOr)) { + pop.emitError("VPTO pipe bridge failed to build the consumer tile " + "template token for TPOP"); + hadError = true; + continue; + } + if (!checkSplitConsistency(pop, pop.getSplit(), "TPOP")) { + continue; + } + spec.addUniqueField(pop, kBridgeSpecConsumerTileKey, *consumerTokOr); + spec.addUniqueField(pop, kBridgeSpecEntryPopKey, entry->entry); + builder.setInsertionPoint(pop); + BridgeCallOp call = builder.create( + pop.getLoc(), /*results=*/TypeRange{builder.getI64Type()}, + /*callee=*/entry->entry, /*storage_size_callee=*/nullptr, + /*args=*/ValueRange{pop.getPipeHandle()}); + popAddresses[pop.getTile()] = call.getResults().front(); + pop.erase(); + } + + // Phase 3: tile_buf_addr -> bridge_inttoptr on the resolved address. + for (TileBufAddrOp addr : addrs) { + Value address = resolveTileAddress(addr.getSrc(), builder, popAddresses); + if (!address) { + addr.emitError( + "VPTO pipe bridge requires tile_buf_addr sources to be a planned " + "alloc_tile or a declare_tile rebound by tpop"); + hadError = true; + continue; + } + builder.setInsertionPoint(addr); + BridgeIntToPtrOp pointer = builder.create( + addr.getLoc(), addr.getDst().getType(), address); + addr.getDst().replaceAllUsesWith(pointer.getResult()); + addr.erase(); + } + + // Phase 4: tpush -> bridge push call on the planned alloc_tile address. + for (TPushOp push : pushes) { + const BridgeWhitelistEntry *entry = routeOp(push); + if (!entry) { + continue; + } + auto alloc = push.getTile().getDefiningOp(); + if (!alloc || !alloc.getAddr()) { + push.emitError("VPTO pipe bridge TPUSH requires a tile from an " + "alloc_tile with a planned address"); + hadError = true; + continue; + } + auto producerTileTy = dyn_cast(push.getTile().getType()); + if (!producerTileTy) { + push.emitError("VPTO pipe bridge TPUSH tile must be a tile_buf"); + hadError = true; + continue; + } + auto producerTokOr = buildBridgeTileToken(producerTileTy); + if (failed(producerTokOr)) { + push.emitError("VPTO pipe bridge failed to build the producer tile " + "template token for TPUSH"); + hadError = true; + continue; + } + if (!checkSplitConsistency(push, push.getSplit(), "TPUSH")) { + continue; + } + spec.addUniqueField(push, kBridgeSpecProducerTileKey, *producerTokOr); + spec.addUniqueField(push, kBridgeSpecEntryPushKey, entry->entry); + builder.setInsertionPoint(push); + emitVoidBridgeCall(builder, push.getLoc(), entry->entry, + ValueRange{push.getPipeHandle(), alloc.getAddr()}); + push.erase(); + } + + // Phase 5: tfree -> bridge free call. + for (TFreeOp free : frees) { + const BridgeWhitelistEntry *entry = routeOp(free); + if (!entry) { + continue; + } + if (free.getEntry()) { + free.emitError("VPTO pipe bridge TFREE supports the pipe-entry form " + "without a tile operand"); + hadError = true; + continue; + } + if (!checkSplitConsistency(free, free.getSplit(), "TFREE")) { + continue; + } + spec.addUniqueField(free, kBridgeSpecEntryFreeKey, entry->entry); + builder.setInsertionPoint(free); + emitVoidBridgeCall(builder, free.getLoc(), entry->entry, + ValueRange{free.getPipeHandle()}); + free.erase(); + } + + // Phase 6: erase tile handles whose consumers are all bridged now. + for (AllocTileOp alloc : allocs) { + if (!alloc.use_empty()) { + alloc.emitError("VPTO pipe bridge: alloc_tile still has users after " + "pipe family lowering"); + hadError = true; + continue; + } + alloc.erase(); + } + for (DeclareTileOp decl : decls) { + if (!decl.use_empty()) { + decl.emitError("VPTO pipe bridge: declare_tile still has users after " + "pipe family lowering"); + hadError = true; + continue; + } + decl.erase(); + } + + // Store the per-function specialization on the function itself; the + // module-level wrapper generation pass merges the per-function specs + // deterministically. The family pass instances may run concurrently, + // so they must not write the shared module attribute directly. + if (!hadError && !spec.hadError()) { + spec.store(func); + } + + if (hadError || spec.hadError()) { + signalPassFailure(); + } + } + +private: + /// Capability check for the pipe bridge. The concrete configuration + /// (slot_size/slot_num/flag_base/nosplit) is read from the op attributes + /// and flows into the generated wrapper; only genuinely unsupported forms + /// are rejected here. + static bool isSupportedPipeCapability(InitializeL2LPipeOp init) { + int8_t dirMask = init.getDirMask(); + if (dirMask != 1 && dirMask != 2) + return false; + if (init.getAccPushEpilogueAttr()) + return false; + auto localAddrTy = dyn_cast(init.getLocalAddr().getType()); + if (!localAddrTy || localAddrTy.getWidth() != 32) + return false; + return true; + } +}; + +} // namespace + +std::unique_ptr createPTOLowerPipeFamilyOpsPass() { + return std::make_unique(); +} + +} // namespace pto +} // namespace mlir diff --git a/lib/PTO/Transforms/PTOToEmitC.cpp b/lib/PTO/Transforms/PTOToEmitC.cpp index 1a35699d48..8e2de69026 100644 --- a/lib/PTO/Transforms/PTOToEmitC.cpp +++ b/lib/PTO/Transforms/PTOToEmitC.cpp @@ -19,6 +19,7 @@ #include "PTO/IR/PTOTypeUtils.h" #include "PTO/IR/PTOSyncUtils.h" #include "PTO/Transforms/MemoryConsistencyAttrs.h" +#include "PTO/Transforms/PTOCppTokens.h" #include "PTO/Transforms/Passes.h" #include "Utils.h" @@ -426,54 +427,9 @@ static bool isF8E8M0ElemType(Type elemTy) { } static std::string getEmitCScalarTypeToken(Type elemTy) { - if (pto::isPTOFloat8E4M3LikeType(elemTy)) { - return "float8_e4m3_t"; - } - if (pto::isPTOFloat8E5M2LikeType(elemTy)) { - return "float8_e5m2_t"; - } - if (isF8E8M0ElemType(elemTy)) { - return "float8_e8m0_t"; - } - if (isa(elemTy)) { - return "hifloat8_t"; - } - if (isa(elemTy)) { - return "float4_e1m2x2_t"; - } - if (isa(elemTy)) { - return "float4_e2m1x2_t"; - } - if (elemTy.isF16()) { - return "half"; - } - if (elemTy.isBF16()) { - return "bfloat16_t"; - } - if (elemTy.isF32()) { - return "float"; - } - if (elemTy.isF64()) { - return "double"; - } - if (elemTy.isInteger(8)) { - return (elemTy.isSignlessInteger(8) || elemTy.isSignedInteger(8)) ? "int8_t" - : "uint8_t"; - } - if (elemTy.isInteger(16)) { - return (elemTy.isSignlessInteger(16) || elemTy.isSignedInteger(16)) - ? "int16_t" - : "uint16_t"; - } - if (elemTy.isInteger(32)) { - return (elemTy.isSignlessInteger(32) || elemTy.isSignedInteger(32)) - ? "int32_t" - : "uint32_t"; - } - if (elemTy.isInteger(64)) { - return cast(elemTy).isUnsigned() ? "uint64_t" : "int64_t"; - } - return "float"; + // The scalar spelling table is shared with the VPTO bridge wrapper + // generator (PTOCppTokens); the EmitC output context needs no qualifier. + return pto::getPTOCppElementTypeToken(elemTy); } static emitc::PointerType getEmitCPointerType(MLIRContext *ctx, @@ -727,31 +683,22 @@ static const char *scalingRoleToken(Type elemTy, return "TileType::Scaling"; } -static const char *tileRoleToken(Attribute memorySpace, +static std::string tileRoleToken(Attribute memorySpace, std::optional elemType = std::nullopt, std::optional configAttr = std::nullopt) { if (auto asAttr = dyn_cast_or_null(memorySpace)) { - switch (asAttr.getAddressSpace()) { - case pto::AddressSpace::VEC: - return "TileType::Vec"; - case pto::AddressSpace::MAT: - return "TileType::Mat"; - case pto::AddressSpace::LEFT: - return "TileType::Left"; - case pto::AddressSpace::RIGHT: - return "TileType::Right"; - case pto::AddressSpace::ACC: - return "TileType::Acc"; - case pto::AddressSpace::BIAS: - return "TileType::Bias"; - case pto::AddressSpace::SCALING: - if (elemType && configAttr) - return scalingRoleToken(*elemType, *configAttr); - return "TileType::Scaling"; - case pto::AddressSpace::GM: - case pto::AddressSpace::Zero: + auto as = asAttr.getAddressSpace(); + // The MX scale tiles refine the SCALING role by layout before the + // shared mapping is consulted. + if (as == pto::AddressSpace::SCALING && elemType && configAttr) + return scalingRoleToken(*elemType, *configAttr); + // Global-memory and default address spaces fall back to the vector + // role on the EmitC path; the shared mapping rejects them. + if (as == pto::AddressSpace::GM || as == pto::AddressSpace::Zero) return "TileType::Vec"; - } + if (auto tok = pto::getPTOCppTileTypeToken(as, /*qualifier=*/""); + succeeded(tok)) + return *tok; } return "TileType::Vec"; } @@ -1122,49 +1069,27 @@ static Value castSignlessIntToUnsignedSameWidth(ConversionPatternRewriter &rewri static bool needsA5NoSplitVectorGuard(Operation *op); static FailureOr getTileSplitToken(int64_t split) { - switch (split) { - case 0: - return std::string("TileSplitAxis::TILE_NO_SPLIT"); - case 1: - return std::string("TileSplitAxis::TILE_UP_DOWN"); - case 2: - return std::string("TileSplitAxis::TILE_LEFT_RIGHT"); - case 3: - return std::string("TileSplitAxis::TILE_UP_DOWN_ODD"); - case 4: - return std::string("TileSplitAxis::TILE_LEFT_RIGHT_ODD"); - default: - return failure(); - } + return pto::getPTOCppTileSplitToken(split, /*qualifier=*/""); } static FailureOr getTPipeDirectionToken(bool isL2G2L, int8_t dirMask, PTOArch targetArch) { - if (dirMask == 1) { - if (isL2G2L && targetArch == PTOArch::A5) + // The A5 "_GM" variants only apply to L2G2L pipes; the local pipe + // directions come from the shared mapping. + if (isL2G2L && targetArch == PTOArch::A5) { + if (dirMask == 1) return std::string("Direction::DIR_C2V_GM"); - return std::string("Direction::DIR_C2V"); - } - if (dirMask == 2) { - if (isL2G2L && targetArch == PTOArch::A5) + if (dirMask == 2) return std::string("Direction::DIR_V2C_GM"); - return std::string("Direction::DIR_V2C"); } - if (dirMask == 3) - return std::string("Direction::DIR_BOTH"); - return failure(); + return pto::getPTOCppDirectionToken(dirMask, /*qualifier=*/""); } static std::string buildTPipeToken(int32_t flagBase, llvm::StringRef dirTok, int32_t slotSize, int32_t slotNum, int32_t localSlotNum, bool nosplit) { - std::string token = "TPipe<" + std::to_string(flagBase) + ", " + dirTok.str() + - ", " + std::to_string(slotSize) + ", " + - std::to_string(slotNum); - token += ", " + std::to_string(localSlotNum); - token += nosplit ? ", true" : ", false"; - token += ">"; - return token; + return pto::renderTPipeSpelling(flagBase, dirTok, slotSize, slotNum, + localSlotNum, nosplit, /*qualifier=*/""); } static FailureOr buildTPipeTokenFromInitOp(Operation *op, @@ -4552,23 +4477,17 @@ static Value materializeGlobalTensorDataPointer( } static std::string tileBufBLayoutToken(pto::TileBufConfigAttr configAttr) { - std::string blTok = "BLayout::RowMajor"; - if (auto blAttr = dyn_cast(configAttr.getBLayout())) { - if (static_cast(blAttr.getValue()) == 1) - blTok = "BLayout::ColMajor"; - } - return blTok; + auto tok = pto::getPTOCppBLayoutToken(getTileBufBLayoutValue(configAttr), + /*qualifier=*/""); + assert(succeeded(tok) && "closed BLayout enum set"); + return *tok; } static std::string tileBufSLayoutToken(pto::TileBufConfigAttr configAttr) { - std::string slTok = "SLayout::NoneBox"; - if (auto slAttr = dyn_cast(configAttr.getSLayout())) { - int32_t slVal = static_cast(slAttr.getValue()); - slTok = (slVal == 1) ? "SLayout::RowMajor" - : (slVal == 2) ? "SLayout::ColMajor" - : "SLayout::NoneBox"; - } - return slTok; + auto tok = pto::getPTOCppSLayoutToken(getTileBufSLayoutValue(configAttr), + /*qualifier=*/""); + assert(succeeded(tok) && "closed SLayout enum set"); + return *tok; } static std::string tileBufPadToken(pto::TileBufConfigAttr configAttr) { diff --git a/lib/PTO/Transforms/VPTOBridgeLowering.cpp b/lib/PTO/Transforms/VPTOBridgeLowering.cpp new file mode 100644 index 0000000000..e1faecda0a --- /dev/null +++ b/lib/PTO/Transforms/VPTOBridgeLowering.cpp @@ -0,0 +1,328 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- VPTOBridgeLowering.cpp - generic C++ interface bridge lowering ----===// +//===----------------------------------------------------------------------===// +// +// Generic bridge lowering pass of the VPTO C++ interface bridge. It knows +// nothing about individual PTO-ISA interface families: it validates each +// pto.bridge_call against the bridge whitelist and mechanically lowers it +// into a call to the wrapper entry, materializing the wrapper declaration +// at module level. Entries that carry `storage_size_callee` additionally +// synthesize the stateful-object pattern (size query + stack storage) so +// family passes can express "construct a template object on the kernel +// stack" without emitting LLVM dialect ops themselves. +// +// The whitelist is also the routing check of last resort: any op still +// present in the IR that the whitelist routes to a wrapper entry was +// missed by its family pass, and is rejected with a diagnostic instead of +// silently flowing into the regular LLVM emission path. +// +//===----------------------------------------------------------------------===// + +#include "PTO/IR/PTO.h" +#include "PTO/Transforms/VPTOBridgeWhitelist.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/IR/Builders.h" +#include "mlir/Pass/Pass.h" +#include "mlir/Transforms/DialectConversion.h" +#include "llvm/ADT/STLExtras.h" +#include "llvm/ADT/StringMap.h" +#include "llvm/ADT/StringSet.h" + +namespace mlir { +namespace pto { + +#define GEN_PASS_DECL_VPTOBRIDGELOWERING +#define GEN_PASS_DEF_VPTOBRIDGELOWERING +#include "PTO/Transforms/Passes.h.inc" + +namespace { + +/// Converts the carrier types a bridge op may hold. These rules mirror the +/// PipeType/PtrType entries of the VPTO type converter +/// (VPTOCANN900LLVMEmitter.cpp convertVPTOType); the bridge lowering runs +/// before that converter and must agree with it so values flow into the +/// remaining PTO ops without extra casts. +class BridgeTypeConverter final : public TypeConverter { +public: + explicit BridgeTypeConverter(MLIRContext *context) { + addConversion([](Type type) -> Type { + if (isa(type)) { + return LLVM::LLVMPointerType::get(type.getContext()); + } + if (auto ptrType = dyn_cast(type)) { + return LLVM::LLVMPointerType::get( + type.getContext(), + static_cast(ptrType.getMemorySpace().getAddressSpace())); + } + return type; + }); + addSourceMaterialization(materializeBridgeCast); + addTargetMaterialization(materializeBridgeCast); + } + +private: + static std::optional materializeBridgeCast(OpBuilder &builder, + Type resultType, + ValueRange inputs, + Location loc) { + if (inputs.size() != 1) { + return std::nullopt; + } + return builder + .create(loc, TypeRange{resultType}, inputs) + .getResult(0); + } +}; + +struct BridgeLoweringState { + const BridgeWhitelist &whitelist; + llvm::StringSet<> declaredEntries; +}; + +/// Creates the module-level private declaration of a wrapper entry the +/// first time it is called. +static void ensureWrapperDecl(ModuleOp module, BridgeLoweringState &state, + PatternRewriter &rewriter, StringRef callee, + TypeRange argTypes, TypeRange resultTypes) { + if (state.declaredEntries.contains(callee)) { + return; + } + state.declaredEntries.insert(callee); + OpBuilder::InsertionGuard guard(rewriter); + rewriter.setInsertionPointToStart(&module.getBodyRegion().front()); + auto decl = rewriter.create( + module.getLoc(), callee, + FunctionType::get(module.getContext(), argTypes, resultTypes)); + decl.setPrivate(); +} + +/// Validates the fully assembled call argument list against the whitelist +/// ABI. Emits a diagnostic and returns failure on any mismatch. +static LogicalResult validateAbi(Operation *op, const BridgeWhitelistEntry &entry, + ValueRange callArgs) { + if (callArgs.size() != entry.abi.size()) { + return op->emitError() + << "VPTO bridge call to '" << entry.entry << "' passes " + << callArgs.size() << " argument(s), whitelist ABI declares " + << entry.abi.size(); + } + for (auto [index, arg] : llvm::enumerate(callArgs)) { + const BridgeAbiArg &abiArg = entry.abi[index]; + if (!bridgeAbiTypeMatches(abiArg.type, arg.getType())) { + return op->emitError() + << "VPTO bridge call to '" << entry.entry << "' argument #" + << index << " has type " << arg.getType() + << ", whitelist ABI declares '" << abiArg.type << "'"; + } + } + return success(); +} + +class LowerBridgeCallPattern final : public OpConversionPattern { +public: + LowerBridgeCallPattern(TypeConverter &converter, MLIRContext *context, + BridgeLoweringState &state) + : OpConversionPattern(converter, context), state(state) {} + + LogicalResult + matchAndRewrite(BridgeCallOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + Location loc = op.getLoc(); + StringRef callee = op.getCalleeAttr().getValue(); + const BridgeWhitelistEntry *entry = state.whitelist.findEntry(callee); + if (!entry) { + return op.emitError() + << "VPTO bridge call to wrapper entry '" << callee + << "' is not declared in the bridge whitelist"; + } + ModuleOp module = op->getParentOfType(); + ValueRange operands = adaptor.getArgs(); + + // Stateful entries synthesize their own storage: query the wrapper for + // the object size and alloca it on the kernel stack. The storage value + // replaces the bridge call result, which must be the only result. + bool hasStorage = op.getStorageSizeCalleeAttr() != nullptr; + SmallVector callArgs; + Value storage; + if (hasStorage) { + if (op.getNumResults() != 1) { + return op.emitError() + << "VPTO bridge call with storage_size_callee must have " + "exactly one result (the storage handle)"; + } + StringRef sizeCallee = op.getStorageSizeCalleeAttr().getValue(); + if (!state.whitelist.findEntry(sizeCallee)) { + return op.emitError() + << "VPTO bridge storage size callee '" << sizeCallee + << "' is not declared in the bridge whitelist"; + } + Value size = rewriter.create(loc, sizeCallee, + rewriter.getI64Type(), + ValueRange{}) + .getResult(0); + ensureWrapperDecl(module, state, rewriter, sizeCallee, /*argTypes=*/{}, + /*resultTypes=*/{rewriter.getI64Type()}); + storage = rewriter.create( + loc, LLVM::LLVMPointerType::get(rewriter.getContext()), + rewriter.getI8Type(), size, /*alignment=*/8); + callArgs.push_back(storage); + } + callArgs.append(operands.begin(), operands.end()); + + if (failed(validateAbi(op, *entry, callArgs))) { + return failure(); + } + + SmallVector resultTypes; + for (Type resultType : op.getResultTypes()) { + Type converted = getTypeConverter()->convertType(resultType); + if (!converted) { + return op.emitError() + << "VPTO bridge call result type " << resultType + << " has no bridge conversion"; + } + resultTypes.push_back(converted); + } + + func::CallOp call = rewriter.create( + loc, callee, TypeRange(resultTypes), ValueRange(callArgs)); + ensureWrapperDecl(module, state, rewriter, callee, + llvm::map_to_vector<4>(callArgs, + [](Value arg) { return arg.getType(); }), + TypeRange(resultTypes)); + + if (hasStorage) { + rewriter.replaceOp(op, storage); + return success(); + } + if (call.getNumResults() == 0) { + rewriter.eraseOp(op); + return success(); + } + rewriter.replaceOp(op, call.getResults()); + return success(); + } + +private: + BridgeLoweringState &state; +}; + +class LowerBridgeIntToPtrPattern final + : public OpConversionPattern { +public: + LowerBridgeIntToPtrPattern(TypeConverter &converter, MLIRContext *context) + : OpConversionPattern(converter, context) {} + + LogicalResult + matchAndRewrite(BridgeIntToPtrOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + Type convertedResult = getTypeConverter()->convertType(op.getResult().getType()); + if (!convertedResult || !isa(convertedResult)) { + return op.emitError() + << "VPTO bridge inttoptr requires a result type that converts " + "to an LLVM pointer, got " + << op.getResult().getType(); + } + rewriter.replaceOpWithNewOp(op, convertedResult, + adaptor.getAddr()); + return success(); + } +}; + +struct VPTOBridgeLoweringPass final + : public impl::VPTOBridgeLoweringBase { + MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(VPTOBridgeLoweringPass) + + void runOnOperation() override { + ModuleOp module = getOperation(); + bool hasBridgeOps = false; + module.walk([&](Operation *op) { + if (isa(op)) { + hasBridgeOps = true; + } + }); + + // The whitelist always resolves through the formal chain (pass option, + // PTOAS_VPTO_BRIDGE_WHITELIST, built-in default), so this pass always + // validates; `whitelistName` is only for diagnostics. + std::string whitelistName; + FailureOr whitelistOr = + loadBridgeWhitelist(whitelistPath, llvm::errs(), &whitelistName); + if (failed(whitelistOr)) { + signalPassFailure(); + return; + } + BridgeWhitelist whitelist = std::move(*whitelistOr); + + // Routing check: an op the whitelist routes to a wrapper entry must + // have been rewritten into bridge ops by the pass owning its lowering + // channel. Leftovers mean that pass was skipped or missed the op; + // reject them here instead of letting them flow into the regular + // emission path. The diagnostic names the channel so the reader knows + // which pass to look at. + llvm::StringMap routedOps; + for (const BridgeWhitelistEntry &entry : whitelist.bridgeOps) { + if (entry.op != BridgeWhitelist::kInternalOp) { + routedOps[entry.op] = &entry; + } + } + bool leftoversFound = false; + module.walk([&](Operation *op) { + auto it = routedOps.find(op->getName().getStringRef()); + if (it == routedOps.end()) { + return; + } + op->emitError() + << "VPTO bridge: '" << it->first() + << "' is routed to wrapper entry '" << it->second->entry + << "' by the bridge whitelist '" << whitelistName + << "' but was not lowered into a pto.bridge_call by " + << (it->second->isDeclarative() + ? "the declarative bridge lowering" + : "its family pass"); + leftoversFound = true; + }); + if (leftoversFound) { + signalPassFailure(); + return; + } + + if (!hasBridgeOps) { + return; + } + + BridgeTypeConverter converter(&getContext()); + ConversionTarget target(getContext()); + target.addIllegalOp(); + // Everything the patterns create (func.call, llvm.alloca, private + // declarations) must be legal on the target, otherwise the conversion + // driver rejects the generated operations and rolls the pattern back. + target.markUnknownOpDynamicallyLegal( + [](Operation *op) { return true; }); + + RewritePatternSet patterns(&getContext()); + BridgeLoweringState state{whitelist}; + patterns.add(converter, &getContext(), state); + patterns.add(converter, &getContext()); + if (failed(applyPartialConversion(module, target, std::move(patterns)))) { + signalPassFailure(); + } + } +}; + +} // namespace + +std::unique_ptr createVPTOBridgeLoweringPass() { + return std::make_unique(); +} + +} // namespace pto +} // namespace mlir diff --git a/lib/PTO/Transforms/VPTOBridgeTokens.cpp b/lib/PTO/Transforms/VPTOBridgeTokens.cpp new file mode 100644 index 0000000000..25891364bf --- /dev/null +++ b/lib/PTO/Transforms/VPTOBridgeTokens.cpp @@ -0,0 +1,106 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- VPTOBridgeTokens.cpp - C++ template token building ----------------===// +//===----------------------------------------------------------------------===// +// +// Implementation of the bridge-side PTO-ISA C++ template token builders. See +// include/PTO/Transforms/VPTOBridgeTokens.h. The IR-fact -> C++ spelling +// mapping rules are shared with the EmitC backend through PTOCppTokens; +// this file holds the bridge assembly rules (fully qualified spellings and +// the NoneBox trailing-argument omission). +// +//===----------------------------------------------------------------------===// + +#include "PTO/Transforms/VPTOBridgeTokens.h" +#include "PTO/IR/PTO.h" +#include "PTO/IR/PTOTypeUtils.h" +#include "PTO/Transforms/PTOCppTokens.h" +#include "llvm/ADT/Twine.h" +#include + +using namespace mlir; +using namespace mlir::pto; + +namespace { + +/// The bridge wrapper is a standalone translation unit, so every pto-isa +/// spelling is emitted fully qualified. +constexpr llvm::StringLiteral kBridgeQualifier = "pto::"; + +} // namespace + +std::string pto::buildBridgeElementTypeToken(Type elementType) { + return getPTOCppElementTypeToken(elementType); +} + +FailureOr pto::buildBridgePipeToken(InitializeL2LPipeOp init) { + IntegerAttr flagBaseAttr = init.getFlagBaseAttr(); + if (!flagBaseAttr) + return failure(); + auto dirTok = getPTOCppDirectionToken(init.getDirMask(), kBridgeQualifier); + if (failed(dirTok)) + return failure(); + + // The local-to-local pipe always uses a localSlotNum of 2 (see EmitC's + // buildTPipeTokenFromInitOp for the InitializeL2LPipeOp case). + constexpr int32_t localSlotNum = 2; + bool nosplit = init.getNosplitAttr() && init.getNosplitAttr().getValue(); + + return renderTPipeSpelling( + static_cast(flagBaseAttr.getInt()), *dirTok, + init.getSlotSize(), init.getSlotNum(), localSlotNum, nosplit, + kBridgeQualifier); +} + +FailureOr pto::buildBridgeTileSplitToken(int64_t split) { + return getPTOCppTileSplitToken(split, kBridgeQualifier); +} + +FailureOr pto::buildBridgeTileToken(TileBufType tile) { + auto addressSpaceAttr = + dyn_cast_or_null(tile.getMemorySpace()); + if (!addressSpaceAttr) + return failure(); + auto tileTypeTok = getPTOCppTileTypeToken(addressSpaceAttr.getAddressSpace(), + kBridgeQualifier); + if (failed(tileTypeTok)) + return failure(); + + ArrayRef shape = tile.getShape(); + ArrayRef validShape = tile.getValidShape(); + if (shape.size() != 2 || validShape.size() != 2) + return failure(); + + auto bLayoutTok = getPTOCppBLayoutToken( + static_cast(tile.getBLayoutValueI32()), kBridgeQualifier); + if (failed(bLayoutTok)) + return failure(); + + std::string token = + "pto::Tile<" + *tileTypeTok + ", " + + buildBridgeElementTypeToken(tile.getElementType()) + ", " + + std::to_string(shape[0]) + ", " + std::to_string(shape[1]) + ", " + + *bLayoutTok + ", " + std::to_string(validShape[0]) + ", " + + std::to_string(validShape[1]); + + // Boxed storage layouts carry the inner-fractal template arguments; the + // default NoneBox layout relies on the Tile template defaults, matching the + // hand-written wrapper specializations. + int32_t sLayoutValue = tile.getSLayoutValueI32(); + if (sLayoutValue != 0) { + auto sLayoutTok = getPTOCppSLayoutToken(static_cast(sLayoutValue), + kBridgeQualifier); + if (failed(sLayoutTok)) + return failure(); + token += ", " + *sLayoutTok + ", " + + std::to_string(tile.getSFractalSizeI32()); + } + token += ">"; + return token; +} diff --git a/lib/PTO/Transforms/VPTOBridgeWhitelist.cpp b/lib/PTO/Transforms/VPTOBridgeWhitelist.cpp new file mode 100644 index 0000000000..6fee5715e7 --- /dev/null +++ b/lib/PTO/Transforms/VPTOBridgeWhitelist.cpp @@ -0,0 +1,537 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- VPTOBridgeWhitelist.cpp - C++ bridge whitelist ---------------------===// +//===----------------------------------------------------------------------===// +// +// YAML parsing and semantic validation for the VPTO C++ interface bridge +// whitelist (see include/PTO/Transforms/VPTOBridgeWhitelist.h). +// +//===----------------------------------------------------------------------===// + +#include "PTO/Transforms/VPTOBridgeWhitelist.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "llvm/ADT/DenseSet.h" +#include "llvm/ADT/StringSet.h" +#include "llvm/Support/MemoryBuffer.h" +#include "llvm/Support/YAMLTraits.h" +#include +#include + +using namespace mlir; +using namespace mlir::pto; + +namespace llvm { +namespace yaml { + +template <> struct MappingTraits { + static void mapping(IO &io, BridgeAbiArg &arg) { + io.mapOptional("type", arg.type); + io.mapOptional("operand", arg.operand, (int64_t)-1); + io.mapOptional("arg", arg.arg); + io.mapOptional("role", arg.role); + } +}; + +template <> struct MappingTraits { + static void mapping(IO &io, BridgeTmplMapField &field) { + io.mapRequired("source", field.source); + io.mapRequired("field", field.field); + io.mapRequired("target", field.target); + io.mapOptional("enum_type", field.enumType); + io.mapOptional("omit_value", field.omitValue); + } +}; + +template <> struct MappingTraits { + static void mapping(IO &io, BridgeWhitelistEntry &entry) { + io.mapRequired("op", entry.op); + io.mapRequired("wrapper", entry.wrapper); + io.mapOptional("lowering", entry.lowering, + std::string(BridgeWhitelistEntry::kLoweringDeclarative)); + io.mapOptional("entry", entry.entry); + io.mapOptional("call", entry.call); + io.mapOptional("tmpl_args", entry.tmplArgs); + io.mapOptional("abi", entry.abi); + io.mapOptional("storage_size_entry", entry.storageSizeEntry); + io.mapOptional("tmpl_map", entry.tmplMap); + } +}; + +template <> struct MappingTraits { + static void mapping(IO &io, BridgeWrapperDecl &decl) { + io.mapRequired("name", decl.name); + io.mapRequired("includes", decl.includes); + io.mapOptional("core", decl.core); + } +}; + +template <> struct MappingTraits { + static void mapping(IO &io, BridgeWhitelist &whitelist) { + io.mapOptional("wrappers", whitelist.wrappers); + io.mapRequired("bridge_ops", whitelist.bridgeOps); + } +}; + +} // namespace yaml +} // namespace llvm + +LLVM_YAML_IS_SEQUENCE_VECTOR(BridgeAbiArg) +LLVM_YAML_IS_SEQUENCE_VECTOR(BridgeTmplMapField) +LLVM_YAML_IS_SEQUENCE_VECTOR(BridgeWhitelistEntry) +LLVM_YAML_IS_SEQUENCE_VECTOR(BridgeWrapperDecl) + +namespace { + +/// tmpl_map `source` tokens accepted for the pipe wrapper. A source names +/// the IR producer of a template argument: the pipe init op attributes or a +/// tile operand's type. +bool isPipeTmplMapSource(StringRef source) { + return source == kPipeInitTmplMapSource || source == kTileTmplMapSource; +} + +} // namespace + +bool pto::isSupportedBridgeAbiType(llvm::StringRef token) { + return token == "ptr" || token == "i64" || token == "i32"; +} + +bool pto::bridgeAbiTypeMatches(llvm::StringRef token, Type type) { + if (token == "ptr") { + return isa(type); + } + if (token == "i64") { + return type.isInteger(64); + } + if (token == "i32") { + return type.isInteger(32); + } + return false; +} + +std::string pto::deriveDefaultBridgeEntry(llvm::StringRef opName) { + constexpr llvm::StringLiteral kTileWorldOpPrefix = "pto.t"; + if (!opName.consume_front(kTileWorldOpPrefix)) { + opName.consume_front("pto."); + } + std::string name = ("pto_vpto_" + opName).str(); + constexpr llvm::StringLiteral kNamePrefix = "pto_vpto_"; + std::replace(name.begin() + kNamePrefix.size(), name.end(), '.', '_'); + return name; +} + +std::string pto::deriveDefaultBridgeCall(llvm::StringRef opName) { + // Unlike the entry name, the interface call keeps the tile-world `t` + // mnemonic lead (pto::TADD, pto::TMATMUL), so only the dialect prefix + // is stripped. + opName.consume_front("pto."); + if (opName.empty()) { + return {}; + } + std::string name; + name.reserve(opName.size()); + for (char c : opName) { + name.push_back(c == '.' ? '_' + : static_cast(llvm::toUpper(c))); + } + return "pto::" + name; +} + +std::string pto::bridgeRoleParamName(llvm::StringRef role) { + std::string name; + name.reserve(role.size()); + bool upperNext = false; + for (char c : role) { + if (c == '_') { + upperNext = true; + continue; + } + name.push_back(upperNext && c >= 'a' && c <= 'z' + ? static_cast(c - 'a' + 'A') + : c); + upperNext = false; + } + return name; +} + +std::string pto::bridgeRoleTypedefTarget(llvm::StringRef role) { + std::string target; + target.reserve(role.size()); + bool upperNext = true; + for (char c : role) { + if (c == '_') { + upperNext = true; + continue; + } + target.push_back(upperNext && c >= 'a' && c <= 'z' + ? static_cast(c - 'a' + 'A') + : c); + upperNext = false; + } + return target; +} + +FailureOr +pto::parseBridgeWhitelistFromBuffer(llvm::StringRef content, + llvm::StringRef sourceName, + llvm::raw_ostream &diagOS) { + BridgeWhitelist whitelist; + llvm::yaml::Input input(content); + input >> whitelist; + if (std::error_code error = input.error()) { + diagOS << "VPTO bridge whitelist: cannot parse '" << sourceName + << "': " << error.message() << "\n"; + return failure(); + } + + llvm::StringSet<> seenEntries; + llvm::StringSet<> seenOps; + for (BridgeWhitelistEntry &entry : whitelist.bridgeOps) { + if (entry.op.empty() || entry.wrapper.empty()) { + diagOS << "VPTO bridge whitelist: entry with op='" << entry.op + << "', wrapper='" << entry.wrapper << "', entry='" << entry.entry + << "' has an empty required field in '" << sourceName << "'\n"; + return failure(); + } + if (entry.lowering != BridgeWhitelistEntry::kLoweringDeclarative && + entry.lowering != BridgeWhitelistEntry::kLoweringCustom) { + diagOS << "VPTO bridge whitelist: entry '" << entry.entry + << "' declares unsupported lowering '" << entry.lowering + << "' in '" << sourceName << "' (supported: declarative, " + "custom)\n"; + return failure(); + } + // Routed declarative entries default their entry name from the op + // name; custom entries and wrapper-internal helpers are named by hand + // because no mechanical rule covers them. + const bool declarativeChannel = + entry.isDeclarative() && entry.op != BridgeWhitelist::kInternalOp; + if (entry.entry.empty()) { + if (declarativeChannel) { + entry.entry = deriveDefaultBridgeEntry(entry.op); + } else { + diagOS << "VPTO bridge whitelist: entry with op='" << entry.op + << "' declares no entry name in '" << sourceName + << "' (only declarative entries default it from the op " + "name)\n"; + return failure(); + } + } + if (!seenEntries.insert(entry.entry).second) { + diagOS << "VPTO bridge whitelist: duplicate wrapper entry '" + << entry.entry << "' in '" << sourceName << "'\n"; + return failure(); + } + if (entry.op != BridgeWhitelist::kInternalOp && + !seenOps.insert(entry.op).second) { + diagOS << "VPTO bridge whitelist: duplicate routed op '" << entry.op + << "' in '" << sourceName << "'\n"; + return failure(); + } + // The call spelling and template arguments belong to the generic + // declarative renderer; entries owned by a dedicated pass (or a + // wrapper-internal helper) must not carry them. + if (!declarativeChannel && (!entry.call.empty() || !entry.tmplArgs.empty())) { + diagOS << "VPTO bridge whitelist: entry '" << entry.entry + << "' declares call/tmpl_args but is not lowered through the " + "declarative channel in '" + << sourceName << "'\n"; + return failure(); + } + if (declarativeChannel) { + // The call spelling follows the op-name convention unless declared; + // a derivation the interface does not follow is overridden with an + // explicit `call`, and a wrong one fails loudly when the generated + // wrapper source is compiled. + if (entry.call.empty()) { + entry.call = deriveDefaultBridgeCall(entry.op); + } + if (entry.call.empty()) { + diagOS << "VPTO bridge whitelist: declarative entry '" << entry.entry + << "' declares no call spelling in '" << sourceName + << "' (the generic renderer needs the C++ call the entry " + "body emits)\n"; + return failure(); + } + for (const std::string &tmplArg : entry.tmplArgs) { + if (tmplArg.empty()) { + diagOS << "VPTO bridge whitelist: declarative entry '" + << entry.entry << "' has an empty tmpl_args item in '" + << sourceName << "'\n"; + return failure(); + } + // A qualified spelling is a literal template argument; anything + // else must name an attr tmpl_map row of this entry whose + // collected spec token feeds the slot. + if (llvm::StringRef(tmplArg).contains("::")) { + continue; + } + bool attrFieldDeclared = false; + for (const BridgeTmplMapField &field : entry.tmplMap) { + if (field.source == kAttrTmplMapSource && field.field == tmplArg) { + attrFieldDeclared = true; + break; + } + } + if (!attrFieldDeclared) { + diagOS << "VPTO bridge whitelist: declarative entry '" + << entry.entry << "' tmpl_args item '" << tmplArg + << "' is neither a qualified literal nor an attr " + "tmpl_map field of the entry in '" + << sourceName << "'\n"; + return failure(); + } + } + } + // Declarative entries bind every abi argument to an IR operand position + // and a template role; the role set names the entry's tile typedefs. + // Wrapper-internal helpers are never routed from an IR op, so they + // carry no operand bindings and the `lowering` value is meaningless + // for them. + llvm::StringSet<> declarativeRoles; + llvm::DenseSet declarativeOperands; + if (declarativeChannel) { + for (BridgeAbiArg &arg : entry.abi) { + if (arg.operand < 0 || arg.role.empty()) { + // The declarative channel is the default, so the likeliest cause + // is an entry that needs a dedicated pass but never opted out. + diagOS << "VPTO bridge whitelist: declarative entry '" << entry.entry + << "' has an abi argument without operand/role binding " + "in '" + << sourceName + << "' (entries owned by a dedicated family pass must " + "declare 'lowering: custom')\n"; + return failure(); + } + // The parameter name is a rendering concern: it defaults to the + // lowerCamelCase of the role, tying the wrapper parameter to its + // tile typedef. + if (arg.arg.empty()) { + arg.arg = bridgeRoleParamName(arg.role); + } + if (!declarativeOperands.insert(arg.operand).second) { + diagOS << "VPTO bridge whitelist: declarative entry '" << entry.entry + << "' binds operand #" << arg.operand + << " more than once in '" << sourceName << "'\n"; + return failure(); + } + declarativeRoles.insert(arg.role); + // Tile addresses are the only carrier the declarative channel + // emits, so the type defaults to i64 when omitted. + if (arg.type.empty()) { + arg.type = "i64"; + } + } + } + for (const BridgeAbiArg &arg : entry.abi) { + if (!isSupportedBridgeAbiType(arg.type)) { + diagOS << "VPTO bridge whitelist: unsupported ABI type token '" + << arg.type << "' for entry '" << entry.entry << "' in '" + << sourceName << "' (supported: ptr, i64, i32)\n"; + return failure(); + } + } + for (const BridgeTmplMapField &field : entry.tmplMap) { + if (field.source.empty() || field.field.empty() || + field.target.empty()) { + diagOS << "VPTO bridge whitelist: tmpl_map row of entry '" + << entry.entry << "' has an empty source/field/target in '" + << sourceName << "'\n"; + return failure(); + } + if (declarativeChannel) { + // Declarative tile typedefs derive from the abi roles, so the only + // tmpl_map rows the channel accepts map enum attributes; a tile + // row here is the legacy spelling and would name a typedef target + // the role-driven renderer never emits. + if (field.source != kAttrTmplMapSource) { + diagOS << "VPTO bridge whitelist: tmpl_map row of entry '" + << entry.entry << "' uses source '" << field.source + << "', but declarative entries only accept 'attr' rows " + "(tile typedefs derive from the abi roles) in '" + << sourceName << "'\n"; + return failure(); + } + if (field.enumType.empty()) { + diagOS << "VPTO bridge whitelist: tmpl_map attr row of entry '" + << entry.entry << "' lacks enum_type in '" << sourceName + << "'\n"; + return failure(); + } + } else if (entry.wrapper == "pipe" && + !isPipeTmplMapSource(field.source)) { + diagOS << "VPTO bridge whitelist: tmpl_map row of entry '" + << entry.entry << "' uses unknown pipe-wrapper source '" + << field.source << "' in '" << sourceName + << "' (supported: pipe.init, tile)\n"; + return failure(); + } + } + } + for (const BridgeWhitelistEntry &entry : whitelist.bridgeOps) { + if (!entry.storageSizeEntry.empty() && + !whitelist.findEntry(entry.storageSizeEntry)) { + diagOS << "VPTO bridge whitelist: entry '" << entry.entry + << "' declares storage_size_entry '" << entry.storageSizeEntry + << "' which is not a declared wrapper entry in '" << sourceName + << "'\n"; + return failure(); + } + } + // Wrapper declarations feed the generic declarative renderer: every + // declared wrapper must own at least one declarative routed entry and no + // custom entry (custom wrappers own a dedicated renderer). + llvm::StringSet<> seenWrappers; + for (const BridgeWrapperDecl &decl : whitelist.wrappers) { + if (decl.name.empty()) { + diagOS << "VPTO bridge whitelist: wrapper declaration with an empty " + "name in '" + << sourceName << "'\n"; + return failure(); + } + if (!seenWrappers.insert(decl.name).second) { + diagOS << "VPTO bridge whitelist: duplicate wrapper declaration '" + << decl.name << "' in '" << sourceName << "'\n"; + return failure(); + } + if (decl.includes.empty()) { + diagOS << "VPTO bridge whitelist: wrapper '" << decl.name + << "' declares no includes in '" << sourceName << "'\n"; + return failure(); + } + for (const std::string &include : decl.includes) { + if (include.empty()) { + diagOS << "VPTO bridge whitelist: wrapper '" << decl.name + << "' has an empty include in '" << sourceName << "'\n"; + return failure(); + } + } + if (!decl.core.empty() && decl.core != kBridgeWrapperCoreCube && + decl.core != kBridgeWrapperCoreVec && + decl.core != kBridgeWrapperCoreBoth) { + diagOS << "VPTO bridge whitelist: wrapper '" << decl.name + << "' declares unsupported core '" << decl.core << "' in '" + << sourceName << "' (supported: cube, vec, both; omit it to " + "derive the guard from the routed tile kinds)\n"; + return failure(); + } + if (whitelist.wrapperHasCustomEntry(decl.name)) { + diagOS << "VPTO bridge whitelist: wrapper '" << decl.name + << "' is declared in the wrappers section but carries " + "'lowering: custom' entries owned by a dedicated renderer " + "in '" + << sourceName << "'\n"; + return failure(); + } + bool hasDeclarativeEntry = false; + for (const BridgeWhitelistEntry &entry : whitelist.bridgeOps) { + if (entry.wrapper == decl.name && entry.isDeclarative() && + entry.op != BridgeWhitelist::kInternalOp) { + hasDeclarativeEntry = true; + break; + } + } + if (!hasDeclarativeEntry) { + diagOS << "VPTO bridge whitelist: wrapper '" << decl.name + << "' is declared in the wrappers section but no declarative " + "entry routes into it in '" + << sourceName << "'\n"; + return failure(); + } + } + return whitelist; +} + +FailureOr +pto::parseBridgeWhitelist(llvm::StringRef path, llvm::raw_ostream &diagOS) { + auto bufferOr = llvm::MemoryBuffer::getFile(path); + if (!bufferOr) { + diagOS << "VPTO bridge whitelist: cannot read '" << path + << "': " << bufferOr.getError().message() << "\n"; + return failure(); + } + return parseBridgeWhitelistFromBuffer(bufferOr.get()->getBuffer(), path, + diagOS); +} + +/// The built-in default whitelist covering the wrappers bridged today: +/// The default bridge whitelist currently contains only the pipe family. +/// Pipe entries opt out of the declarative channel because storage lifecycle +/// and TPOP address rebinding require a dedicated lowering pass. +static constexpr llvm::StringLiteral kDefaultBridgeWhitelistYaml = R"yaml( +wrappers: +bridge_ops: + - op: pto.initialize_l2l_pipe + wrapper: pipe + lowering: custom # storage lifecycle: owned by the pipe family pass + entry: pto_vpto_pipe_init + storage_size_entry: pto_vpto_pipe_size + abi: + - type: ptr # storage, synthesized by the bridge lowering + - type: i32 # consumer local buffer address + tmpl_map: + - source: pipe.init + field: pipe + target: Pipe + - op: pto.tpush + wrapper: pipe + lowering: custom # consumes the family pass storage SSA value + entry: pto_vpto_pipe_push + abi: + - type: ptr # storage + - type: i64 # producer tile address + tmpl_map: + - source: tile + field: tile + target: ProducerTile + - op: pto.tpop + wrapper: pipe + lowering: custom # rebinds the tile address to the bridge call result + entry: pto_vpto_pipe_pop + abi: + - type: ptr # storage + tmpl_map: + - source: tile + field: tile + target: ConsumerTile + - op: pto.tfree + wrapper: pipe + lowering: custom # consumes the family pass storage SSA value + entry: pto_vpto_pipe_free + abi: + - type: ptr # storage + - op: internal # wrapper-internal helper, not routed from an IR op + wrapper: pipe + entry: pto_vpto_pipe_size + abi: [] +)yaml"; + +std::string pto::resolveBridgeWhitelistPath(llvm::StringRef optionValue) { + if (!optionValue.empty()) { + return std::string(optionValue); + } + if (const char *envPath = std::getenv("PTOAS_VPTO_BRIDGE_WHITELIST")) { + return envPath; + } + return {}; +} + +FailureOr +pto::loadBridgeWhitelist(llvm::StringRef optionValue, + llvm::raw_ostream &diagOS, std::string *sourceName) { + std::string path = resolveBridgeWhitelistPath(optionValue); + if (sourceName) { + *sourceName = + path.empty() ? kBuiltinBridgeWhitelistSource.str() : path; + } + if (!path.empty()) { + return parseBridgeWhitelist(path, diagOS); + } + return parseBridgeWhitelistFromBuffer( + kDefaultBridgeWhitelistYaml, kBuiltinBridgeWhitelistSource, diagOS); +} diff --git a/lib/PTO/Transforms/VPTOBridgeWrapperGen.cpp b/lib/PTO/Transforms/VPTOBridgeWrapperGen.cpp new file mode 100644 index 0000000000..8251a8d68c --- /dev/null +++ b/lib/PTO/Transforms/VPTOBridgeWrapperGen.cpp @@ -0,0 +1,644 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +//===- VPTOBridgeWrapperGen.cpp - bridge wrapper source generation -------===// +//===----------------------------------------------------------------------===// +// +// Wrapper generation pass of the VPTO C++ interface bridge. The family +// passes collect the wrapper specialization (C++ template tokens built from +// the op attributes/operand types plus the whitelist wrapper entry names) +// into per-function attributes; this pass merges them into the module spec +// and renders it into the complete bridge wrapper C++ source, stored in the +// `pto.vpto.bridge.wrapper_source` module attribute. Object emission then +// compiles the source with Bisheng (once per core kind, selected by the +// __DAV_CUBE__/__DAV_VEC__ guards) and links the bitcode into the device +// modules, replacing the former hand-written wrapper translation unit. +// +// Rendering dispatches on the wrapper the module used. A wrapper whose +// entries carry `lowering: custom` owns a dedicated renderer that knows +// its family semantics (today only pipe: Pipe/Tile typedefs, a +// placement-new init entry, a sizeof size entry, and the producer/consumer +// entries placed on the cores implied by the pipe direction). Every other +// wrapper renders through the generic declarative renderer, driven +// entirely by the whitelist: the wrappers section supplies the includes +// and the core guard, and each used entry its call spelling, template +// arguments and abi-role tile typedefs. Adding a mechanically mapped +// interface therefore needs only a whitelist registration. +// +//===----------------------------------------------------------------------===// + +#include "PTO/Transforms/VPTOBridgeTokens.h" +#include "PTO/Transforms/VPTOBridgeWhitelist.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/IR/Builders.h" +#include "mlir/IR/BuiltinAttributes.h" +#include "mlir/IR/BuiltinOps.h" +#include "mlir/Pass/Pass.h" +#include "llvm/ADT/ArrayRef.h" +#include "llvm/ADT/STLExtras.h" +#include "llvm/ADT/StringRef.h" +#include "llvm/ADT/StringSet.h" +#include "llvm/Support/raw_ostream.h" +#include +#include + +namespace mlir { +namespace pto { + +#define GEN_PASS_DECL_VPTOBRIDGEWRAPPERGEN +#define GEN_PASS_DEF_VPTOBRIDGEWRAPPERGEN +#include "PTO/Transforms/Passes.h.inc" + +namespace { + +/// `wrapper` whitelist values owning a dedicated renderer: entries of such +/// a wrapper carry `lowering: custom` and their family semantics (storage +/// lifecycle, address rebinding) are rendered by hand-written code below +/// instead of the generic declarative renderer. Every other wrapper is +/// whitelist declared and renders generically. +constexpr llvm::StringLiteral kPipeWrapper = "pipe"; + +/// The pipe bridge specialization fields read from the spec DictionaryAttr. +/// The tile typedefs are not read here: they come from the pipe entries' +/// tmpl_map declarations. +struct BridgePipeSpec { + StringRef pipe; + StringRef split; + StringRef entryInit; + StringRef entrySize; + StringRef entryPush; + StringRef entryPop; + StringRef entryFree; +}; + +/// One `using = ;` typedef rendered into the wrapper. +struct BridgeTypedefDecl { + std::string target; + std::string token; +}; + +/// A {spec key -> spec struct field} row used to fill the family spec +/// structs from the spec attribute. +struct BridgeSpecField { + llvm::StringLiteral key; + StringRef *field; +}; + +/// Renders the complete bridge wrapper source for the pipe specialization. +/// For a C2V pipe the cube core produces (push) and the vector core consumes +/// (pop/free); a V2C pipe swaps the roles. Each core section is guarded so +/// that compiling the same source per core kind yields exactly the entries +/// of that core. The typedef section is rendered from the tmpl_map-driven +/// declarations; the body references the fixed Pipe/ProducerTile/ +/// ConsumerTile names. +FailureOr +renderPipeBridgeSource(const BridgePipeSpec &spec, + ArrayRef typedefs) { + bool cubeProduces; + if (spec.pipe.contains("pto::Direction::DIR_C2V")) { + cubeProduces = true; + } else if (spec.pipe.contains("pto::Direction::DIR_V2C")) { + cubeProduces = false; + } else { + return failure(); + } + + std::string source; + llvm::raw_string_ostream os(source); + + os << "// Generated by ptoas (pto-emit-vpto-bridge-wrapper). Do not edit.\n" + << "// VPTO pipe bridge wrapper for " << spec.pipe << ".\n" + << "#include \n" + << "#include \n" + << "#include \n" + << "#include \n" + << "#include \n" + << "#include \n" + << "\n" + << "[aicore] inline void *operator new(size_t, void *ptr) noexcept { " + "return ptr; }\n" + << "\n"; + for (const BridgeTypedefDecl &decl : typedefs) { + os << "using " << decl.target << " = " << decl.token << ";\n"; + } + os << "\n" + << "extern \"C\" [aicore] void " << spec.entryInit + << "(void *storage, uint32_t localBuffer) {\n" + << " new (storage) Pipe(nullptr, localBuffer, 0);\n" + << "}\n" + << "\n" + << "extern \"C\" [aicore] size_t " << spec.entrySize + << "() { return sizeof(Pipe); }\n" + << "\n"; + + auto renderPush = [&]() { + os << "extern \"C\" [aicore] void " << spec.entryPush + << "(void *storage, uint64_t producerAddress) {\n" + << " auto &pipe = *reinterpret_cast(storage);\n" + << " ProducerTile tile;\n" + << " pto::TASSIGN_IMPL(tile, producerAddress);\n" + << " pto::TPUSH(pipe, tile);\n" + << "}\n"; + }; + auto renderPop = [&]() { + os << "extern \"C\" [aicore] uint64_t " << spec.entryPop + << "(void *storage) {\n" + << " auto &pipe = *reinterpret_cast(storage);\n" + << " ConsumerTile tile;\n" + << " pto::TPOP(pipe, tile);\n" + << " pipe_barrier(PIPE_ALL);\n" + << " return reinterpret_cast(tile.data());\n" + << "}\n"; + }; + auto renderFree = [&]() { + os << "extern \"C\" [aicore] void " << spec.entryFree + << "(void *storage) {\n" + << " auto &pipe = *reinterpret_cast(storage);\n" + << " pto::TFREE(pipe);\n" + << "}\n"; + }; + + os << "#ifdef __DAV_CUBE__\n"; + if (cubeProduces) { + renderPush(); + } else { + renderPop(); + os << "\n"; + renderFree(); + } + os << "#endif\n" + << "\n" + << "#ifdef __DAV_VEC__\n"; + if (cubeProduces) { + renderPop(); + os << "\n"; + renderFree(); + } else { + renderPush(); + } + os << "#endif\n"; + + os.flush(); + return source; +} + +/// Returns the C++ parameter spelling of an abi carrier type. +static llvm::StringLiteral bridgeAbiParamType(llvm::StringRef type) { + if (type == "i64") + return "uint64_t"; + if (type == "i32") + return "uint32_t"; + return "void *"; +} + +/// Renders the complete bridge wrapper source of a wrapper declared in the +/// whitelist wrappers section: the declaration supplies the includes and +/// the core guard, and every used declarative entry its abi-bound tile +/// typedefs (named by the CamelCase of the role), the TASSIGN bindings and +/// the declared call. Tile typedef targets sort by name so the rendered +/// source is stable under whitelist reordering; the merged spec guarantees +/// one token per role. A tmpl_args item either is a qualified literal or +/// names an attr spec key, and the whole template argument list is omitted +/// when the spec carries no token for one (e.g. an Unspecified phase). +FailureOr +renderDeclarativeBridgeSource(ModuleOp module, + const BridgeWhitelist &whitelist, + const BridgeWrapperDecl &decl, + DictionaryAttr specAttr, + const llvm::StringSet<> &usedEntries) { + // Tile typedefs: one per abi role the used entries bind. A role token + // missing from the spec means the declarative lowering never routed the + // op; say so instead of rendering an undefined typedef. + std::map typedefTokens; + for (const BridgeWhitelistEntry &entry : whitelist.bridgeOps) { + if (entry.wrapper != decl.name || !entry.isDeclarative() || + !usedEntries.count(entry.entry)) + continue; + for (const BridgeAbiArg &arg : entry.abi) { + std::string target = bridgeRoleTypedefTarget(arg.role); + if (typedefTokens.count(target)) + continue; + auto value = specAttr.getAs(arg.role); + if (!value || value.getValue().empty()) { + module.emitError() + << "VPTO bridge: whitelist entry '" << entry.entry + << "' binds the role '" << arg.role + << "', but no tile token was collected for it; the declarative " + "bridge lowering must run before wrapper generation"; + return failure(); + } + typedefTokens.emplace(target, value.getValue().str()); + } + } + // Attribute tmpl_map rows never render a typedef: their spec tokens are + // constant values (e.g. an AccPhase enumerator), not types; they feed the + // entry's tmpl_args instead. A missing spec token omits the whole + // template argument list at call render time below. + + std::string source; + llvm::raw_string_ostream os(source); + + os << "// Generated by ptoas (pto-emit-vpto-bridge-wrapper). Do not edit.\n" + << "// VPTO bridge wrapper for '" << decl.name << "'.\n" + << "#include \n"; + for (const std::string &include : decl.includes) { + os << "#include <" << include << ">\n"; + } + os << "#include \n" + << "\n"; + for (const auto &typedefToken : typedefTokens) { + os << "using " << typedefToken.first << " = " << typedefToken.second + << ";\n"; + } + os << "\n"; + + // Core guard: declared in the wrappers section, or derived from the tile + // kinds the declarative lowering collected under the reserved + // `core.` spec key. A wrapper whose used entries collected no + // tile has nothing to derive from and must declare `core`. + llvm::StringRef core = decl.core; + if (core.empty()) { + auto coreToken = specAttr.getAs("core." + decl.name); + if (!coreToken || coreToken.getValue().empty()) { + module.emitError() + << "VPTO bridge: wrapper '" << decl.name + << "' declares no core and the declarative lowering collected " + "no tile kind for it; declare 'core: cube|vec|both' in the " + "wrappers section"; + return failure(); + } + core = coreToken.getValue(); + } + const bool guardCube = core == kBridgeWrapperCoreCube; + const bool guardVec = core == kBridgeWrapperCoreVec; + if (guardCube) + os << "#ifdef __DAV_CUBE__\n"; + else if (guardVec) + os << "#ifdef __DAV_VEC__\n"; + + bool firstEntry = true; + for (const BridgeWhitelistEntry &entry : whitelist.bridgeOps) { + if (entry.wrapper != decl.name || !entry.isDeclarative() || + !usedEntries.count(entry.entry)) + continue; + if (!firstEntry) + os << "\n"; + firstEntry = false; + os << "extern \"C\" [aicore] void " << entry.entry << "("; + llvm::interleaveComma(entry.abi, os, [&](const BridgeAbiArg &arg) { + os << bridgeAbiParamType(arg.type) << " " << arg.arg << "Address"; + }); + os << ") {\n"; + for (const BridgeAbiArg &arg : entry.abi) { + os << " " << bridgeRoleTypedefTarget(arg.role) << " " << arg.arg + << ";\n"; + } + for (const BridgeAbiArg &arg : entry.abi) { + os << " pto::TASSIGN_IMPL(" << arg.arg << ", " << arg.arg + << "Address);\n"; + } + // Template arguments: literals render as declared; spec-backed items + // drop the whole list when their token was omitted. + llvm::SmallVector tmplTokens; + bool renderTmplArgs = true; + for (const std::string &item : entry.tmplArgs) { + if (llvm::StringRef(item).contains("::")) { + tmplTokens.push_back(item); + continue; + } + auto value = specAttr.getAs(item); + if (!value || value.getValue().empty()) { + renderTmplArgs = false; + break; + } + tmplTokens.push_back(value.getValue()); + } + os << " " << entry.call; + if (renderTmplArgs && !tmplTokens.empty()) { + os << "<"; + llvm::interleaveComma(tmplTokens, os); + os << ">"; + } + os << "("; + llvm::interleaveComma(entry.abi, os, + [&](const BridgeAbiArg &arg) { os << arg.arg; }); + os << ");\n" + << "}\n"; + } + if (guardCube || guardVec) + os << "#endif\n"; + + os.flush(); + return source; +} + +/// Merges the per-function bridge specs collected by the family pass into the +/// module-level spec attribute. Identical fields deduplicate; a field with +/// two different values (e.g. two pipe configurations) is a conflict. The +/// functions are visited in module order, so the merge is deterministic. +static LogicalResult mergeFuncSpecsIntoModule(ModuleOp module) { + SmallVector specFuncs; + for (auto func : module.getOps()) { + if (func->getAttrOfType(kBridgeFuncSpecAttrName)) + specFuncs.push_back(func); + } + if (specFuncs.empty()) + return success(); + + SmallVector merged; + for (func::FuncOp func : specFuncs) { + auto funcSpec = func->getAttrOfType(kBridgeFuncSpecAttrName); + for (NamedAttribute field : funcSpec) { + bool found = false; + for (NamedAttribute existing : merged) { + if (existing.getName() != field.getName()) + continue; + if (existing.getValue() != field.getValue()) { + func.emitError() + << "VPTO bridge: conflicting bridge specialization across " + "functions; the family configurations must be identical"; + return failure(); + } + found = true; + break; + } + if (!found) + merged.push_back(field); + } + func->removeAttr(kBridgeFuncSpecAttrName); + } + + module->setAttr(kBridgeSpecAttrName, + DictionaryAttr::get(module.getContext(), merged)); + return success(); +} + +/// Maps a tmpl_map `source` of a custom-channel entry to the spec keys +/// carrying its collected token. The pipe `tile` source is split by the +/// routed op: a push entry binds the producer tile, a pop entry the +/// consumer tile. Unknown sources are rejected at whitelist parse time. +/// Declarative wrappers never reach here: their tile typedefs derive from +/// the abi roles. +static llvm::SmallVector +tmplMapSourceSpecKeys(const BridgeWhitelistEntry &entry, + llvm::StringRef source) { + if (source == kPipeInitTmplMapSource) + return {kBridgeSpecPipeKey}; + if (source == kTileTmplMapSource) { + if (entry.op == "pto.tpush") + return {kBridgeSpecProducerTileKey}; + if (entry.op == "pto.tpop") + return {kBridgeSpecConsumerTileKey}; + return {kBridgeSpecProducerTileKey, kBridgeSpecConsumerTileKey}; + } + return {}; +} + +/// Collects the wrapper entry names the module spec actually uses (the +/// entry.* field values), so whitelist processing only touches entries the +/// module bridged. New families only need to collect fields under the +/// entry.* prefix; no per-family key list is maintained here. +static llvm::StringSet<> collectUsedEntries(DictionaryAttr specAttr) { + llvm::StringSet<> usedEntries; + for (NamedAttribute attr : specAttr) { + if (!attr.getName().getValue().starts_with("entry.")) + continue; + if (auto value = dyn_cast(attr.getValue())) + usedEntries.insert(value.getValue()); + } + return usedEntries; +} + +/// Builds the typedef declarations of a custom-channel wrapper from the +/// tmpl_map rows of the entries the module uses, in whitelist order, and +/// validates that every declared template slot is covered by a token the +/// family pass collected into the spec. A target may be declared by several +/// entries; the first declaration wins and the merged spec guarantees the +/// tokens are identical. +static FailureOr> +buildCustomTypedefDecls(ModuleOp module, const BridgeWhitelist &whitelist, + llvm::StringRef wrapper, DictionaryAttr specAttr, + const llvm::StringSet<> &usedEntries) { + SmallVector decls; + llvm::StringSet<> seenTargets; + for (const BridgeWhitelistEntry &entry : whitelist.bridgeOps) { + if (entry.wrapper != wrapper || !usedEntries.count(entry.entry)) + continue; + for (const BridgeTmplMapField &row : entry.tmplMap) { + for (llvm::StringLiteral key : tmplMapSourceSpecKeys(entry, row.source)) { + auto value = specAttr.getAs(key); + if (!value || value.getValue().empty()) { + module.emitError() + << "VPTO bridge: whitelist entry '" << entry.entry + << "' declares tmpl_map target '" << row.target + << "' from source '" << row.source + << "', but no token was collected for it"; + return failure(); + } + if (seenTargets.insert(row.target).second) + decls.push_back({row.target, value.getValue().str()}); + } + } + } + return decls; +} + +struct VPTOBridgeWrapperGenPass final + : public impl::VPTOBridgeWrapperGenBase { + MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(VPTOBridgeWrapperGenPass) + + void runOnOperation() override { + ModuleOp module = getOperation(); + + // Merge the per-function specs into the module spec first: the family + // pass instances may have run concurrently and only write their own + // function attribute, so this single-threaded pass owns the module + // attribute. + if (failed(mergeFuncSpecsIntoModule(module))) { + signalPassFailure(); + return; + } + + auto specAttr = + module->getAttrOfType(kBridgeSpecAttrName); + if (!specAttr) { + // No bridge specialization was collected; nothing to generate. + return; + } + + // Consume the whitelist: the tmpl_map declarations of the entries this + // module uses must be covered by the collected specialization, and they + // drive the wrapper typedef sections. + FailureOr whitelistOr = + loadBridgeWhitelist(whitelistPath, llvm::errs()); + if (failed(whitelistOr)) { + signalPassFailure(); + return; + } + llvm::StringSet<> usedEntries = collectUsedEntries(specAttr); + + // Which wrapper source to render is a whitelist fact, not something to + // infer from which spec keys happen to be present: every entry the + // module used names its wrapper, and entries sharing a wrapper render + // into one translation unit. Deriving the set here means a newly + // bridged declarative interface needs no edit to this pass at all, and + // the diagnostics below name whatever wrappers the whitelist actually + // declares. + SmallVector usedWrappers; + for (const BridgeWhitelistEntry &entry : whitelistOr->bridgeOps) { + if (!usedEntries.count(entry.entry)) + continue; + if (!llvm::is_contained(usedWrappers, StringRef(entry.wrapper))) + usedWrappers.push_back(entry.wrapper); + } + if (usedWrappers.size() != 1) { + InFlightDiagnostic diag = module.emitError(); + if (usedWrappers.empty()) { + diag << "VPTO bridge: the collected specialization names no wrapper " + "entry declared in the bridge whitelist"; + } else { + diag << "VPTO bridge: mixing the "; + for (auto [index, wrapper] : llvm::enumerate(usedWrappers)) { + if (index) + diag << " and "; + diag << "'" << wrapper << "'"; + } + diag << " bridge wrappers in one module is not supported yet"; + } + signalPassFailure(); + return; + } + StringRef usedWrapper = usedWrappers.front(); + + FailureOr source = failure(); + if (whitelistOr->wrapperHasCustomEntry(usedWrapper)) { + if (usedWrapper == kPipeWrapper) { + FailureOr> typedefsOr = + buildCustomTypedefDecls(module, *whitelistOr, usedWrapper, + specAttr, usedEntries); + if (failed(typedefsOr)) { + signalPassFailure(); + return; + } + SmallVector typedefs = std::move(*typedefsOr); + // The pipe entry bodies reference fixed typedef names, so every + // name they use must be rendered by a tmpl_map declaration. + llvm::StringSet<> declTargets; + for (const BridgeTypedefDecl &decl : typedefs) + declTargets.insert(decl.target); + auto requireTypedefTargets = [&](ArrayRef targets) { + bool ok = true; + for (StringRef target : targets) { + if (declTargets.count(target)) + continue; + module.emitError() + << "VPTO bridge: no tmpl_map row renders the '" << target + << "' typedef the wrapper entry bodies need; declare it in " + "the whitelist entry"; + ok = false; + } + return ok; + }; + // All pipe spec fields are mandatory: a pipe bridge kernel always + // carries the full init/size/push/pop/free entry set with a single + // pipe and tile pair configuration. + BridgePipeSpec spec; + BridgeSpecField fields[] = { + {kBridgeSpecPipeKey, &spec.pipe}, + {kBridgeSpecSplitKey, &spec.split}, + {kBridgeSpecEntryInitKey, &spec.entryInit}, + {kBridgeSpecEntrySizeKey, &spec.entrySize}, + {kBridgeSpecEntryPushKey, &spec.entryPush}, + {kBridgeSpecEntryPopKey, &spec.entryPop}, + {kBridgeSpecEntryFreeKey, &spec.entryFree}, + }; + bool ok = true; + // The tile typedefs are rendered from the tmpl_map declarations; + // the renderer only needs to know the tokens were collected. + for (llvm::StringLiteral key : + {kBridgeSpecProducerTileKey, kBridgeSpecConsumerTileKey}) { + auto value = specAttr.getAs(key); + if (!value || value.getValue().empty()) { + module.emitError() + << "VPTO pipe bridge spec is missing the '" << key + << "' field; the pipe family pass must collect it before " + "wrapper generation"; + ok = false; + break; + } + } + if (ok) { + for (const auto &field : fields) { + auto value = specAttr.getAs(field.key); + if (!value || value.getValue().empty()) { + module.emitError() + << "VPTO pipe bridge spec is missing the '" << field.key + << "' field; the pipe family pass must collect it before " + "wrapper generation"; + ok = false; + break; + } + *field.field = value.getValue(); + } + } + if (ok) + ok = requireTypedefTargets({"Pipe", "ProducerTile", "ConsumerTile"}); + if (ok) { + source = renderPipeBridgeSource(spec, typedefs); + if (failed(source)) { + module.emitError() + << "VPTO pipe bridge: cannot render the wrapper source; the " + "pipe token must carry a C2V or V2C direction"; + } + } + } else { + // The whitelist routed these ops into a custom wrapper this pass + // has no dedicated renderer for. Say so instead of falling through + // to the generic renderer, which cannot express family semantics. + module.emitError() + << "VPTO bridge: whitelist entries name the custom wrapper '" + << usedWrapper + << "', which has no dedicated renderer in the bridge wrapper " + "generator (available: " + << kPipeWrapper << ")"; + } + } else { + const BridgeWrapperDecl *decl = whitelistOr->findWrapper(usedWrapper); + if (!decl) { + module.emitError() + << "VPTO bridge: wrapper '" << usedWrapper + << "' routes declarative entries but has no declaration in the " + "whitelist wrappers section (declare its name, includes and " + "core there)"; + } else { + source = renderDeclarativeBridgeSource(module, *whitelistOr, *decl, + specAttr, usedEntries); + } + } + if (failed(source)) { + signalPassFailure(); + return; + } + + OpBuilder builder(module); + module->setAttr(kBridgeWrapperSourceAttrName, + builder.getStringAttr(*source)); + module->removeAttr(kBridgeSpecAttrName); + } +}; + +} // namespace + +std::unique_ptr createVPTOBridgeWrapperGenPass() { + return std::make_unique(); +} + +} // namespace pto +} // namespace mlir diff --git a/lib/PTO/Transforms/VPTOCANN900LLVMEmitter.cpp b/lib/PTO/Transforms/VPTOCANN900LLVMEmitter.cpp index 0019778dd1..ca975c72b2 100644 --- a/lib/PTO/Transforms/VPTOCANN900LLVMEmitter.cpp +++ b/lib/PTO/Transforms/VPTOCANN900LLVMEmitter.cpp @@ -188,7 +188,10 @@ static Type convertVPTOType(Type type, Builder &builder) { if (isa(type)) { return VectorType::get({32}, builder.getI8Type()); } - if (isa(type)) { + if (isa(type)) { + return builder.getI64Type(); + } + if (isa(type)) { return LLVM::LLVMPointerType::get(builder.getContext()); } if (auto ptrType = dyn_cast(type)) { @@ -247,7 +250,7 @@ static bool hasVPTOConvertibleType(Type type) { return false; } if (isa(type) || + pto::StructType, pto::PipeType, pto::TileBufType>(type) || pto::isPTOLowPrecisionType(type)) return true; if (auto vecType = dyn_cast(type)) { @@ -11371,6 +11374,9 @@ static void configureVPTOOpLoweringTarget(ConversionTarget &target, LLVM::LLVMDialect, func::FuncDialect, scf::SCFDialect>(); target.addLegalOp(); + target.addIllegalOp(); target.addIllegalOp(); kernelModulePM.addPass(std::make_unique()); + kernelModulePM.addPass(pto::createVPTOBridgeLoweringPass()); kernelModulePM.addPass(std::make_unique()); kernelModulePM.addPass(std::make_unique()); kernelModulePM.addPass( diff --git a/lib/PTO/Transforms/VPTOLLVMEmitter.cpp b/lib/PTO/Transforms/VPTOLLVMEmitter.cpp index 811ccdb26d..27e94101ec 100644 --- a/lib/PTO/Transforms/VPTOLLVMEmitter.cpp +++ b/lib/PTO/Transforms/VPTOLLVMEmitter.cpp @@ -14203,6 +14203,7 @@ static LogicalResult runPipeline(ModuleOp module, const std::string &march, pm.enableVerifier(); auto &kernelModulePM = pm.nest(); kernelModulePM.addPass(std::make_unique()); + kernelModulePM.addPass(pto::createVPTOBridgeLoweringPass()); kernelModulePM.addPass(std::make_unique(march)); kernelModulePM.addPass(std::make_unique()); kernelModulePM.addPass( diff --git a/lib/PTO/Transforms/VPTOSplitCVModule.cpp b/lib/PTO/Transforms/VPTOSplitCVModule.cpp index 785daf605b..515a975993 100644 --- a/lib/PTO/Transforms/VPTOSplitCVModule.cpp +++ b/lib/PTO/Transforms/VPTOSplitCVModule.cpp @@ -14,6 +14,7 @@ #include "mlir/IR/SymbolTable.h" #include "mlir/Pass/Pass.h" #include "llvm/ADT/STLExtras.h" +#include "llvm/ADT/StringSet.h" namespace mlir { namespace pto { @@ -38,6 +39,22 @@ static bool hasKernelKindChildModule(ModuleOp module) { static bool hasCVSections(ModuleOp module); +static std::optional +getFunctionKernelKind(func::FuncOp funcOp) { + auto attr = funcOp->getAttrOfType( + FunctionKernelKindAttr::name); + if (!attr) + return std::nullopt; + return attr.getKernelKind(); +} + +static bool hasFunctionKernelKind(ModuleOp module, FunctionKernelKind kind) { + return llvm::any_of(module.getOps(), [&](func::FuncOp funcOp) { + auto functionKind = getFunctionKernelKind(funcOp); + return functionKind && *functionKind == kind; + }); +} + static bool isVPTOBackendModule(ModuleOp module) { auto backend = module->getAttrOfType("pto.backend"); return backend && backend.getValue() == "vpto"; @@ -302,12 +319,42 @@ static void rewriteSectionsForKind(ModuleOp module, FunctionKernelKind kind) { } } +static void pruneFunctionsForKind(ModuleOp module, FunctionKernelKind kind) { + llvm::StringSet<> removedSymbols; + SmallVector eraseFuncs; + for (func::FuncOp funcOp : module.getOps()) { + auto functionKind = getFunctionKernelKind(funcOp); + if (functionKind && *functionKind != kind) { + removedSymbols.insert(funcOp.getSymName()); + eraseFuncs.push_back(funcOp); + } + } + + SmallVector eraseCalls; + module.walk([&](func::CallOp callOp) { + if (removedSymbols.contains(callOp.getCallee())) + eraseCalls.push_back(callOp); + }); + for (func::CallOp callOp : eraseCalls) + callOp.erase(); + for (func::FuncOp funcOp : eraseFuncs) + funcOp.erase(); +} + static ModuleOp cloneModuleForKind(ModuleOp source, FunctionKernelKind kind, - OpBuilder &builder) { + OpBuilder &builder, + bool functionKindInput = false) { auto cloned = cast(source->clone()); cloned->setAttr(FunctionKernelKindAttr::name, FunctionKernelKindAttr::get(cloned.getContext(), kind)); - eraseSectionSplitCandidatesWithoutSectionKind(cloned, kind); + pruneFunctionsForKind(cloned, kind); + cloned.walk([&](func::FuncOp funcOp) { + if (!funcOp.isDeclaration()) + funcOp->setAttr(FunctionKernelKindAttr::name, + FunctionKernelKindAttr::get(funcOp.getContext(), kind)); + }); + if (!functionKindInput) + eraseSectionSplitCandidatesWithoutSectionKind(cloned, kind); rewriteSectionsForKind(cloned, kind); builder.insert(cloned); return cloned; @@ -327,6 +374,37 @@ static LogicalResult materializeExplicitKernelKindSections(ModuleOp module) { return success(); } +/// Wraps `module` in a fresh outer module holding one cloned child per +/// requested kernel kind, then replaces `module`'s body and attributes with +/// the outer module's. `functionKindInput` selects the per-function +/// kernel-kind pruning path in cloneModuleForKind instead of the +/// section-based one. +static void splitIntoKindModules(ModuleOp module, bool needVector, + bool needCube, bool functionKindInput) { + SmallVector outerAttrs; + outerAttrs.reserve(module->getAttrs().size()); + for (NamedAttribute attr : module->getAttrs()) { + if (attr.getName() != SymbolTable::getSymbolAttrName()) { + outerAttrs.push_back(attr); + } + } + + auto outer = ModuleOp::create(module.getLoc()); + outer->setAttrs(DictionaryAttr::get(module.getContext(), outerAttrs)); + OpBuilder builder(outer.getBody(), outer.getBody()->end()); + if (needVector) { + cloneModuleForKind(module, FunctionKernelKind::Vector, builder, + functionKindInput); + } + if (needCube) { + cloneModuleForKind(module, FunctionKernelKind::Cube, builder, + functionKindInput); + } + + module.getBodyRegion().takeBody(outer.getBodyRegion()); + module->setAttrs(outer->getAttrs()); +} + static LogicalResult splitCVModule(ModuleOp module) { flattenSingleUnpartitionedChild(module); if (hasKernelKind(module)) { @@ -343,6 +421,15 @@ static LogicalResult splitCVModule(ModuleOp module) { } return success(); } + bool hasVectorFunctions = + hasFunctionKernelKind(module, FunctionKernelKind::Vector); + bool hasCubeFunctions = + hasFunctionKernelKind(module, FunctionKernelKind::Cube); + if (hasVectorFunctions || hasCubeFunctions) { + splitIntoKindModules(module, hasVectorFunctions, hasCubeFunctions, + /*functionKindInput=*/true); + return success(); + } if (!hasCVSections(module)) { return success(); } @@ -358,26 +445,8 @@ static LogicalResult splitCVModule(ModuleOp module) { return success(); } - SmallVector outerAttrs; - outerAttrs.reserve(module->getAttrs().size()); - for (NamedAttribute attr : module->getAttrs()) { - if (attr.getName() != SymbolTable::getSymbolAttrName()) { - outerAttrs.push_back(attr); - } - } - - auto outer = ModuleOp::create(module.getLoc()); - outer->setAttrs(DictionaryAttr::get(module.getContext(), outerAttrs)); - OpBuilder builder(outer.getBody(), outer.getBody()->end()); - if (needVector) { - cloneModuleForKind(module, FunctionKernelKind::Vector, builder); - } - if (needCube) { - cloneModuleForKind(module, FunctionKernelKind::Cube, builder); - } - - module.getBodyRegion().takeBody(outer.getBodyRegion()); - module->setAttrs(outer->getAttrs()); + splitIntoKindModules(module, needVector, needCube, + /*functionKindInput=*/false); return success(); } diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist-decl-missing-wrappers.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist-decl-missing-wrappers.yaml new file mode 100644 index 0000000000..57f2f45b0b --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist-decl-missing-wrappers.yaml @@ -0,0 +1,13 @@ +# Diagnostic fixture: a valid declarative entry whose wrapper has no +# declaration in the wrappers section. The entry parses and lowers, but +# the generic renderer needs the declaration's includes and core guard, so +# wrapper generation rejects the module with a diagnostic naming the +# missing section. +bridge_ops: + - op: pto.tmatmul + wrapper: matmul + call: pto::TMATMUL + abi: + - {operand: 2, arg: dst, role: result_tile} + - {operand: 0, arg: lhs, role: left_tile} + - {operand: 1, arg: rhs, role: right_tile} diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist-legacy-family-lowering.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist-legacy-family-lowering.yaml new file mode 100644 index 0000000000..8670dde9e8 --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist-legacy-family-lowering.yaml @@ -0,0 +1,13 @@ +# Diagnostic fixture: `lowering: family` was the pre-flip spelling of the +# opt-out channel, back when it was also the default. It is no longer an +# accepted value -- the opt-out is spelled `lowering: custom` -- and a +# whitelist still using it must be rejected with the supported set listed, +# not silently reinterpreted. +bridge_ops: + - op: pto.initialize_l2l_pipe + wrapper: pipe + lowering: family + entry: pto_vpto_pipe_init + abi: + - type: ptr + - type: i32 diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist-missing-custom-tag.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist-missing-custom-tag.yaml new file mode 100644 index 0000000000..9b2c5ff7e9 --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist-missing-custom-tag.yaml @@ -0,0 +1,14 @@ +# Diagnostic fixture: a pipe entry that needs the dedicated family pass but +# never opted out of the default declarative channel. Because `lowering` +# defaults to `declarative`, the entry is validated as a declarative one and +# rejected at parse time for its unbound abi arguments -- with the missing +# `lowering: custom` tag named -- instead of parsing clean and only failing +# after lowering, at the leftover routing check, against a family pass that +# was never going to claim it. +bridge_ops: + - op: pto.initialize_l2l_pipe + wrapper: pipe + entry: pto_vpto_pipe_init + abi: + - type: ptr + - type: i32 diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist-no-routing.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist-no-routing.yaml new file mode 100644 index 0000000000..4fc74e1690 --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist-no-routing.yaml @@ -0,0 +1,6 @@ +# Explicitly empty VPTO bridge whitelist: routes no IR op to the C++ +# interface bridge. Since Phase 4 ships a built-in default whitelist +# (which routes pto.tmatmul to the TMATMUL bridge), tests that exercise the +# regular tile-op expansion path (tmatmul -> pto.mad) must opt out of the +# bridge routing with this document. +bridge_ops: [] diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist-tmpl-missing-field.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist-tmpl-missing-field.yaml new file mode 100644 index 0000000000..7c94b2df8b --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist-tmpl-missing-field.yaml @@ -0,0 +1,13 @@ +# tmpl_map diagnostic fixture: the row below lacks the mandatory 'target' +# key, which the whitelist parser must reject. +bridge_ops: + - op: pto.initialize_l2l_pipe + wrapper: pipe + lowering: custom # storage lifecycle: owned by the pipe family pass + entry: pto_vpto_pipe_init + abi: + - type: ptr + - type: i32 + tmpl_map: + - source: pipe.init + field: slot_size diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist-tmpl-unknown-source.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist-tmpl-unknown-source.yaml new file mode 100644 index 0000000000..5e7b02b138 --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist-tmpl-unknown-source.yaml @@ -0,0 +1,15 @@ +# tmpl_map diagnostic fixture: the row below uses a source token outside the +# pipe-family known set (pipe.init, tile), which the whitelist parser must +# reject. +bridge_ops: + - op: pto.initialize_l2l_pipe + wrapper: pipe + lowering: custom # storage lifecycle: owned by the pipe family pass + entry: pto_vpto_pipe_init + abi: + - type: ptr + - type: i32 + tmpl_map: + - source: pipe.pop + field: split + target: TPipe::slotSize diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist-unknown-wrapper.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist-unknown-wrapper.yaml new file mode 100644 index 0000000000..2e9691d6f1 --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist-unknown-wrapper.yaml @@ -0,0 +1,34 @@ +# Diagnostic fixture: a whitelist whose routed ops name a `wrapper` value the +# bridge wrapper generator has no renderer for. The generator dispatches its +# renderer on the wrapper names of the entries the module used, so instead of +# falling through into some other wrapper's renderer (and failing on its +# missing spec fields) it must reject the module with a diagnostic naming the +# unknown wrapper and the available ones. +bridge_ops: + - op: pto.initialize_l2l_pipe + wrapper: pipe_unfinished_port + lowering: custom + entry: pto_vpto_pipe_init + storage_size_entry: pto_vpto_pipe_size + abi: + - type: ptr + - type: i32 + tmpl_map: + - source: pipe.init + field: pipe + target: Pipe + - op: pto.tpush + wrapper: pipe_unfinished_port + lowering: custom + entry: pto_vpto_pipe_push + abi: + - type: ptr + - type: i64 + tmpl_map: + - source: tile + field: tile + target: ProducerTile + - op: internal + wrapper: pipe_unfinished_port + entry: pto_vpto_pipe_size + abi: [] diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist-variant.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist-variant.yaml new file mode 100644 index 0000000000..7afaf873a4 --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist-variant.yaml @@ -0,0 +1,46 @@ +# Variant VPTO pipe bridge whitelist with declarative tmpl_map rows, used by +# the Phase 2 bridge spec/wrapper generation lit tests. The tmpl_map rows are +# declarative IR-field -> C++-template-slot mappings validated at parse time; +# the token construction itself lives in VPTOBridgeTokens. +bridge_ops: + - op: pto.initialize_l2l_pipe + wrapper: pipe + lowering: custom # storage lifecycle: owned by the pipe family pass + entry: pto_vpto_pipe_init + storage_size_entry: pto_vpto_pipe_size + abi: + - type: ptr # storage, synthesized by the bridge lowering + - type: i32 # consumer local buffer address + tmpl_map: + - source: pipe.init + field: slot_size + target: TPipe::slotSize + - source: pipe.init + field: slot_num + target: TPipe::slotNum + - source: tile + field: rows + target: Tile::Rows + - op: pto.tpush + wrapper: pipe + lowering: custom # consumes the family pass storage SSA value + entry: pto_vpto_pipe_push + abi: + - type: ptr # storage + - type: i64 # producer tile address + - op: pto.tpop + wrapper: pipe + lowering: custom # rebinds the tile address to the bridge call result + entry: pto_vpto_pipe_pop + abi: + - type: ptr # storage + - op: pto.tfree + wrapper: pipe + lowering: custom # consumes the family pass storage SSA value + entry: pto_vpto_pipe_free + abi: + - type: ptr # storage + - op: internal # wrapper-internal helper, not routed from an IR op + wrapper: pipe + entry: pto_vpto_pipe_size + abi: [] diff --git a/test/lit/vpto/Inputs/vpto-bridge-whitelist.yaml b/test/lit/vpto/Inputs/vpto-bridge-whitelist.yaml new file mode 100644 index 0000000000..804efcb51d --- /dev/null +++ b/test/lit/vpto/Inputs/vpto-bridge-whitelist.yaml @@ -0,0 +1,34 @@ +# VPTO C++ interface bridge whitelist used by the lit bridge tests. Mirrors +# the fifo-tile-data-consume case whitelist. +bridge_ops: + - op: pto.initialize_l2l_pipe + wrapper: pipe + lowering: custom # storage lifecycle: owned by the pipe family pass + entry: pto_vpto_pipe_init + storage_size_entry: pto_vpto_pipe_size + abi: + - type: ptr # storage, synthesized by the bridge lowering + - type: i32 # consumer local buffer address + - op: pto.tpush + wrapper: pipe + lowering: custom # consumes the family pass storage SSA value + entry: pto_vpto_pipe_push + abi: + - type: ptr # storage + - type: i64 # producer tile address + - op: pto.tpop + wrapper: pipe + lowering: custom # rebinds the tile address to the bridge call result + entry: pto_vpto_pipe_pop + abi: + - type: ptr # storage + - op: pto.tfree + wrapper: pipe + lowering: custom # consumes the family pass storage SSA value + entry: pto_vpto_pipe_free + abi: + - type: ptr # storage + - op: internal # wrapper-internal helper, not routed from an IR op + wrapper: pipe + entry: pto_vpto_pipe_size + abi: [] diff --git a/test/lit/vpto/cube/expand_tile_op_tilelang_tmatmul.pto b/test/lit/vpto/cube/expand_tile_op_tilelang_tmatmul.pto index 6c5fe2e23c..39a36dbfba 100644 --- a/test/lit/vpto/cube/expand_tile_op_tilelang_tmatmul.pto +++ b/test/lit/vpto/cube/expand_tile_op_tilelang_tmatmul.pto @@ -8,8 +8,10 @@ // Guard: TileLang cube tmatmul expansion on the dav-c310-cube VPTO path must // inline the template body and lower the tile op to cube MAD instructions. -// -// RUN: ptoas --pto-arch=a5 --pto-backend=vpto --enable-tile-op-expand --emit-vpto %s -o - 2>/dev/null | FileCheck %s +// The empty bridge whitelist opts this test out of the built-in default +// routing (which would send pto.tmatmul to the TMATMUL C++ bridge). + +// RUN: env PTOAS_VPTO_BRIDGE_WHITELIST=%S/../Inputs/vpto-bridge-whitelist-no-routing.yaml ptoas --pto-arch=a5 --pto-backend=vpto --enable-tile-op-expand --emit-vpto %s -o - 2>/dev/null | FileCheck %s // CHECK-LABEL: func.func @TMATMUL // CHECK-NOT: pto.tmatmul ins diff --git a/test/lit/vpto/legacy_aicore_kernel_attr.pto b/test/lit/vpto/legacy_aicore_kernel_attr.pto index 25a7c6e386..7793666c4e 100644 --- a/test/lit/vpto/legacy_aicore_kernel_attr.pto +++ b/test/lit/vpto/legacy_aicore_kernel_attr.pto @@ -26,10 +26,10 @@ module attributes {pto.target_arch = "a5"} { // SPLIT: module attributes {{.*}}pto.kernel_kind = #pto.kernel_kind // SPLIT: func.func @legacy_attr_kernel -// SPLIT-SAME: attributes {pto.aicore} +// SPLIT-SAME: attributes {pto.aicore, pto.kernel_kind = #pto.kernel_kind} // SPLIT: module attributes {{.*}}pto.kernel_kind = #pto.kernel_kind // SPLIT: func.func @legacy_attr_kernel -// SPLIT-SAME: attributes {pto.aicore} +// SPLIT-SAME: attributes {pto.aicore, pto.kernel_kind = #pto.kernel_kind} // LLVM-DAG: llvm.func @legacy_attr_kernel_mix_aiv // LLVM-DAG: llvm.func @legacy_attr_kernel_mix_aic diff --git a/test/lit/vpto/ptodsl_tileop_split_simt_entry.pto b/test/lit/vpto/ptodsl_tileop_split_simt_entry.pto index 6160fd167d..6b05bd2932 100644 --- a/test/lit/vpto/ptodsl_tileop_split_simt_entry.pto +++ b/test/lit/vpto/ptodsl_tileop_split_simt_entry.pto @@ -10,9 +10,11 @@ // CHECK: module attributes {{.*}}pto.kernel_kind = #pto.kernel_kind // CHECK: func.func @vector_kernel -// CHECK: func.func @vector_simt() attributes {no_inline, pto.simt_entry} +// CHECK-SAME: attributes {pto.kernel, pto.kernel_kind = #pto.kernel_kind} +// CHECK: func.func @vector_simt() attributes {no_inline, pto.kernel_kind = #pto.kernel_kind, pto.simt_entry} // CHECK: module attributes {{.*}}pto.kernel_kind = #pto.kernel_kind // CHECK: func.func @vector_kernel +// CHECK-SAME: attributes {pto.kernel, pto.kernel_kind = #pto.kernel_kind} // CHECK-NOT: func.func @vector_simt() module attributes {pto.target_arch = "a5"} { diff --git a/test/lit/vpto/vpto_bridge_declarative_unrouted_passthrough.pto b/test/lit/vpto/vpto_bridge_declarative_unrouted_passthrough.pto new file mode 100644 index 0000000000..5ef8ba869c --- /dev/null +++ b/test/lit/vpto/vpto_bridge_declarative_unrouted_passthrough.pto @@ -0,0 +1,69 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Declarative routing is whitelist driven and must never over-lower: an op +// the whitelist does not route keeps flowing through the regular tile-op +// expansion path (tmatmul -> pto.mad), and an op routed to the family +// channel is left for its family pass. Both forms of pass-through are +// checked here against the empty-routing and the pipe-family whitelists. + +// RUN: pto-test-opt %s "-pto-lower-declarative-bridge-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist-no-routing.yaml" | FileCheck %s --check-prefix=NOROUTE +// RUN: pto-test-opt %s "-pto-lower-declarative-bridge-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist-variant.yaml" | FileCheck %s --check-prefix=FAMILY + +// NOROUTE-LABEL: func.func private @unrouted_matmul_cube +// NOROUTE: pto.tmatmul +// NOROUTE-NOT: pto.bridge_call +// NOROUTE-NOT: pto.vpto.bridge.func_spec + +// FAMILY-LABEL: func.func private @family_routed_pipe_cube +// FAMILY: pto.tpush +// FAMILY-NOT: pto.bridge_call +// FAMILY-NOT: pto.vpto.bridge.func_spec + +module attributes {pto.target_arch = "a5"} { + func.func private @unrouted_matmul_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %c1024 = arith.constant 1024 : i64 + %c2048 = arith.constant 2048 : i64 + %lhs = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %rhs = pto.alloc_tile addr = %c1024 + : !pto.tile_buf + %acc = pto.alloc_tile addr = %c2048 + : !pto.tile_buf + pto.tmatmul ins(%lhs, %rhs + : !pto.tile_buf, + !pto.tile_buf) + outs(%acc : !pto.tile_buf) + return + } + + func.func private @family_routed_pipe_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_default_whitelist_lowering.pto b/test/lit/vpto/vpto_bridge_default_whitelist_lowering.pto new file mode 100644 index 0000000000..104c0d7707 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_default_whitelist_lowering.pto @@ -0,0 +1,41 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Phase 4 formal whitelist channel: with neither the whitelist-path pass +// option nor PTOAS_VPTO_BRIDGE_WHITELIST configured, the family passes and +// the wrapper generation fall back to the built-in default whitelist +// shipped with ptoas, so pipe family ops are bridged out of the box. The +// env var is explicitly unset so the test exercises the built-in default +// even on machines exporting a whitelist. + +// RUN: env -u PTOAS_VPTO_BRIDGE_WHITELIST pto-test-opt %s -pto-lower-pipe-family-ops | FileCheck %s + +// CHECK-LABEL: func.func private @default_whitelist_cube +// CHECK: pto.bridge_call "pto_vpto_pipe_init"{{.*}}storage_size_callee = "pto_vpto_pipe_size" +// CHECK: pto.bridge_call "pto_vpto_pipe_push" +// CHECK-NOT: pto.alloc_tile +// CHECK-NOT: pto.tpush + +module { + func.func private @default_whitelist_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_emitc_token_parity.pto b/test/lit/vpto/vpto_bridge_emitc_token_parity.pto new file mode 100644 index 0000000000..b2976f9d9e --- /dev/null +++ b/test/lit/vpto/vpto_bridge_emitc_token_parity.pto @@ -0,0 +1,85 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Token parity guard (design decision 3): the EmitC backend and the VPTO +// bridge each own a copy of the "IR op attributes -> C++ template token" +// logic (PTOToEmitC.cpp vs VPTOBridgeTokens.cpp). For the same pipe +// configuration the two copies must render field-identical template +// arguments. This test feeds one module through both paths and pins both +// renderings side by side so any drift fails loudly. +// +// The pipe configuration deliberately uses a non-default variant +// (flag_base=8, slot_size=2048, slot_num=4) so hardcoded defaults cannot +// pass. The BRIDGE run uses the built-in default whitelist (env unset) so +// the Phase 4 formal channel is exercised too. +// +// Known deliberate textual differences (not drift): +// - the bridge tokens carry the `pto::` namespace prefix; +// - the bridge NoneBox tile token omits the trailing +// `SLayout::NoneBox, , PadValue::Null, CompactMode::Null` +// arguments and relies on the Tile template defaults, matching the +// hand-written wrapper specializations; +// - the split token is EmitC-only; the bridge enforces split=1. +// +// Both runs use --pto-level=level3 so the shared input may carry explicit +// planned addresses on alloc_tile: the bridge TPUSH requires them, while +// the default level rejects them. + +// RUN: ptoas --pto-arch=a5 --pto-level=level3 %s 2>&1 | FileCheck %s --check-prefix=EMITC +// RUN: env -u PTOAS_VPTO_BRIDGE_WHITELIST pto-test-opt %s -pto-lower-pipe-family-ops -pto-emit-vpto-bridge-wrapper | FileCheck %s --check-prefix=BRIDGE + +module { + func.func @parity_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 2048, slot_num = 4, flag_base = 8, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } + + func.func @parity_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 2048, slot_num = 4, flag_base = 8, nosplit = false} (%addr : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + pto.tfree(%pipe : !pto.pipe) {split = 1} + return + } + + func.func @entry(%addr: i32) attributes {pto.entry} { + call @parity_cube(%addr) : (i32) -> () + call @parity_vector(%addr) : (i32) -> () + return + } +} + +// EMITC: TPipe<8, Direction::DIR_C2V, 2048, 4, 2, false> +// EMITC: TPUSH, Tile, TileSplitAxis::TILE_UP_DOWN> +// EMITC: TPOP, Tile, TileSplitAxis::TILE_UP_DOWN> + +// BRIDGE: using Pipe = pto::TPipe<8, pto::Direction::DIR_C2V, 2048, 4, 2, false>; +// BRIDGE: using ProducerTile = pto::Tile; +// BRIDGE: using ConsumerTile = pto::Tile; diff --git a/test/lit/vpto/vpto_bridge_pipe_family_lowering.pto b/test/lit/vpto/vpto_bridge_pipe_family_lowering.pto new file mode 100644 index 0000000000..fe77a6d4ee --- /dev/null +++ b/test/lit/vpto/vpto_bridge_pipe_family_lowering.pto @@ -0,0 +1,88 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// RUN: pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist.yaml" | FileCheck %s --check-prefix=FAMILY +// RUN: pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist.yaml" "-vpto-bridge-lowering=whitelist-path=%S/Inputs/vpto-bridge-whitelist.yaml" | FileCheck %s --check-prefix=LOWER + +// FAMILY-LABEL: func.func private @bridge_cube +// FAMILY: pto.bridge_call "pto_vpto_pipe_init"{{.*}}storage_size_callee = "pto_vpto_pipe_size" +// FAMILY-SAME: -> !pto.pipe +// FAMILY: pto.bridge_call "pto_vpto_pipe_push" +// FAMILY-NOT: pto.alloc_tile +// FAMILY-NOT: pto.tpush + +// FAMILY-LABEL: func.func private @bridge_vector +// FAMILY: pto.bridge_call "pto_vpto_pipe_init" +// FAMILY: pto.bridge_call "pto_vpto_pipe_pop" +// FAMILY-SAME: -> i64 +// FAMILY: pto.bridge_inttoptr +// FAMILY: pto.bridge_call "pto_vpto_pipe_free" +// FAMILY-NOT: pto.declare_tile +// FAMILY-NOT: pto.tpop +// FAMILY-NOT: pto.tile_buf_addr +// FAMILY-NOT: pto.tfree + +// LOWER-DAG: func.func private @pto_vpto_pipe_size() -> i64 +// LOWER-DAG: func.func private @pto_vpto_pipe_init(!llvm.ptr, i32) +// LOWER-DAG: func.func private @pto_vpto_pipe_push(!llvm.ptr, i64) +// LOWER-DAG: func.func private @pto_vpto_pipe_pop(!llvm.ptr) -> i64 +// LOWER-DAG: func.func private @pto_vpto_pipe_free(!llvm.ptr) + +// LOWER-LABEL: func.func private @bridge_cube +// LOWER: call @pto_vpto_pipe_size() : () -> i64 +// LOWER: llvm.alloca +// LOWER: call @pto_vpto_pipe_init +// LOWER: call @pto_vpto_pipe_push +// LOWER-NOT: pto.bridge_call + +// LOWER-LABEL: func.func private @bridge_vector +// LOWER: call @pto_vpto_pipe_size() : () -> i64 +// LOWER: llvm.alloca +// LOWER: call @pto_vpto_pipe_init +// LOWER: call @pto_vpto_pipe_pop +// LOWER: llvm.inttoptr +// LOWER: call @pto_vpto_pipe_free +// LOWER-NOT: pto.bridge_call +// LOWER-NOT: pto.bridge_inttoptr + +module { + func.func private @bridge_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } + + func.func private @bridge_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + pto.tfree(%pipe : !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_pipe_family_skip_pipeless.pto b/test/lit/vpto/vpto_bridge_pipe_family_skip_pipeless.pto new file mode 100644 index 0000000000..f12e7164fe --- /dev/null +++ b/test/lit/vpto/vpto_bridge_pipe_family_skip_pipeless.pto @@ -0,0 +1,35 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// The pipe family pass must leave pipe-less functions untouched: their tile +// handles keep flowing through the regular lowering, and no whitelist is +// required. + +// RUN: pto-test-opt %s -pto-lower-pipe-family-ops | FileCheck %s + +// CHECK-LABEL: func.func @pipeless_tile_handles +// CHECK: pto.alloc_tile +// CHECK: pto.tile_buf_addr +// CHECK-NOT: pto.bridge_call +// CHECK-NOT: pto.bridge_inttoptr + +module { + func.func @pipeless_tile_handles() attributes {pto.entry} { + %tile = pto.alloc_tile + : !pto.tile_buf + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + return + } +} diff --git a/test/lit/vpto/vpto_bridge_pipe_loop_consume.pto b/test/lit/vpto/vpto_bridge_pipe_loop_consume.pto new file mode 100644 index 0000000000..3fb2ea7383 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_pipe_loop_consume.pto @@ -0,0 +1,98 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Correctness boundary probe turned regression: a FIFO consumed inside an +// scf.for loop (TPUSH per producer iteration, TPOP + address consumption + +// TFREE per consumer iteration). The IR holds a single pop op executed +// repeatedly at runtime, so the bridge must rebind the consumer tile on +// every iteration: the pop call result is a loop-body SSA value feeding a +// loop-body bridge_inttoptr, which pins each iteration to its own FIFO +// slot. The EmitC run anchors the reference loop shape (tile declared +// outside the loop, TPOP rebinding it inside) so the bridge lowering stays +// semantically aligned with the production path. +// +// RUN: env -u PTOAS_VPTO_BRIDGE_WHITELIST pto-test-opt %s -pto-lower-pipe-family-ops -pto-emit-vpto-bridge-wrapper | FileCheck %s --check-prefix=BRIDGE +// RUN: ptoas --pto-arch=a5 --pto-level=level3 %s 2>&1 | FileCheck %s --check-prefix=EMITC + +module { + func.func @probe_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c4 = arith.constant 4 : index + %addr0 = arith.constant 0 : i64 + %pipe_addr = arith.constant 0 : i32 + %acc = pto.alloc_tile addr = %addr0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 2048, slot_num = 4, + flag_base = 8, nosplit = false} (%pipe_addr : i32) -> !pto.pipe + scf.for %i = %c0 to %c4 step %c1 { + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + } + return + } + + func.func @probe_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c4 = arith.constant 4 : index + %addr0 = arith.constant 0 : i32 + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 2048, slot_num = 4, + flag_base = 8, nosplit = false} (%addr0 : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + scf.for %i = %c0 to %c4 step %c1 { + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + %off = arith.index_cast %i : index to i64 + %off_f32 = arith.sitofp %off : i64 to f32 + pto.store_scalar %off_f32, %ub[%i] : !pto.ptr, f32 + pto.tfree(%pipe : !pto.pipe) {split = 1} + } + return + } +} + +// The per-iteration pop address must stay inside the loop body: the +// bridge_call result feeds the bridge_inttoptr and the store in the same +// scf.for region, so each iteration consumes its own FIFO slot. + +// BRIDGE-LABEL: func.func @probe_cube +// BRIDGE: scf.for +// BRIDGE: pto.bridge_call "pto_vpto_pipe_push" + +// BRIDGE-LABEL: func.func @probe_vector +// BRIDGE: scf.for +// BRIDGE: [[POP:%.+]] = pto.bridge_call "pto_vpto_pipe_pop"{{.*}}-> i64 +// BRIDGE: [[PTR:%.+]] = pto.bridge_inttoptr [[POP]] : i64 -> !pto.ptr +// BRIDGE: pto.store_scalar {{%.+}}, [[PTR]] +// BRIDGE: pto.bridge_call "pto_vpto_pipe_free" + +// EMITC-LABEL: AICORE void probe_cube +// EMITC: for ( +// EMITC: TPUSH + +// EMITC-LABEL: AICORE void probe_vector +// EMITC: for ( +// EMITC: TPOP +// EMITC: TFREE diff --git a/test/lit/vpto/vpto_bridge_pipe_split_left_right.pto b/test/lit/vpto/vpto_bridge_pipe_split_left_right.pto new file mode 100644 index 0000000000..8894fa14e7 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_pipe_split_left_right.pto @@ -0,0 +1,65 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Non-unit split on the pipe entries: tpush/tpop/tfree carry split = 2, +// which maps to pto::TileSplitAxis::TILE_LEFT_RIGHT. The family pass fixes +// the split token on the first bridged op of the function and every later +// op checks against it; the wrapper renders the shared TileSplitAxis into +// all three entry bodies. + +// RUN: pto-test-opt %s -pto-lower-pipe-family-ops | FileCheck %s --check-prefix=SPEC +// RUN: pto-test-opt %s -pto-lower-pipe-family-ops -pto-emit-vpto-bridge-wrapper | FileCheck %s --check-prefix=WRAP + +// SPEC-LABEL: func.func private @bridge_cube +// SPEC-SAME: pto.vpto.bridge.func_spec = { +// SPEC-SAME: split = "pto::TileSplitAxis::TILE_LEFT_RIGHT" + +// SPEC-LABEL: func.func private @bridge_vector +// SPEC-SAME: pto.vpto.bridge.func_spec = { +// SPEC-SAME: split = "pto::TileSplitAxis::TILE_LEFT_RIGHT" + +// WRAP: module attributes {{{.*}}pto.vpto.bridge.wrapper_source = " +// WRAP-SAME: pto::TPUSH(pipe, tile); +// WRAP-SAME: pto::TPOP(pipe, tile); +// WRAP-SAME: pto::TFREE(pipe); + +module { + func.func private @bridge_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 2} + return + } + + func.func private @bridge_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 2} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + pto.tfree(%pipe : !pto.pipe) {split = 2} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_pipe_split_mismatch_diag.pto b/test/lit/vpto/vpto_bridge_pipe_split_mismatch_diag.pto new file mode 100644 index 0000000000..9a0e9ae0f8 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_pipe_split_mismatch_diag.pto @@ -0,0 +1,38 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// The wrapper renders one shared TileSplitAxis for the push/pop/free +// entries, so all bridged pipe ops of a function must agree on the split. +// The first bridged op fixes the token; a later op with a different split +// is diagnosed instead of silently rendering the wrong axis. + +// RUN: not pto-test-opt %s -pto-lower-pipe-family-ops 2>&1 | FileCheck %s + +// CHECK: VPTO pipe bridge TFREE split 2 does not match the split 1 already bridged in this function; the wrapper renders one shared TileSplitAxis + +module { + func.func private @bridge_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + pto.tfree(%pipe : !pto.pipe) {split = 2} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_pop_rebind_diag.pto b/test/lit/vpto/vpto_bridge_pop_rebind_diag.pto new file mode 100644 index 0000000000..5f1878cfef --- /dev/null +++ b/test/lit/vpto/vpto_bridge_pop_rebind_diag.pto @@ -0,0 +1,57 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Sequential rebind boundary: two TPOPs rebinding the same declared tile +// (the unrolled form of streaming consumption) are not supported by the +// bridge yet. The family pass must reject the second pop with a diagnostic +// instead of emitting a silently wrong address (each pop would otherwise +// overwrite the tile's single recorded slot address) or crashing on the +// duplicated spec fields. +// +// RUN: env -u PTOAS_VPTO_BRIDGE_WHITELIST not pto-test-opt %s -pto-lower-pipe-family-ops 2>&1 | FileCheck %s + +// CHECK: error: VPTO pipe bridge supports at most one TPOP per declared tile; sequential rebind consumption is not supported yet + +module { + func.func @probe_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %addr0 = arith.constant 0 : i32 + %c0 = arith.constant 0 : index + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 2048, slot_num = 4, + flag_base = 8, nosplit = false} (%addr0 : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub1 = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + %v1 = arith.constant 1.0 : f32 + pto.store_scalar %v1, %ub1[%c0] : !pto.ptr, f32 + pto.tfree(%pipe : !pto.pipe) {split = 1} + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub2 = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + %v2 = arith.constant 2.0 : f32 + pto.store_scalar %v2, %ub2[%c0] : !pto.ptr, f32 + pto.tfree(%pipe : !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_spec_config_matrix.pto b/test/lit/vpto/vpto_bridge_spec_config_matrix.pto new file mode 100644 index 0000000000..6ac43c8c70 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_spec_config_matrix.pto @@ -0,0 +1,88 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Phase 2 configuration matrix: the pipe family pass must no longer reject +// non-default pipe configurations. The slot_size/slot_num/flag_base/nosplit +// attributes and the tile shapes flow verbatim into the collected bridge +// specialization as C++ template tokens. The whitelist carries declarative +// tmpl_map rows, which the pass consumes without changing the collected +// fields. Each function stores its specialization in a function attribute; +// the wrapper generation pass merges the identical per-function specs and +// renders the variant tokens into the wrapper source. + +// RUN: pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist-variant.yaml" | FileCheck %s --check-prefix=SPEC +// RUN: pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist-variant.yaml" -pto-emit-vpto-bridge-wrapper | FileCheck %s --check-prefix=WRAP + +// SPEC-LABEL: func.func private @variant_bridge_cube +// SPEC-SAME: pto.vpto.bridge.func_spec = { +// SPEC-SAME: entry.init = "pto_vpto_pipe_init" +// SPEC-SAME: pipe = "pto::TPipe<8, pto::Direction::DIR_C2V, 2048, 4, 2, true>" +// SPEC-SAME: producer_tile = "pto::Tile" +// SPEC: pto.bridge_call "pto_vpto_pipe_init"{{.*}}storage_size_callee = "pto_vpto_pipe_size" +// SPEC-SAME: -> !pto.pipe +// SPEC: pto.bridge_call "pto_vpto_pipe_push" +// SPEC-SAME: : !pto.pipe, i64 +// SPEC-NOT: pto.alloc_tile +// SPEC-NOT: pto.tpush + +// SPEC-LABEL: func.func private @variant_bridge_vector +// SPEC-SAME: pto.vpto.bridge.func_spec = { +// SPEC-SAME: consumer_tile = "pto::Tile" +// SPEC-SAME: pipe = "pto::TPipe<8, pto::Direction::DIR_C2V, 2048, 4, 2, true>" +// SPEC-SAME: split = "pto::TileSplitAxis::TILE_UP_DOWN" +// SPEC: pto.bridge_call "pto_vpto_pipe_init" +// SPEC: pto.bridge_call "pto_vpto_pipe_pop" +// SPEC-SAME: -> i64 +// SPEC: pto.bridge_inttoptr +// SPEC: pto.bridge_call "pto_vpto_pipe_free" +// SPEC-NOT: pto.declare_tile +// SPEC-NOT: pto.tpop +// SPEC-NOT: pto.tfree + +// WRAP: module attributes {{{.*}}pto.vpto.bridge.wrapper_source = " +// WRAP-SAME: using Pipe = pto::TPipe<8, pto::Direction::DIR_C2V, 2048, 4, 2, true>; +// WRAP-SAME: using ProducerTile = pto::Tile; +// WRAP-SAME: using ConsumerTile = pto::Tile; +// WRAP-NOT: pto.vpto.bridge.func_spec + +module { + func.func private @variant_bridge_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 2048, slot_num = 4, flag_base = 8, nosplit = true} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } + + func.func private @variant_bridge_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 2048, slot_num = 4, flag_base = 8, nosplit = true} (%addr : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + pto.tfree(%pipe : !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_spec_conflict_diag.pto b/test/lit/vpto/vpto_bridge_spec_conflict_diag.pto new file mode 100644 index 0000000000..4edddcb47e --- /dev/null +++ b/test/lit/vpto/vpto_bridge_spec_conflict_diag.pto @@ -0,0 +1,54 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Cross-function spec merging: identical pipe configurations deduplicate +// into a single module spec, but two functions initializing the pipe with +// different slot configurations are a conflict the wrapper generation pass +// must reject with a diagnostic (a single wrapper cannot serve two +// specializations). + +// RUN: not pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist.yaml" -pto-emit-vpto-bridge-wrapper 2>&1 | FileCheck %s + +// CHECK: error: VPTO bridge: conflicting bridge specialization across functions; the family configurations must be identical + +module { + func.func private @conflict_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } + + func.func private @conflict_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 256, slot_num = 16, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + pto.tfree(%pipe : !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_whitelist_default_channel_diag.pto b/test/lit/vpto/vpto_bridge_whitelist_default_channel_diag.pto new file mode 100644 index 0000000000..86220ffc2e --- /dev/null +++ b/test/lit/vpto/vpto_bridge_whitelist_default_channel_diag.pto @@ -0,0 +1,43 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// The `lowering` whitelist field defaults to the generic declarative +// channel, so an interface family that is mechanically mapped costs no pass +// code and no tag, while an entry owned by a dedicated family pass has to +// opt out with `lowering: custom`. +// +// MISSING covers the payoff of that polarity: an entry that forgets the +// opt-out is validated as declarative, and its unbound abi arguments are +// rejected at whitelist parse time with the missing tag named. Before the +// default was flipped this file parsed clean and only failed after lowering, +// blaming a family pass that never existed. +// +// LEGACY pins the vocabulary: `family`, the pre-flip spelling of the opt-out, +// is not silently reinterpreted. + +// RUN: not pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist-missing-custom-tag.yaml" 2>&1 | FileCheck %s --check-prefix=MISSING +// RUN: not pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist-legacy-family-lowering.yaml" 2>&1 | FileCheck %s --check-prefix=LEGACY + +// MISSING: VPTO bridge whitelist: declarative entry 'pto_vpto_pipe_init' has an abi argument without operand/role binding +// MISSING-SAME: vpto-bridge-whitelist-missing-custom-tag.yaml +// MISSING-SAME: (entries owned by a dedicated family pass must declare 'lowering: custom') + +// LEGACY: VPTO bridge whitelist: entry 'pto_vpto_pipe_init' declares unsupported lowering 'family' +// LEGACY-SAME: (supported: declarative, custom) + +module { + func.func private @default_channel_diag(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tfree(%pipe : !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_whitelist_residual_diag.pto b/test/lit/vpto/vpto_bridge_whitelist_residual_diag.pto new file mode 100644 index 0000000000..b2d602c213 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_whitelist_residual_diag.pto @@ -0,0 +1,35 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Whitelist-driven routing check: an op the whitelist routes to a wrapper +// entry but the family pass did not lower must be rejected with a +// diagnostic, not silently emitted as LLVM IR. + +// RUN: not pto-test-opt %s "-vpto-bridge-lowering=whitelist-path=%S/Inputs/vpto-bridge-whitelist.yaml" 2>&1 | FileCheck %s + +// CHECK: error: VPTO bridge: 'pto.initialize_l2l_pipe' is routed to wrapper entry 'pto_vpto_pipe_init' +// CHECK: error: VPTO bridge: 'pto.tpush' is routed to wrapper entry 'pto_vpto_pipe_push' + +module { + func.func private @unlowered_pipe_ops(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_whitelist_unknown_wrapper_diag.pto b/test/lit/vpto/vpto_bridge_whitelist_unknown_wrapper_diag.pto new file mode 100644 index 0000000000..d8b83a9591 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_whitelist_unknown_wrapper_diag.pto @@ -0,0 +1,38 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Wrapper-generator dispatch: the renderer is selected by the whitelist +// `wrapper` field, not by sniffing which spec keys were collected. Entries +// carrying `lowering: custom` name a wrapper that owns a dedicated +// renderer; naming one no renderer knows must be rejected with a +// diagnostic that names the custom wrapper and the available renderers, +// instead of falling through into another wrapper's renderer and failing +// on unrelated missing spec fields. + +// RUN: not pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist-unknown-wrapper.yaml" "-pto-emit-vpto-bridge-wrapper=whitelist-path=%S/Inputs/vpto-bridge-whitelist-unknown-wrapper.yaml" 2>&1 | FileCheck %s + +// CHECK: error: VPTO bridge: whitelist entries name the custom wrapper 'pipe_unfinished_port', which has no dedicated renderer in the bridge wrapper generator (available: pipe) + +module { + func.func private @unknown_wrapper(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_wrapper_source_c2v.pto b/test/lit/vpto/vpto_bridge_wrapper_source_c2v.pto new file mode 100644 index 0000000000..de166f4052 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_wrapper_source_c2v.pto @@ -0,0 +1,71 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Wrapper generation for a C2V pipe: the wrapper-gen pass renders the spec +// collected by the family pass into `pto.vpto.bridge.wrapper_source`, drops +// the spec attribute, and the rendered source instantiates the pipe/tile +// tokens from the IR configuration (not from any hardcoded default). The +// cube core produces (push) and the vector core consumes (pop/free), each +// guarded by its core macro. The StringAttr prints on a single line with +// \0A escapes, hence the single-line {{.*}} patterns. + +// RUN: pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist.yaml" -pto-emit-vpto-bridge-wrapper | FileCheck %s + +// CHECK-NOT: pto.vpto.bridge.spec +// CHECK: module attributes {{{.*}}pto.vpto.bridge.wrapper_source = " +// CHECK-NOT: pto.vpto.bridge.func_spec +// CHECK-SAME: using Pipe = pto::TPipe<0, pto::Direction::DIR_C2V, 1024, 8, 2, false>; +// CHECK-SAME: using ProducerTile = pto::Tile; +// CHECK-SAME: using ConsumerTile = pto::Tile; +// CHECK-SAME: extern \22C\22 [aicore] void pto_vpto_pipe_init(void *storage, uint32_t localBuffer) +// CHECK-SAME: extern \22C\22 [aicore] size_t pto_vpto_pipe_size() { return sizeof(Pipe); } +// CHECK-SAME: #ifdef __DAV_CUBE__ +// CHECK-SAME: pto::TPUSH(pipe, tile); +// CHECK-SAME: #endif +// CHECK-SAME: #ifdef __DAV_VEC__ +// CHECK-SAME: pto::TPOP(pipe, tile); +// CHECK-SAME: pipe_barrier(PIPE_ALL); +// CHECK-SAME: pto::TFREE(pipe); +// CHECK-SAME: #endif + +module { + func.func private @bridge_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %acc = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%acc, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } + + func.func private @bridge_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + pto.tfree(%pipe : !pto.pipe) {split = 1} + return + } +} diff --git a/test/lit/vpto/vpto_bridge_wrapper_source_v2c.pto b/test/lit/vpto/vpto_bridge_wrapper_source_v2c.pto new file mode 100644 index 0000000000..668fdf10f3 --- /dev/null +++ b/test/lit/vpto/vpto_bridge_wrapper_source_v2c.pto @@ -0,0 +1,65 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software; you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// You may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// https://www.huawei.com/ +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS +// OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A +// PARTICULAR PURPOSE. +// See LICENSE for more details. + +// Wrapper generation for a V2C pipe: the roles swap compared to C2V. The +// vector core produces (push) and the cube core consumes (pop/free), so the +// rendered source must place the push entry under __DAV_VEC__ and the +// pop/free entries under __DAV_CUBE__, with the producer/consumer tile +// typedefs bound accordingly. + +// RUN: pto-test-opt %s "-pto-lower-pipe-family-ops=whitelist-path=%S/Inputs/vpto-bridge-whitelist.yaml" -pto-emit-vpto-bridge-wrapper | FileCheck %s + +// CHECK: module attributes {{{.*}}pto.vpto.bridge.wrapper_source = " +// CHECK-NOT: pto.vpto.bridge.func_spec +// CHECK-SAME: using Pipe = pto::TPipe<0, pto::Direction::DIR_V2C, 1024, 8, 2, false>; +// CHECK-SAME: using ProducerTile = pto::Tile; +// CHECK-SAME: using ConsumerTile = pto::Tile; +// CHECK-SAME: #ifdef __DAV_CUBE__ +// CHECK-SAME: pto::TPOP(pipe, tile); +// CHECK-SAME: pto::TFREE(pipe); +// CHECK-SAME: #endif +// CHECK-SAME: #ifdef __DAV_VEC__ +// CHECK-SAME: pto::TPUSH(pipe, tile); +// CHECK-SAME: #endif + +module attributes {pto.target_arch = "a5"} { + func.func private @v2c_bridge_vector(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %tile = pto.alloc_tile addr = %c0 + : !pto.tile_buf + %pipe = pto.initialize_l2l_pipe{dir_mask = 2, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + pto.tpush(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + return + } + + func.func private @v2c_bridge_cube(%addr: i32) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %pipe = pto.initialize_l2l_pipe{dir_mask = 2, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%addr : i32) -> !pto.pipe + %tile = pto.declare_tile + -> !pto.tile_buf + pto.tpop(%tile, %pipe + : !pto.tile_buf, + !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + pto.tfree(%pipe : !pto.pipe) {split = 1} + return + } +} diff --git a/test/vpto/cases/kernels/cube-matmul-bridge/compare.py b/test/vpto/cases/kernels/cube-matmul-bridge/compare.py new file mode 100644 index 0000000000..c1391455e3 --- /dev/null +++ b/test/vpto/cases/kernels/cube-matmul-bridge/compare.py @@ -0,0 +1,46 @@ +#!/usr/bin/env python3 +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + + +import os +import sys + +import numpy as np + + +def compare_bin(golden_path: str, output_path: str) -> bool: + if not os.path.exists(golden_path) or not os.path.exists(output_path): + return False + golden = np.fromfile(golden_path, dtype=np.float32) + output = np.fromfile(output_path, dtype=np.float32) + if golden.shape != output.shape: + print(f"[ERROR] shape mismatch: {golden.shape} vs {output.shape}") + return False + if np.allclose(golden, output, atol=1e-2, rtol=1e-2): + return True + diff = np.where(np.abs(golden - output) > (1e-2 + 1e-2 * np.abs(golden)))[0] + idx = int(diff[0]) if diff.size else 0 + print(f"[ERROR] first mismatch at idx={idx}: golden={float(golden[idx])}, out={float(output[idx])}") + return False + + +def main() -> None: + strict = os.getenv("COMPARE_STRICT", "1") != "0" + ok = compare_bin("golden_v3.bin", "v3.bin") + if not ok: + if strict: + print("[ERROR] compare failed") + sys.exit(2) + print("[WARN] compare failed (non-gating)") + return + print("[INFO] compare passed") + + +if __name__ == "__main__": + main() diff --git a/test/vpto/cases/kernels/cube-matmul-bridge/golden.py b/test/vpto/cases/kernels/cube-matmul-bridge/golden.py new file mode 100644 index 0000000000..c9f56dae66 --- /dev/null +++ b/test/vpto/cases/kernels/cube-matmul-bridge/golden.py @@ -0,0 +1,43 @@ +#!/usr/bin/env python3 +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + + +import argparse +from pathlib import Path +import numpy as np + +M = 16 +N = 16 +K = 16 + + +def generate(output_dir: Path) -> None: + a = np.eye(M, K, dtype=np.float16) + b = (np.arange(K * N, dtype=np.float16) * np.float16(0.25)).reshape(K, N) + c = np.zeros((M, N), dtype=np.float32) + a_f32 = a.astype(np.float32, copy=False) + b_f32 = b.astype(np.float32, copy=False) + golden_c = a_f32 @ b_f32 + + output_dir.mkdir(parents=True, exist_ok=True) + a.reshape(-1).tofile(output_dir / "v1.bin") + b.reshape(-1).tofile(output_dir / "v2.bin") + c.reshape(-1).tofile(output_dir / "v3.bin") + golden_c.reshape(-1).tofile(output_dir / "golden_v3.bin") + + +def main() -> None: + parser = argparse.ArgumentParser() + parser.add_argument("--output-dir", type=Path, default=Path(".")) + args = parser.parse_args() + generate(args.output_dir) + + +if __name__ == "__main__": + main() diff --git a/test/vpto/cases/kernels/cube-matmul-bridge/kernel.pto b/test/vpto/cases/kernels/cube-matmul-bridge/kernel.pto new file mode 100644 index 0000000000..f29542d287 --- /dev/null +++ b/test/vpto/cases/kernels/cube-matmul-bridge/kernel.pto @@ -0,0 +1,74 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Phase 3 second interface family e2e: the cube-side tmatmul is routed to +// the generated C++ bridge wrapper (TMATMUL) instead of the regular +// tile-op expansion to pto.mad. The three operand tiles carry planned L0 +// addresses matching the MTE load/store targets; the result leaves L0C via +// the fixpipe path. A successful compare proves the bridge path computes +// the same result as the golden matmul. +module attributes {pto.target_arch = "a5"} { + func.func @cube_matmul_bridge_kernel(%a: !pto.ptr, + %b: !pto.ptr, + %out: !pto.ptr) attributes {pto.entry} { + func.call @cube_matmul_bridge_cube(%a, %b, %out) + : (!pto.ptr, !pto.ptr, !pto.ptr) -> () + return + } + + func.func private @cube_matmul_bridge_cube( + %a: !pto.ptr, %b: !pto.ptr, %out: !pto.ptr) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0 = arith.constant 0 : i64 + %c1 = arith.constant 1 : i64 + %c16 = arith.constant 16 : i64 + %c512 = arith.constant 512 : i64 + + %l1a = pto.castptr %c0 : i64 -> !pto.ptr + %l1b = pto.castptr %c512 : i64 -> !pto.ptr + + %lhs = pto.alloc_tile + : !pto.tile_buf + %rhs = pto.alloc_tile + : !pto.tile_buf + %acc = pto.alloc_tile + : !pto.tile_buf + %l0a = pto.tile_buf_addr %lhs + : !pto.tile_buf + -> !pto.ptr + %l0b = pto.tile_buf_addr %rhs + : !pto.tile_buf + -> !pto.ptr + %l0c = pto.tile_buf_addr %acc + : !pto.tile_buf + -> !pto.ptr + + pto.mte_gm_l1 %a, %l1a, %c512 nburst(%c1, %c0, %c0) loop(%c1, %c0, %c0) + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, loop i64, i64, i64 + pto.mte_gm_l1 %b, %l1b, %c512 nburst(%c1, %c0, %c0) loop(%c1, %c0, %c0) + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, loop i64, i64, i64 + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.mte_l1_l0a %l1a, %l0a, %c16, %c16, %c0, %c0 + : !pto.ptr, !pto.ptr, i64, i64, i64, i64 + pto.mte_l1_l0b %l1b, %l0b, %c16, %c16, %c0, %c0 {transpose = true} + : !pto.ptr, !pto.ptr, i64, i64, i64, i64 + pto.set_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.tmatmul ins(%lhs, %rhs + : !pto.tile_buf, + !pto.tile_buf) + outs(%acc : !pto.tile_buf) + pto.set_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.wait_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.mte_l0c_gm %l0c, %out, %c16, %c16, %c16, %c16, %c0, %c0, nz2nd + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + pto.barrier #pto.pipe + return + } +} diff --git a/test/vpto/cases/kernels/cube-matmul-bridge/launch.cpp b/test/vpto/cases/kernels/cube-matmul-bridge/launch.cpp new file mode 100644 index 0000000000..a65e18a235 --- /dev/null +++ b/test/vpto/cases/kernels/cube-matmul-bridge/launch.cpp @@ -0,0 +1,49 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +#ifndef __VEC_SCOPE__ +#define __VEC_SCOPE__ +#endif + +#if defined(__CCE_AICORE__) && defined(__NPU_ARCH__) && (__NPU_ARCH__ == 2201) +typedef struct { unsigned char v; } hifloat8_t; +typedef struct { unsigned char v; } float8_e4m3_t; +typedef struct { unsigned char v; } float8_e5m2_t; +typedef struct { unsigned char v; } float8_e8m0_t; +typedef struct { unsigned char v; } float4_e1m2x2_t; +typedef struct { unsigned char v; } float4_e2m1x2_t; +#endif + +#include + +#if defined(__CCE_AICORE__) && defined(PTOAS_ENABLE_CCE_PRINT) +#include +#endif + +#if !defined(__CCE_AICORE__) && !defined(TMRGSORT_HPP) +struct MrgSortExecutedNumList { + uint16_t mrgSortList0; + uint16_t mrgSortList1; + uint16_t mrgSortList2; + uint16_t mrgSortList3; +}; +#endif + +#ifndef __CPU_SIM +#include "acl/acl.h" +#endif + +extern "C" __global__ [aicore] void cube_matmul_bridge_kernel(__gm__ __fp16 *a, + __gm__ __fp16 *b, + __gm__ float *c); + +void LaunchCube_matmul_bridge_kernel(__fp16 *a, __fp16 *b, float *c, void *stream) { + cube_matmul_bridge_kernel<<<1, nullptr, stream>>>((__gm__ __fp16 *)a, + (__gm__ __fp16 *)b, + (__gm__ float *)c); +} diff --git a/test/vpto/cases/kernels/cube-matmul-bridge/main.cpp b/test/vpto/cases/kernels/cube-matmul-bridge/main.cpp new file mode 100644 index 0000000000..25d0182781 --- /dev/null +++ b/test/vpto/cases/kernels/cube-matmul-bridge/main.cpp @@ -0,0 +1,127 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +#include "test_common.h" +#include "acl/acl.h" +#include +#include +#include + +using namespace PtoTestCommon; + +#ifndef TMRGSORT_HPP +struct MrgSortExecutedNumList { + uint16_t mrgSortList0; + uint16_t mrgSortList1; + uint16_t mrgSortList2; + uint16_t mrgSortList3; +}; +#endif + +#define ACL_CHECK(expr) \ + do { \ + const aclError _ret = (expr); \ + if (_ret != ACL_SUCCESS) { \ + std::fprintf(stderr, "[ERROR] %s failed: %d (%s:%d)\n", #expr, \ + (int)_ret, __FILE__, __LINE__); \ + const char *_recent = aclGetRecentErrMsg(); \ + if (_recent != nullptr && _recent[0] != '\0') \ + std::fprintf(stderr, "[ERROR] RecentErrMsg: %s\n", _recent); \ + rc = 1; \ + goto cleanup; \ + } \ + } while (0) + +#define FILE_CHECK(expr, path) \ + do { \ + if (!(expr)) { \ + std::fprintf(stderr, "[ERROR] file operation failed: %s (%s:%d)\n", \ + path, __FILE__, __LINE__); \ + rc = 1; \ + goto cleanup; \ + } \ + } while (0) + +void LaunchCube_matmul_bridge_kernel(__fp16 *a, __fp16 *b, float *c, void *stream); + +int main() { + constexpr size_t kM = 16; + constexpr size_t kN = 16; + constexpr size_t kK = 16; + constexpr size_t aElem = kM * kK; + constexpr size_t bElem = kK * kN; + constexpr size_t cElem = kM * kN; + + constexpr size_t aSize = aElem * sizeof(__fp16); + constexpr size_t bSize = bElem * sizeof(__fp16); + constexpr size_t cSize = cElem * sizeof(float); + + __fp16 *aHost = nullptr; + __fp16 *bHost = nullptr; + float *cHost = nullptr; + __fp16 *aDevice = nullptr; + __fp16 *bDevice = nullptr; + float *cDevice = nullptr; + + int rc = 0; + bool aclInited = false; + bool deviceSet = false; + int deviceId = 0; + aclrtStream stream = nullptr; + size_t inputSize = 0; + + ACL_CHECK(aclInit(nullptr)); + aclInited = true; + if (const char *envDevice = std::getenv("ACL_DEVICE_ID")) + deviceId = std::atoi(envDevice); + ACL_CHECK(aclrtSetDevice(deviceId)); + deviceSet = true; + ACL_CHECK(aclrtCreateStream(&stream)); + + ACL_CHECK(aclrtMallocHost((void **)(&aHost), aSize)); + ACL_CHECK(aclrtMallocHost((void **)(&bHost), bSize)); + ACL_CHECK(aclrtMallocHost((void **)(&cHost), cSize)); + ACL_CHECK(aclrtMalloc((void **)&aDevice, aSize, ACL_MEM_MALLOC_HUGE_FIRST)); + ACL_CHECK(aclrtMalloc((void **)&bDevice, bSize, ACL_MEM_MALLOC_HUGE_FIRST)); + ACL_CHECK(aclrtMalloc((void **)&cDevice, cSize, ACL_MEM_MALLOC_HUGE_FIRST)); + + inputSize = aSize; + FILE_CHECK(ReadFile("./v1.bin", inputSize, aHost, aSize) && inputSize == aSize, + "./v1.bin"); + inputSize = bSize; + FILE_CHECK(ReadFile("./v2.bin", inputSize, bHost, bSize) && inputSize == bSize, + "./v2.bin"); + inputSize = cSize; + FILE_CHECK(ReadFile("./v3.bin", inputSize, cHost, cSize) && inputSize == cSize, + "./v3.bin"); + + ACL_CHECK(aclrtMemcpy(aDevice, aSize, aHost, aSize, ACL_MEMCPY_HOST_TO_DEVICE)); + ACL_CHECK(aclrtMemcpy(bDevice, bSize, bHost, bSize, ACL_MEMCPY_HOST_TO_DEVICE)); + ACL_CHECK(aclrtMemcpy(cDevice, cSize, cHost, cSize, ACL_MEMCPY_HOST_TO_DEVICE)); + + LaunchCube_matmul_bridge_kernel(aDevice, bDevice, cDevice, stream); + ACL_CHECK(aclrtSynchronizeStream(stream)); + + ACL_CHECK(aclrtMemcpy(cHost, cSize, cDevice, cSize, ACL_MEMCPY_DEVICE_TO_HOST)); + FILE_CHECK(WriteFile("./v3.bin", cHost, cSize), "./v3.bin"); + +cleanup: + aclrtFree(aDevice); + aclrtFree(bDevice); + aclrtFree(cDevice); + aclrtFreeHost(aHost); + aclrtFreeHost(bHost); + aclrtFreeHost(cHost); + if (stream != nullptr) + aclrtDestroyStream(stream); + if (deviceSet) + aclrtResetDevice(deviceId); + if (aclInited) + aclFinalize(); + return rc; +} diff --git a/test/vpto/cases/kernels/fifo-tile-data-consume/compare.py b/test/vpto/cases/kernels/fifo-tile-data-consume/compare.py new file mode 100644 index 0000000000..c1391455e3 --- /dev/null +++ b/test/vpto/cases/kernels/fifo-tile-data-consume/compare.py @@ -0,0 +1,46 @@ +#!/usr/bin/env python3 +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + + +import os +import sys + +import numpy as np + + +def compare_bin(golden_path: str, output_path: str) -> bool: + if not os.path.exists(golden_path) or not os.path.exists(output_path): + return False + golden = np.fromfile(golden_path, dtype=np.float32) + output = np.fromfile(output_path, dtype=np.float32) + if golden.shape != output.shape: + print(f"[ERROR] shape mismatch: {golden.shape} vs {output.shape}") + return False + if np.allclose(golden, output, atol=1e-2, rtol=1e-2): + return True + diff = np.where(np.abs(golden - output) > (1e-2 + 1e-2 * np.abs(golden)))[0] + idx = int(diff[0]) if diff.size else 0 + print(f"[ERROR] first mismatch at idx={idx}: golden={float(golden[idx])}, out={float(output[idx])}") + return False + + +def main() -> None: + strict = os.getenv("COMPARE_STRICT", "1") != "0" + ok = compare_bin("golden_v3.bin", "v3.bin") + if not ok: + if strict: + print("[ERROR] compare failed") + sys.exit(2) + print("[WARN] compare failed (non-gating)") + return + print("[INFO] compare passed") + + +if __name__ == "__main__": + main() diff --git a/test/vpto/cases/kernels/fifo-tile-data-consume/golden.py b/test/vpto/cases/kernels/fifo-tile-data-consume/golden.py new file mode 100644 index 0000000000..682fea177d --- /dev/null +++ b/test/vpto/cases/kernels/fifo-tile-data-consume/golden.py @@ -0,0 +1,43 @@ +#!/usr/bin/env python3 +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + + +import argparse +from pathlib import Path +import numpy as np + +M = 16 +N = 16 +K = 16 + + +def generate(output_dir: Path) -> None: + a = np.eye(M, K, dtype=np.float16) + b = np.arange(K * N, dtype=np.float16).reshape(K, N) + c = np.zeros((8, N), dtype=np.float32) + a_f32 = a.astype(np.float32, copy=False) + b_f32 = b.astype(np.float32, copy=False) + golden_c = (a_f32 @ b_f32)[:8, :] + + output_dir.mkdir(parents=True, exist_ok=True) + a.reshape(-1).tofile(output_dir / "v1.bin") + b.reshape(-1).tofile(output_dir / "v2.bin") + c.reshape(-1).tofile(output_dir / "v3.bin") + golden_c.reshape(-1).tofile(output_dir / "golden_v3.bin") + + +def main() -> None: + parser = argparse.ArgumentParser() + parser.add_argument("--output-dir", type=Path, default=Path(".")) + args = parser.parse_args() + generate(args.output_dir) + + +if __name__ == "__main__": + main() diff --git a/test/vpto/cases/kernels/fifo-tile-data-consume/kernel.pto b/test/vpto/cases/kernels/fifo-tile-data-consume/kernel.pto new file mode 100644 index 0000000000..5734daec57 --- /dev/null +++ b/test/vpto/cases/kernels/fifo-tile-data-consume/kernel.pto @@ -0,0 +1,87 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// The output has no Cube-to-GM bypass. It is written only from the Vec Tile +// address returned by TPOP, so compare success proves FIFO Tile consumption. +module attributes {pto.target_arch = "a5"} { + func.func @fifo_tile_data_consume_kernel(%a: !pto.ptr, + %b: !pto.ptr, + %out: !pto.ptr) attributes {pto.entry} { + func.call @fifo_tile_data_consume_cube(%a, %b, %out) + : (!pto.ptr, !pto.ptr, !pto.ptr) -> () + func.call @fifo_tile_data_consume_vector(%a, %b, %out) + : (!pto.ptr, !pto.ptr, !pto.ptr) -> () + return + } + + func.func private @fifo_tile_data_consume_cube( + %a: !pto.ptr, %b: !pto.ptr, %out: !pto.ptr) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0_i32 = arith.constant 0 : i32 + %c0 = arith.constant 0 : i64 + %c1 = arith.constant 1 : i64 + %c16 = arith.constant 16 : i64 + %c512 = arith.constant 512 : i64 + %fifo_addr = pto.import_reserved_buffer {name = "c2v_fifo", peer_func = @fifo_tile_data_consume_vector} -> i32 + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%fifo_addr : i32) -> !pto.pipe + + %l1a = pto.castptr %c0 : i64 -> !pto.ptr + %l1b = pto.castptr %c512 : i64 -> !pto.ptr + %l0a = pto.castptr %c0 : i64 -> !pto.ptr + %l0b = pto.castptr %c0 : i64 -> !pto.ptr + %acc = pto.alloc_tile + : !pto.tile_buf + %l0c = pto.tile_buf_addr %acc + : !pto.tile_buf + -> !pto.ptr + + pto.mte_gm_l1 %a, %l1a, %c512 nburst(%c1, %c0, %c0) loop(%c1, %c0, %c0) + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, loop i64, i64, i64 + pto.mte_gm_l1 %b, %l1b, %c512 nburst(%c1, %c0, %c0) loop(%c1, %c0, %c0) + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, loop i64, i64, i64 + pto.set_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE2", "PIPE_MTE1", "EVENT_ID0"] + pto.mte_l1_l0a %l1a, %l0a, %c16, %c16, %c0, %c0 + : !pto.ptr, !pto.ptr, i64, i64, i64, i64 + pto.mte_l1_l0b %l1b, %l0b, %c16, %c16, %c0, %c0 {transpose = true} + : !pto.ptr, !pto.ptr, i64, i64, i64, i64 + pto.set_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE1", "PIPE_M", "EVENT_ID0"] + pto.mad %l0a, %l0b, %l0c, %c16, %c16, %c16 + : !pto.ptr, !pto.ptr, !pto.ptr, i64, i64, i64 + pto.set_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.wait_flag["PIPE_M", "PIPE_FIX", "EVENT_ID1"] + pto.tpush(%acc, %pipe : !pto.tile_buf, !pto.pipe) {split = 1} + return + } + + func.func private @fifo_tile_data_consume_vector( + %a: !pto.ptr, %b: !pto.ptr, %out: !pto.ptr) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %c0_i32 = arith.constant 0 : i32 + %c0 = arith.constant 0 : i64 + %c1 = arith.constant 1 : i64 + %c512 = arith.constant 512 : i64 + %fifo_addr = pto.reserve_buffer {name = "c2v_fifo", size = 8192, location = #pto.address_space, auto = true} -> i32 + %pipe = pto.initialize_l2l_pipe{dir_mask = 1, slot_size = 1024, slot_num = 8, flag_base = 0, nosplit = false} (%fifo_addr : i32) -> !pto.pipe + %tile = pto.declare_tile -> !pto.tile_buf + pto.tpop(%tile, %pipe : !pto.tile_buf, !pto.pipe) {split = 1} + %ub = pto.tile_buf_addr %tile + : !pto.tile_buf + -> !pto.ptr + %subblock = pto.get_subblock_idx + %is0 = arith.cmpi eq, %subblock, %c0 : i64 + scf.if %is0 { + pto.mte_ub_gm %ub, %out, %c512 nburst(%c1, %c0, %c0) + : !pto.ptr, !pto.ptr, i64, i64, i64, i64 + } + pto.tfree(%pipe : !pto.pipe) {split = 1} + pto.barrier #pto.pipe + return + } +} diff --git a/test/vpto/cases/kernels/fifo-tile-data-consume/launch.cpp b/test/vpto/cases/kernels/fifo-tile-data-consume/launch.cpp new file mode 100644 index 0000000000..935312e25b --- /dev/null +++ b/test/vpto/cases/kernels/fifo-tile-data-consume/launch.cpp @@ -0,0 +1,49 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +#ifndef __VEC_SCOPE__ +#define __VEC_SCOPE__ +#endif + +#if defined(__CCE_AICORE__) && defined(__NPU_ARCH__) && (__NPU_ARCH__ == 2201) +typedef struct { unsigned char v; } hifloat8_t; +typedef struct { unsigned char v; } float8_e4m3_t; +typedef struct { unsigned char v; } float8_e5m2_t; +typedef struct { unsigned char v; } float8_e8m0_t; +typedef struct { unsigned char v; } float4_e1m2x2_t; +typedef struct { unsigned char v; } float4_e2m1x2_t; +#endif + +#include + +#if defined(__CCE_AICORE__) && defined(PTOAS_ENABLE_CCE_PRINT) +#include +#endif + +#if !defined(__CCE_AICORE__) && !defined(TMRGSORT_HPP) +struct MrgSortExecutedNumList { + uint16_t mrgSortList0; + uint16_t mrgSortList1; + uint16_t mrgSortList2; + uint16_t mrgSortList3; +}; +#endif + +#ifndef __CPU_SIM +#include "acl/acl.h" +#endif + +extern "C" __global__ [aicore] void fifo_tile_data_consume_kernel(__gm__ __fp16 *a, + __gm__ __fp16 *b, + __gm__ float *c); + +void LaunchFifo_tile_data_consume_kernel(__fp16 *a, __fp16 *b, float *c, void *stream) { + fifo_tile_data_consume_kernel<<<1, nullptr, stream>>>((__gm__ __fp16 *)a, + (__gm__ __fp16 *)b, + (__gm__ float *)c); +} diff --git a/test/vpto/cases/kernels/fifo-tile-data-consume/main.cpp b/test/vpto/cases/kernels/fifo-tile-data-consume/main.cpp new file mode 100644 index 0000000000..a715d6c966 --- /dev/null +++ b/test/vpto/cases/kernels/fifo-tile-data-consume/main.cpp @@ -0,0 +1,127 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +#include "test_common.h" +#include "acl/acl.h" +#include +#include +#include + +using namespace PtoTestCommon; + +#ifndef TMRGSORT_HPP +struct MrgSortExecutedNumList { + uint16_t mrgSortList0; + uint16_t mrgSortList1; + uint16_t mrgSortList2; + uint16_t mrgSortList3; +}; +#endif + +#define ACL_CHECK(expr) \ + do { \ + const aclError _ret = (expr); \ + if (_ret != ACL_SUCCESS) { \ + std::fprintf(stderr, "[ERROR] %s failed: %d (%s:%d)\n", #expr, \ + (int)_ret, __FILE__, __LINE__); \ + const char *_recent = aclGetRecentErrMsg(); \ + if (_recent != nullptr && _recent[0] != '\0') \ + std::fprintf(stderr, "[ERROR] RecentErrMsg: %s\n", _recent); \ + rc = 1; \ + goto cleanup; \ + } \ + } while (0) + +#define FILE_CHECK(expr, path) \ + do { \ + if (!(expr)) { \ + std::fprintf(stderr, "[ERROR] file operation failed: %s (%s:%d)\n", \ + path, __FILE__, __LINE__); \ + rc = 1; \ + goto cleanup; \ + } \ + } while (0) + +void LaunchFifo_tile_data_consume_kernel(__fp16 *a, __fp16 *b, float *c, void *stream); + +int main() { + constexpr size_t kM = 16; + constexpr size_t kN = 16; + constexpr size_t kK = 16; + constexpr size_t aElem = kM * kK; + constexpr size_t bElem = kK * kN; + constexpr size_t cElem = 8 * kN; + + constexpr size_t aSize = aElem * sizeof(__fp16); + constexpr size_t bSize = bElem * sizeof(__fp16); + constexpr size_t cSize = cElem * sizeof(float); + + __fp16 *aHost = nullptr; + __fp16 *bHost = nullptr; + float *cHost = nullptr; + __fp16 *aDevice = nullptr; + __fp16 *bDevice = nullptr; + float *cDevice = nullptr; + + int rc = 0; + bool aclInited = false; + bool deviceSet = false; + int deviceId = 0; + aclrtStream stream = nullptr; + size_t inputSize = 0; + + ACL_CHECK(aclInit(nullptr)); + aclInited = true; + if (const char *envDevice = std::getenv("ACL_DEVICE_ID")) + deviceId = std::atoi(envDevice); + ACL_CHECK(aclrtSetDevice(deviceId)); + deviceSet = true; + ACL_CHECK(aclrtCreateStream(&stream)); + + ACL_CHECK(aclrtMallocHost((void **)(&aHost), aSize)); + ACL_CHECK(aclrtMallocHost((void **)(&bHost), bSize)); + ACL_CHECK(aclrtMallocHost((void **)(&cHost), cSize)); + ACL_CHECK(aclrtMalloc((void **)&aDevice, aSize, ACL_MEM_MALLOC_HUGE_FIRST)); + ACL_CHECK(aclrtMalloc((void **)&bDevice, bSize, ACL_MEM_MALLOC_HUGE_FIRST)); + ACL_CHECK(aclrtMalloc((void **)&cDevice, cSize, ACL_MEM_MALLOC_HUGE_FIRST)); + + inputSize = aSize; + FILE_CHECK(ReadFile("./v1.bin", inputSize, aHost, aSize) && inputSize == aSize, + "./v1.bin"); + inputSize = bSize; + FILE_CHECK(ReadFile("./v2.bin", inputSize, bHost, bSize) && inputSize == bSize, + "./v2.bin"); + inputSize = cSize; + FILE_CHECK(ReadFile("./v3.bin", inputSize, cHost, cSize) && inputSize == cSize, + "./v3.bin"); + + ACL_CHECK(aclrtMemcpy(aDevice, aSize, aHost, aSize, ACL_MEMCPY_HOST_TO_DEVICE)); + ACL_CHECK(aclrtMemcpy(bDevice, bSize, bHost, bSize, ACL_MEMCPY_HOST_TO_DEVICE)); + ACL_CHECK(aclrtMemcpy(cDevice, cSize, cHost, cSize, ACL_MEMCPY_HOST_TO_DEVICE)); + + LaunchFifo_tile_data_consume_kernel(aDevice, bDevice, cDevice, stream); + ACL_CHECK(aclrtSynchronizeStream(stream)); + + ACL_CHECK(aclrtMemcpy(cHost, cSize, cDevice, cSize, ACL_MEMCPY_DEVICE_TO_HOST)); + FILE_CHECK(WriteFile("./v3.bin", cHost, cSize), "./v3.bin"); + +cleanup: + aclrtFree(aDevice); + aclrtFree(bDevice); + aclrtFree(cDevice); + aclrtFreeHost(aHost); + aclrtFreeHost(bHost); + aclrtFreeHost(cHost); + if (stream != nullptr) + aclrtDestroyStream(stream); + if (deviceSet) + aclrtResetDevice(deviceId); + if (aclInited) + aclFinalize(); + return rc; +} diff --git a/test/vpto/scripts/run_host_vpto_validation.sh b/test/vpto/scripts/run_host_vpto_validation.sh index 06331a012a..583e96b00f 100755 --- a/test/vpto/scripts/run_host_vpto_validation.sh +++ b/test/vpto/scripts/run_host_vpto_validation.sh @@ -362,9 +362,13 @@ build_one_impl() { read -r -a ptoas_args <<< "${PTOAS_FLAGS}" fi + # The bridge wrapper source is generated and compiled by ptoas itself; + # routing uses the built-in default whitelist unless the caller exports + # PTOAS_VPTO_BRIDGE_WHITELIST or the case passes whitelist-path via + # ptoas.flags. log "[$case_name] step 1/4: emit kernel fatobj" "${PTOAS_BIN}" "${ptoas_args[@]}" \ - "${case_dir}/kernel.pto" -o "${kernel_fatobj}" + "${case_dir}/kernel.pto" -o "${kernel_fatobj}" log "[$case_name] step 2/4: build launch object" build_launch_object "${case_dir}" "${launch_obj}" diff --git a/tools/ptoas/ObjectEmission.cpp b/tools/ptoas/ObjectEmission.cpp index cce84cbefa..bb986ff5ee 100644 --- a/tools/ptoas/ObjectEmission.cpp +++ b/tools/ptoas/ObjectEmission.cpp @@ -264,6 +264,23 @@ discoverCppIncludeDirs(llvm::StringRef ascendHome, ptoIsaPath = *env; } + if (ptoIsaPath.empty()) { + // Probe common development checkouts so the bridge wrapper and C++ + // device emission find pto/pto-inst.hpp without extra environment setup. + if (auto home = llvm::sys::Process::GetEnv("HOME")) { + std::string candidates[] = { + joinPath(*home, "pto-isa"), + joinPath(joinPath(*home, "llvm-workspace"), "pto-isa"), + }; + for (const std::string &candidate : candidates) { + if (llvm::sys::fs::is_directory(candidate)) { + ptoIsaPath = candidate; + break; + } + } + } + } + addPTOISAIncludeDirs(includeDirs, ptoIsaPath); addExistingIncludeDir(includeDirs, joinPath(ascendHome, "include")); std::string driverPath = @@ -281,6 +298,24 @@ discoverCppIncludeDirs(llvm::StringRef ascendHome, return includeDirs; } +// Merges externally compiled bridge bitcode (e.g. PTO-ISA template +// instantiation wrappers) into the VPTO device LLVM IR before Bisheng +// compiles the device object. Both inputs are bitcode/IR of the same LLVM +// version, linked with the llvm-link shipped with the Bisheng toolchain. +static bool linkDeviceLLVMBitcode(llvm::StringRef llPath, + llvm::StringRef bridgePath, + llvm::StringRef linkedPath, + const mlir::pto::CANNToolchain &toolchain, + llvm::StringRef stderrPath, + llvm::raw_ostream &diagOS) { + std::string llvmLinkPath = + joinPath(toolchain.bishengCompilerBinDirPath, "llvm-link"); + llvm::SmallVector args = { + llvmLinkPath, llPath.str(), bridgePath.str(), "-o", linkedPath.str()}; + return runCommandWithStderr(llvmLinkPath, args, stderrPath, diagOS, + "device LLVM bridge link"); +} + static bool compileDeviceLLVMToObject(llvm::StringRef llPath, llvm::StringRef outObjPath, llvm::StringRef targetCPU, @@ -325,6 +360,48 @@ static std::string resolveTargetCPU(llvm::Module &module, return getTargetCPU(fallback).str(); } +// Compiles the generated bridge wrapper C++ source to device bitcode for +// one core kind. The wrapper separates the core entries with __DAV_CUBE__ / +// __DAV_VEC__ guards, so compiling the same source once per target yields +// exactly the entries of that core. The command mirrors the former +// hand-written wrapper build (bisheng -O2 -std=c++17 -c -emit-llvm -xcce +// --cce-aicore-only -DREGISTER_BASE). +static bool compileBridgeWrapperToBitcode( + llvm::StringRef wrapperSource, + mlir::pto::ObjectEmissionDeviceTarget target, + const mlir::pto::CANNToolchain &toolchain, + mlir::pto::TempFileRegistry &tempFiles, llvm::StringRef stderrPath, + std::string &outBitcodePath, llvm::raw_ostream &diagOS) { + std::string sourcePath; + if (failed(tempFiles.create("ptoas-vpto-bridge-wrapper", ".cpp", sourcePath, + diagOS))) + return false; + if (!writeTextFile(sourcePath, wrapperSource, diagOS)) + return false; + if (failed(tempFiles.create("ptoas-vpto-bridge-wrapper", ".bc", + outBitcodePath, diagOS))) + return false; + llvm::SmallVector args = { + toolchain.bishengPath, + "-O2", + "-std=c++17", + "-c", + "-emit-llvm", + "-xcce", + "--cce-aicore-only", + std::string("--cce-aicore-arch=") + getTargetCPU(target).str(), + "-DREGISTER_BASE", + }; + for (const std::string &includeDir : toolchain.cppIncludeDirs) { + args.push_back("-I" + includeDir); + } + args.push_back(sourcePath); + args.push_back("-o"); + args.push_back(outBitcodePath); + return runCommandWithStderr(toolchain.bishengPath, args, stderrPath, diagOS, + "bridge wrapper bitcode compilation"); +} + class VPTOFatobjArtifacts { public: explicit VPTOFatobjArtifacts(mlir::pto::TempFileRegistry &tempFiles) @@ -349,6 +426,7 @@ class VPTOFatobjArtifacts { bool emitCubeObject(llvm::Module *module, const mlir::pto::CANNToolchain &toolchain, + llvm::StringRef bridgeBitcodePath, llvm::raw_ostream &diagOS) { if (!module) { return true; @@ -360,11 +438,13 @@ class VPTOFatobjArtifacts { return false; } return succeeded(mlir::pto::emitVPTOCubeDeviceObject( - *module, cubeLLPath, cubeObjPath, toolchain, stderrPath, diagOS)); + *module, cubeLLPath, cubeObjPath, toolchain, bridgeBitcodePath, + stderrPath, diagOS)); } bool emitVectorObject(llvm::Module *module, const mlir::pto::CANNToolchain &toolchain, + llvm::StringRef bridgeBitcodePath, mlir::pto::VFSIMTSizeFixMode vfsimtSizeFixMode, llvm::raw_ostream &diagOS) { if (!module) { @@ -379,8 +459,8 @@ class VPTOFatobjArtifacts { return false; } if (failed(mlir::pto::emitVPTOVectorDeviceObject( - *module, vectorLLPath, rawVectorObjPath, toolchain, stderrPath, - diagOS))) { + *module, vectorLLPath, rawVectorObjPath, toolchain, + bridgeBitcodePath, stderrPath, diagOS))) { return false; } if (vfsimtSizeFixMode == mlir::pto::VFSIMTSizeFixMode::Off) { @@ -404,6 +484,42 @@ class VPTOFatobjArtifacts { return true; } + bool compileBridgeWrapper(llvm::StringRef wrapperSource, + mlir::pto::ObjectEmissionDeviceTarget target, + const mlir::pto::CANNToolchain &toolchain, + std::string &outBitcodePath, + llvm::raw_ostream &diagOS) { + if (wrapperSource.empty()) { + return true; + } + return compileBridgeWrapperToBitcode(wrapperSource, target, toolchain, + tempFiles, stderrPath, + outBitcodePath, diagOS); + } + + // Compiles the generated bridge wrapper source to bitcode for each device + // module present, once per core kind (the wrapper guards its entries with + // __DAV_CUBE__ / __DAV_VEC__). An empty wrapper source compiles nothing. + bool compileBridgeWrappers(llvm::Module *cubeModule, + llvm::Module *vectorModule, + llvm::StringRef wrapperSource, + const mlir::pto::CANNToolchain &toolchain, + std::string &cubeBitcodePath, + std::string &vectorBitcodePath, + llvm::raw_ostream &diagOS) { + if (cubeModule && + !compileBridgeWrapper(wrapperSource, + mlir::pto::ObjectEmissionDeviceTarget::Cube, + toolchain, cubeBitcodePath, diagOS)) + return false; + if (vectorModule && + !compileBridgeWrapper(wrapperSource, + mlir::pto::ObjectEmissionDeviceTarget::Vector, + toolchain, vectorBitcodePath, diagOS)) + return false; + return true; + } + bool mergeDeviceObjects(const mlir::pto::CANNToolchain &toolchain, llvm::raw_ostream &diagOS) { llvm::SmallVector deviceObjPaths; @@ -580,6 +696,27 @@ static bool compileDeviceLLVMToObject(llvm::StringRef llPath, "device LLVM compilation", llPath); } +// Compiles the written device LLVM IR to an object, linking externally +// compiled bridge bitcode into the IR first when provided. The linked IR is +// materialized next to the object file and compiled instead of the original +// IR, so the original .ll stays a faithful dump of the module. +static bool compileDeviceLLVMToObjectWithBridgeLink( + llvm::StringRef llPath, llvm::StringRef bridgeBitcodePath, + llvm::StringRef outObjPath, llvm::StringRef targetCPU, + const mlir::pto::CANNToolchain &toolchain, llvm::StringRef stderrPath, + llvm::raw_ostream &diagOS) { + std::string compileInput = llPath.str(); + if (!bridgeBitcodePath.empty()) { + std::string linkedPath = (outObjPath + ".linked.bc").str(); + if (!linkDeviceLLVMBitcode(llPath, bridgeBitcodePath, linkedPath, + toolchain, stderrPath, diagOS)) + return false; + compileInput = linkedPath; + } + return compileDeviceLLVMToObject(compileInput, outObjPath, targetCPU, + toolchain.bishengPath, stderrPath, diagOS); +} + static bool compileCppDeviceSourceToObject( llvm::StringRef cppPath, llvm::StringRef outObjPath, llvm::StringRef targetCPU, const mlir::pto::CANNToolchain &toolchain, @@ -1074,8 +1211,8 @@ static mlir::LogicalResult applyVPTOLLVMABINames(llvm::Module &module, mlir::LogicalResult mlir::pto::emitVPTOVectorDeviceObject( llvm::Module &module, llvm::StringRef llPath, llvm::StringRef outObjPath, - const CANNToolchain &toolchain, llvm::StringRef stderrPath, - llvm::raw_ostream &diagOS) { + const CANNToolchain &toolchain, llvm::StringRef bridgeBitcodePath, + llvm::StringRef stderrPath, llvm::raw_ostream &diagOS) { if (failed(applyVPTOLLVMABINames( module, toolchain.vptoPublicABISuffix(ObjectEmissionDeviceTarget::Vector), @@ -1085,18 +1222,18 @@ mlir::LogicalResult mlir::pto::emitVPTOVectorDeviceObject( if (failed(writeLLVMModule(module, llPath, diagOS))) { return failure(); } - return compileDeviceLLVMToObject(llPath, outObjPath, - resolveTargetCPU(module, - ObjectEmissionDeviceTarget::Vector), - toolchain.bishengPath, stderrPath, diagOS) + return compileDeviceLLVMToObjectWithBridgeLink( + llPath, bridgeBitcodePath, outObjPath, + resolveTargetCPU(module, ObjectEmissionDeviceTarget::Vector), + toolchain, stderrPath, diagOS) ? success() : failure(); } mlir::LogicalResult mlir::pto::emitVPTOCubeDeviceObject( llvm::Module &module, llvm::StringRef llPath, llvm::StringRef outObjPath, - const CANNToolchain &toolchain, llvm::StringRef stderrPath, - llvm::raw_ostream &diagOS) { + const CANNToolchain &toolchain, llvm::StringRef bridgeBitcodePath, + llvm::StringRef stderrPath, llvm::raw_ostream &diagOS) { if (failed(applyVPTOLLVMABINames( module, toolchain.vptoPublicABISuffix(ObjectEmissionDeviceTarget::Cube), @@ -1106,17 +1243,18 @@ mlir::LogicalResult mlir::pto::emitVPTOCubeDeviceObject( if (failed(writeLLVMModule(module, llPath, diagOS))) { return failure(); } - return compileDeviceLLVMToObject(llPath, outObjPath, - resolveTargetCPU(module, - ObjectEmissionDeviceTarget::Cube), - toolchain.bishengPath, stderrPath, diagOS) + return compileDeviceLLVMToObjectWithBridgeLink( + llPath, bridgeBitcodePath, outObjPath, + resolveTargetCPU(module, ObjectEmissionDeviceTarget::Cube), + toolchain, stderrPath, diagOS) ? success() : failure(); } mlir::LogicalResult mlir::pto::emitFatobjLLVM( llvm::Module *cubeModule, llvm::Module *vectorModule, - llvm::StringRef stubSource, llvm::StringRef outputPath, + llvm::StringRef stubSource, llvm::StringRef bridgeWrapperSource, + llvm::StringRef outputPath, llvm::StringRef moduleId, const CANNToolchain &toolchain, TempFileRegistry &tempFiles, VFSIMTSizeFixMode vfsimtSizeFixMode, llvm::raw_ostream &diagOS) { @@ -1132,11 +1270,20 @@ mlir::LogicalResult mlir::pto::emitFatobjLLVM( if (!artifacts.initCommandLogs(diagOS)) { return failure(); } - if (!artifacts.emitCubeObject(cubeModule, toolchain, diagOS)) { + std::string cubeBridgeBitcodePath; + std::string vectorBridgeBitcodePath; + if (!artifacts.compileBridgeWrappers(cubeModule, vectorModule, + bridgeWrapperSource, toolchain, + cubeBridgeBitcodePath, + vectorBridgeBitcodePath, diagOS)) + return failure(); + if (!artifacts.emitCubeObject(cubeModule, toolchain, cubeBridgeBitcodePath, + diagOS)) { return failure(); } if (!artifacts.emitVectorObject(vectorModule, toolchain, - vfsimtSizeFixMode, diagOS)) { + vectorBridgeBitcodePath, vfsimtSizeFixMode, + diagOS)) { return failure(); } if (!artifacts.mergeDeviceObjects(toolchain, diagOS)) { @@ -1185,7 +1332,8 @@ mlir::LogicalResult mlir::pto::linkFatobjs( mlir::LogicalResult mlir::pto::emitFatobjLLVMWithRuntime( llvm::Module *cubeModule, llvm::Module *vectorModule, - llvm::StringRef stubSource, llvm::ToolOutputFile &outputFile, + llvm::StringRef stubSource, llvm::StringRef bridgeWrapperSource, + llvm::ToolOutputFile &outputFile, VFSIMTSizeFixMode vfsimtSizeFixMode, llvm::raw_ostream &diagOS) { if (!cubeModule && !vectorModule) { @@ -1207,11 +1355,20 @@ mlir::LogicalResult mlir::pto::emitFatobjLLVMWithRuntime( return failure(); } - if (!artifacts.emitCubeObject(cubeModule, *toolchain, diagOS)) { + std::string cubeBridgeBitcodePath; + std::string vectorBridgeBitcodePath; + if (!artifacts.compileBridgeWrappers(cubeModule, vectorModule, + bridgeWrapperSource, *toolchain, + cubeBridgeBitcodePath, + vectorBridgeBitcodePath, diagOS)) + return failure(); + if (!artifacts.emitCubeObject(cubeModule, *toolchain, + cubeBridgeBitcodePath, diagOS)) { return failure(); } if (!artifacts.emitVectorObject(vectorModule, *toolchain, - vfsimtSizeFixMode, diagOS)) { + vectorBridgeBitcodePath, vfsimtSizeFixMode, + diagOS)) { return failure(); } diff --git a/tools/ptoas/ObjectEmission.h b/tools/ptoas/ObjectEmission.h index c4e6436524..f44342bcdb 100644 --- a/tools/ptoas/ObjectEmission.h +++ b/tools/ptoas/ObjectEmission.h @@ -119,17 +119,18 @@ LogicalResult emitFatobjCCE(llvm::StringRef cppSource, LogicalResult emitVPTOVectorDeviceObject( llvm::Module &module, llvm::StringRef llPath, llvm::StringRef outObjPath, - const CANNToolchain &toolchain, llvm::StringRef stderrPath, - llvm::raw_ostream &diagOS); + const CANNToolchain &toolchain, llvm::StringRef bridgeBitcodePath, + llvm::StringRef stderrPath, llvm::raw_ostream &diagOS); LogicalResult emitVPTOCubeDeviceObject( llvm::Module &module, llvm::StringRef llPath, llvm::StringRef outObjPath, - const CANNToolchain &toolchain, llvm::StringRef stderrPath, - llvm::raw_ostream &diagOS); + const CANNToolchain &toolchain, llvm::StringRef bridgeBitcodePath, + llvm::StringRef stderrPath, llvm::raw_ostream &diagOS); LogicalResult emitFatobjLLVM( llvm::Module *cubeModule, llvm::Module *vectorModule, - llvm::StringRef stubSource, llvm::StringRef outputPath, + llvm::StringRef stubSource, llvm::StringRef bridgeWrapperSource, + llvm::StringRef outputPath, llvm::StringRef moduleId, const CANNToolchain &toolchain, TempFileRegistry &tempFiles, VFSIMTSizeFixMode vfsimtSizeFixMode, llvm::raw_ostream &diagOS); @@ -155,6 +156,7 @@ LogicalResult linkFatobjs(llvm::ArrayRef fatobjPaths, LogicalResult emitFatobjLLVMWithRuntime(llvm::Module *cubeModule, llvm::Module *vectorModule, llvm::StringRef stubSource, + llvm::StringRef bridgeWrapperSource, llvm::ToolOutputFile &outputFile, VFSIMTSizeFixMode vfsimtSizeFixMode, llvm::raw_ostream &diagOS); diff --git a/tools/ptoas/VPTOFatobjEmission.h b/tools/ptoas/VPTOFatobjEmission.h index ead2cee152..10cc5d00ec 100644 --- a/tools/ptoas/VPTOFatobjEmission.h +++ b/tools/ptoas/VPTOFatobjEmission.h @@ -19,6 +19,7 @@ inline LogicalResult emitVPTOFatobj(llvm::Module *cubeModule, llvm::ToolOutputFile &outputFile, llvm::raw_ostream &diagOS) { return emitFatobjLLVMWithRuntime(cubeModule, vectorModule, stubSource, + /*bridgeWrapperSource=*/"", outputFile, VFSIMTSizeFixMode::Auto, diagOS); } diff --git a/tools/ptoas/driver.cpp b/tools/ptoas/driver.cpp index 84c8ba1200..80bb58d40a 100644 --- a/tools/ptoas/driver.cpp +++ b/tools/ptoas/driver.cpp @@ -1141,6 +1141,7 @@ static LogicalResult emitVPTOLLVMFatobj( if (failed(mlir::pto::emitFatobjLLVM( jobResult.vptoCubeModule.module.get(), jobResult.vptoVectorModule.module.get(), stubSource, + jobResult.vptoBridgeWrapperSource, outputPath, moduleId, *toolchain, context.getTempFiles(), context.getVFSIMTSizeFixMode(), llvm::errs()))) { return failure(); diff --git a/tools/ptoas/ptoas.cpp b/tools/ptoas/ptoas.cpp index 9ca7c61b1c..21e1b8abbe 100644 --- a/tools/ptoas/ptoas.cpp +++ b/tools/ptoas/ptoas.cpp @@ -11,6 +11,7 @@ #include "PTO/IR/VMIUtils.h" #include "PTO/IR/PTOMultiBuffer.h" #include "PTO/Transforms/VPTOLLVMEmitter.h" +#include "PTO/Transforms/VPTOBridgeTokens.h" #include "PTO/Transforms/Passes.h" #include "PTO/Transforms/BufferizableOpInterfaceImpl.h" #include "VPTOHostStubEmission.h" @@ -3215,6 +3216,15 @@ static int emitVPTOBackendResult(ModuleOp module, PTOASCompileResult &result, } } + // Pick up the rendered bridge wrapper source before LLVM translation and + // drop the carrier attribute so it never leaks into the emitted modules. + std::string bridgeWrapperSource; + if (auto wrapperAttr = module->getAttrOfType( + pto::kBridgeWrapperSourceAttrName)) { + bridgeWrapperSource = wrapperAttr.getValue().str(); + module->removeAttr(pto::kBridgeWrapperSourceAttrName); + } + if (failed( pto::lowerVPTOModuleToLLVMModules(module, options, result.vptoCubeModule, @@ -3225,6 +3235,7 @@ static int emitVPTOBackendResult(ModuleOp module, PTOASCompileResult &result, } result.vptoStubSource = std::move(stubSource); + result.vptoBridgeWrapperSource = std::move(bridgeWrapperSource); result.kind = PTOASCompileResultKind::VPTOObject; return 0; } @@ -3236,6 +3247,14 @@ static LogicalResult runVPTOBackendPipeline(OwningOpRef &module, if (!hasTileOpsToExpand) { pm.addNestedPass(pto::createPTOCanonicalizeIRPass()); } + // The declarative bridge pass runs before the pipe family pass: it erases + // the tile handles its bridged ops consume, so the pipe family pass only + // sees the tile handles that remain in the function. + pm.addNestedPass(pto::createPTOLowerDeclarativeBridgeOpsPass()); + pm.addNestedPass(pto::createPTOLowerPipeFamilyOpsPass()); + // Render the bridge wrapper source from the spec the family pass collected + // before the module is split into per-kind kernel modules. + pm.addPass(pto::createVPTOBridgeWrapperGenPass()); pm.addPass(pto::createVPTOSplitCVModulePass()); pm.addPass(pto::createVPTONormalizeContainerPass()); if (hasTileOpsToExpand) { @@ -3357,10 +3376,6 @@ int mlir::pto::compilePTOASModule( llvm::errs() << "Error: --enable-bufid_sync requires --pto-arch=a5.\n"; return 1; } - if (vptoSchedulerMode != VPTOSchedulerCLIMode::Off && arch != "a5") { - llvm::errs() << "Error: --vpto-scheduler requires --pto-arch=a5.\n"; - return 1; - } module->getOperation()->setAttr("pto.target_arch", mlir::StringAttr::get(module->getContext(), arch)); diff --git a/tools/ptoas/ptoas.h b/tools/ptoas/ptoas.h index 41d08dfda7..0de43af0cf 100644 --- a/tools/ptoas/ptoas.h +++ b/tools/ptoas/ptoas.h @@ -117,6 +117,7 @@ struct PTOASCompileResult { void reset() { textOutput.clear(); vptoStubSource.clear(); + vptoBridgeWrapperSource.clear(); vptoCubeModule.reset(); vptoVectorModule.reset(); kind = PTOASCompileResultKind::Text; @@ -125,6 +126,11 @@ struct PTOASCompileResult { PTOASCompileResultKind kind = PTOASCompileResultKind::Text; std::string textOutput; std::string vptoStubSource; + /// Rendered VPTO bridge wrapper C++ source (from the + /// `pto.vpto.bridge.wrapper_source` module attribute); compiled by object + /// emission and linked into the device modules. Empty when the kernel + /// uses no bridge interface. + std::string vptoBridgeWrapperSource; EmittedLLVMModule vptoCubeModule; EmittedLLVMModule vptoVectorModule; };