Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 0 additions & 9 deletions docker/test_wheel_imports.sh
Original file line number Diff line number Diff line change
Expand Up @@ -219,13 +219,4 @@ grep -q "candidates = " "${CLEAN_ENV_PTO_IR}" || {
echo "Error: clean-environment ptoas smoke output is missing TileLib candidate metadata" >&2
exit 1
}
if ! grep -q "TileLib daemon started successfully" "${CLEAN_ENV_LOG}"; then
echo "Error: TileLib daemon did not report a successful start" >&2
exit 1
fi
if ! grep -q "TileLib daemon stopped" "${CLEAN_ENV_LOG}"; then
echo "Error: TileLib daemon did not report a clean stop" >&2
exit 1
fi

echo "All wheel import tests passed!"
3 changes: 1 addition & 2 deletions docs/designs/ptoas-compiler-dso-wheel-linking.md
Original file line number Diff line number Diff line change
Expand Up @@ -109,8 +109,7 @@ _core
│ ├── ptoas.cpp
│ ├── driver.cpp
│ ├── VPTOHostStubEmission.cpp
│ ├── ObjectEmission.cpp
│ └── TilelangDaemon.cpp
│ └── ObjectEmission.cpp
├── PTOCAPI
└── PTOASPythonCAPI
```
Expand Down
14 changes: 8 additions & 6 deletions docs/designs/ptoas-python-launcher-layout.md
Original file line number Diff line number Diff line change
Expand Up @@ -139,7 +139,7 @@ assemble Python packages from unrelated build directories.
The current archive is built against CPython 3.11 and requires a CPython 3.11
interpreter. `bin/ptoas` adds the archive root to `sys.path`, then uses the same
`ptoas._cli -> ptoas._core` path as the install tree. The packaged `ptodsl/`
tree supports the compiler's default PTODSL TileLib backend; it does not turn
tree supports the compiler's PTODSL TileLib implementation; it does not turn
the archive into a normal pip-installable PTODSL distribution.

Linux archives use package-relative and archive-relative `$ORIGIN` RPATHs;
Expand All @@ -155,8 +155,10 @@ declared installation layout. CTest and direct developer-tree runs must set an
explicit matching `PYTHONPATH`; PTODSL does not guess repository, LLVM build,
or PTOAS install paths at import time.

TileOp expansion remains a lazy, separate daemon process. The PTOAS CLI passes
the packaged PTODSL root and the active Python executable to the native driver,
which starts the daemon only when expansion is required. Keeping the daemon out
of the compiler process also prevents independently packaged MLIR/LLVM Python
bindings from registering runtime state in the native PTOAS process.
TileOp expansion runs in the CLI's existing Python process. `_core.main` creates
one Python-owned MLIR context for the compilation session, and the native driver
borrows that exact context. The in-process TileLib service materializes a source
module in the shared context and clones it into native ownership before the
Python module owner can be released. The packaged MLIR bindings and PTOAS must
therefore remain one ABI-matched runtime rather than independently replaceable
components.
30 changes: 16 additions & 14 deletions docs/designs/ptodsl-tilelib-template-selection-design.md
Original file line number Diff line number Diff line change
Expand Up @@ -2,10 +2,7 @@

## Background

PTOAS currently supports two TileLib backends for VPTO tile-op expansion:

- `tilelang`, the legacy TileLangDSL template implementation.
- `ptodsl`, the PTODSL-native template implementation.
PTOAS uses the PTODSL-native TileLib implementation for VPTO tile-op expansion.

A tile op may have several legal implementations for the same op
name. Those implementations can differ by dtype, layout, memory space,
Expand Down Expand Up @@ -42,24 +39,24 @@ in the ISA and user guide documents, not here.

## Pipeline

The PTODSL TileLib path has two compiler interactions with the Python daemon.
The PTODSL TileLib path has two interactions with the in-process Python service.

```text
TileOp in MLIR
|
| InsertTemplateAttributes
| - reconstruct operand specs from MLIR
| - collect context attributes
| - ask the PTODSL daemon for legal candidates
| - ask the PTODSL service for legal candidates
| - store compact candidate metadata on the TileOp
v
TileOp with candidates attr
|
| ExpandTileOp
| - build a specialization key from current MLIR operands and attrs
| - choose candidate 0 from the compact candidates attr
| - ask the daemon to render that candidate
| - clone the generated helper and replace the TileOp with func.call
| - ask the service to materialize that candidate in the shared context
| - import the generated entry/helpers and replace the TileOp with func.call
v
VPTO-facing IR
```
Expand All @@ -69,6 +66,10 @@ before later passes can make candidate information harder to reconstruct.
`ExpandTileOp` still renders from the current MLIR operands so the helper body
matches the actual operand types and view metadata that survived to expansion.

Both stages are compiler-session passes. They require the `TileLibService`
owned by `PTOASContext`, are constructed explicitly by the PTOAS pipeline, and
are not registered as standalone textual passes.

## Template Metadata

PTODSL template authors register versions through `tilelib.tile_template`.
Expand Down Expand Up @@ -102,7 +103,7 @@ remain in Python metadata for selection, diagnostics, and future tooling.
## Operand Specs

Both `InsertTemplateAttributes` and `ExpandTileOp` reconstruct operand specs
from MLIR. The JSON shape sent to the daemon is deliberately close to
from MLIR. The JSON shape sent to the Python service is deliberately close to
`TileSpec`, `ViewSpec`, `ScalarSpec`, and `VectorSpec`.

| Operand kind | Required metadata |
Expand Down Expand Up @@ -137,7 +138,7 @@ template is considered ported.

## Candidate Legality And Ranking

The daemon loads only the template module for the requested op and target. It
The service loads only the template module for the requested op and target. It
then evaluates each registered candidate:

1. Bind positional MLIR operands to the template parameter names.
Expand All @@ -149,7 +150,7 @@ then evaluates each registered candidate:
7. Run custom constraint predicates.
8. Sort legal candidates by descending priority.

If no candidate is legal, the daemon reports a `NoMatchingTemplate` error with
If no candidate is legal, the service reports a `NoMatchingTemplate` error with
per-candidate reasons. If multiple candidates tie for the highest priority and
no explicit candidate is requested, the registry reports ambiguity rather than
silently picking one.
Expand All @@ -170,7 +171,7 @@ TileOp. Each entry contains:
- `tail`

This attribute is intentionally not a copy of the full Python metadata object.
Legality has already happened in the daemon. The IR only needs a stable list of
Legality has already happened in the service. The IR only needs a stable list of
legal render targets and the small amount of metadata consumed by downstream
passes.

Expand All @@ -180,8 +181,9 @@ Python metadata. Add a field only when a C++ pass or IR-level test consumes it.
## Expansion And Specialization

`ExpandTileOp` uses the first candidate in the compact candidate list. For
PTODSL, it passes the selected candidate name back to the daemon so rendering
cannot accidentally choose a different legal template after the metadata pass.
PTODSL, it passes the selected candidate name back to the service so
materialization cannot accidentally choose a different legal template after the
metadata pass.

The specialization key deduplicates generated helpers inside one module. It
must include every input that can change the rendered helper body:
Expand Down
3 changes: 3 additions & 0 deletions include/PTO/Transforms/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -14,5 +14,8 @@
set(LLVM_TARGET_DEFINITIONS Passes.td)
mlir_tablegen(Passes.h.inc -gen-pass-decls -name PTO)

set(LLVM_TARGET_DEFINITIONS TileLibPasses.td)
mlir_tablegen(TileLibPasses.h.inc -gen-pass-decls -name PTOTileLib)

# [关键] 重命名为 PTOPassesIncGen (去掉 MLIR 前缀)
add_public_tablegen_target(PTOPassesIncGen)
8 changes: 4 additions & 4 deletions include/PTO/Transforms/Passes.h
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,7 @@
#include "llvm/ADT/StringRef.h"
#include "mlir/Pass/Pass.h"
#include "PTO/IR/PTODialect.h"
#include "PTO/Transforms/TileLibService.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
Expand Down Expand Up @@ -127,11 +128,10 @@ std::unique_ptr<Pass> createVMILegalizeArithSelectPass();
std::unique_ptr<Pass> createVMILowerUnifiedToLegacyPass();
std::unique_ptr<Pass> createVMINormalizeSignlessIntToUnsignedPass();
std::unique_ptr<Pass> createVMIToVPTOPass();
std::unique_ptr<Pass> createInsertTemplateAttributesPass();
std::unique_ptr<Pass> createInsertTemplateAttributesPass(
const InsertTemplateAttributesOptions &options);
std::unique_ptr<Pass> createExpandTileOpPass();
std::unique_ptr<Pass> createExpandTileOpPass(const ExpandTileOpOptions &options);
std::shared_ptr<TileLibService> tileLibService);
std::unique_ptr<Pass>
createExpandTileOpPass(std::shared_ptr<TileLibService> tileLibService);
std::unique_ptr<Pass> createFoldTileBufIntrinsicsPass();
std::unique_ptr<Pass> createFoldTileBufIntrinsicsPass(llvm::StringRef foldMode);
std::unique_ptr<Pass> createPTOCanonicalizeIRPass();
Expand Down
76 changes: 1 addition & 75 deletions include/PTO/Transforms/Passes.td
Original file line number Diff line number Diff line change
Expand Up @@ -507,84 +507,10 @@ def PTOResolveReservedBuffers : Pass<"pto-resolve-reserved-buffers", "ModuleOp">
];
}

def InsertTemplateAttributes
: Pass<"pto-insert-template-attributes", "ModuleOp"> {
let summary = "Attach legal PTODSL template candidates to tile operations";
let description = [{
Queries the PTODSL TileLib daemon for legal template candidates and stores
the compact candidate list on each tile operation as the `candidates`
attribute. Each candidate contains only id, name, loop_depth, postupdate,
and tail metadata.
}];
let constructor = "mlir::pto::createInsertTemplateAttributesPass()";
let dependentDialects = [
"mlir::pto::PTODialect",
"mlir::func::FuncDialect"
];
let options = [
Option<"pythonExe", "python-exe", "std::string",
/*default=*/"\"python3\"",
"Python executable for TileLib metadata invocation">,
Option<"daemonSocketPath", "daemon-socket-path", "std::string",
/*default=*/"\"\"",
"Path to the PTODSL TileLib daemon Unix socket">,
Option<"tileLibPkgPath", "tile-lib-pkg-path", "std::string",
/*default=*/"\"\"",
"PYTHONPATH root for PTODSL">,
Option<"daemonHelperModule", "daemon-helper-module", "std::string",
/*default=*/"\"ptodsl.tilelib.serving.helper\"",
"Python module used for daemon metadata RPC calls">
];
}

def ExpandTileOp : Pass<"pto-expand-tile-op", "ModuleOp"> {
let summary = "Expand tile ops into calls to TileLib template functions";
let description = [{
Expands tile-level operations (pto.tadd, pto.tsub, etc.) by invoking the
selected Python TileLib backend to instantiate template libraries. The
generated template functions use tile_buf parameters and contain
vector-level implementations (pto.vecscope, pto.vlds, pto.vadd,
pto.vsts, etc.).

Each tile op is replaced by a func.call to the generated template function,
with tile_buf operands passed directly (no type bridging).

After this pass, the Inline pass inlines template bodies, and
FoldTileBufIntrinsics resolves tile_buf_addr / tile_valid_rows /
tile_valid_cols.
}];
let constructor = "mlir::pto::createExpandTileOpPass()";
let dependentDialects = [
"mlir::pto::PTODialect",
"mlir::memref::MemRefDialect",
"mlir::arith::ArithDialect",
"mlir::func::FuncDialect",
"mlir::scf::SCFDialect",
"mlir::vector::VectorDialect"
];
let options = [
Option<"pythonExe", "python-exe", "std::string",
/*default=*/"\"python3\"",
"Python executable for TileLib invocation">,
Option<"daemonSocketPath", "daemon-socket-path", "std::string",
/*default=*/"\"\"",
"Path to Unix domain socket for daemon RPC">,
Option<"tileLibBackend", "tile-lib-backend", "std::string",
/*default=*/"\"ptodsl\"",
"TileLib backend: ptodsl">,
Option<"tileLibPkgPath", "tile-lib-pkg-path", "std::string",
/*default=*/"\"\"",
"PYTHONPATH root for the selected TileLib backend">,
Option<"daemonHelperModule", "daemon-helper-module", "std::string",
/*default=*/"\"ptodsl.tilelib.serving.helper\"",
"Python module used for daemon helper RPC calls">
];
}

def FoldTileBufIntrinsics : Pass<"pto-fold-tile-buf-intrinsics", "mlir::func::FuncOp"> {
let summary = "Fold structured-view intrinsics after template inlining";
let description = [{
After TileLang DSL template functions are inlined, the IR contains
After PTODSL template functions are inlined, the IR contains
structured-view intrinsics whose operands are now bound to concrete values.

This pass resolves them:
Expand Down
62 changes: 62 additions & 0 deletions include/PTO/Transforms/TileLibPasses.td
Original file line number Diff line number Diff line change
@@ -0,0 +1,62 @@
// 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.

// TileLib passes require a live compiler-session service and are intentionally
// excluded from the global textual pass registry. Their constructor strings
// suppress generated default factories; PTOAS creates them through the
// service-injecting factories declared in Passes.h.

#ifndef MLIR_DIALECT_PTO_TILELIB_PASSES
#define MLIR_DIALECT_PTO_TILELIB_PASSES

include "mlir/Pass/PassBase.td"

def InsertTemplateAttributes
: Pass<"pto-insert-template-attributes", "ModuleOp"> {
let summary = "Attach legal PTODSL template candidates to tile operations";
let description = [{
Queries the compiler's in-process PTODSL TileLib service for legal template
candidates and stores the compact candidate list on each tile operation as
the `candidates` attribute. Each candidate contains only id, name,
loop_depth, postupdate, and tail metadata.
}];
let constructor = "mlir::pto::createInsertTemplateAttributesPass()";
let dependentDialects = [
"mlir::pto::PTODialect",
"mlir::func::FuncDialect"
];
}

def ExpandTileOp : Pass<"pto-expand-tile-op", "ModuleOp"> {
let summary = "Expand tile ops into calls to TileLib template functions";
let description = [{
Expands tile-level operations (pto.tadd, pto.tsub, etc.) by asking the
compiler's in-process PTODSL TileLib service to instantiate template
libraries in the current MLIRContext. The generated template functions use
tile_buf parameters and contain vector-level implementations (pto.vecscope,
pto.vlds, pto.vadd, pto.vsts, etc.).

Each tile op is replaced by a func.call to the generated template function,
with tile_buf operands passed directly (no type bridging).

After this pass, the Inline pass inlines template bodies, and
FoldTileBufIntrinsics resolves tile_buf_addr / tile_valid_rows /
tile_valid_cols.
}];
let constructor = "mlir::pto::createExpandTileOpPass()";
let dependentDialects = [
"mlir::pto::PTODialect",
"mlir::memref::MemRefDialect",
"mlir::arith::ArithDialect",
"mlir::func::FuncDialect",
"mlir::scf::SCFDialect",
"mlir::vector::VectorDialect"
];
}

#endif // MLIR_DIALECT_PTO_TILELIB_PASSES
56 changes: 56 additions & 0 deletions include/PTO/Transforms/TileLibService.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,56 @@
// 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 MLIR_DIALECT_PTO_TRANSFORMS_TILELIBSERVICE_H
#define MLIR_DIALECT_PTO_TRANSFORMS_TILELIBSERVICE_H

#include "mlir/IR/BuiltinOps.h"
#include "mlir/Support/LogicalResult.h"

#include "llvm/ADT/STLFunctionalExtras.h"
#include "llvm/ADT/StringRef.h"

#include <string>

namespace mlir::pto {

/// Pure-data request used by the in-process TileLib materializer. The JSON
/// fields are request data only; generated MLIR never crosses this interface as
/// text. Keeping this interface independent of pybind11 allows transform passes
/// to remain usable from native tests and non-Python hosts.
struct TileLibMaterializationRequest {
std::string target;
std::string op;
std::string operandSpecsJson;
std::string contextAttrsJson;
std::string candidateId;
};

using TileLibMaterializationCallback =
llvm::function_ref<LogicalResult(ModuleOp source, StringRef entrySymbol)>;

/// Synchronous handoff for a materialized TileLib implementation. The source
/// module is borrowed and remains owned by the service for the duration of the
/// callback. Consumers must clone/import any operations they need before the
/// callback returns.
class TileLibService {
public:
virtual ~TileLibService() = default;

virtual FailureOr<std::string>
getMetadata(const TileLibMaterializationRequest &request) = 0;

virtual LogicalResult
materialize(const TileLibMaterializationRequest &request,
MLIRContext &context,
TileLibMaterializationCallback callback) = 0;
};

} // namespace mlir::pto

#endif // MLIR_DIALECT_PTO_TRANSFORMS_TILELIBSERVICE_H
Loading