Skip to content
Open
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
3 changes: 3 additions & 0 deletions include/PTO/Transforms/Passes.h
Original file line number Diff line number Diff line change
Expand Up @@ -107,6 +107,9 @@ LogicalResult validateVPTOEmissionIR(ModuleOp module,
llvm::raw_ostream *diagOS = nullptr);
std::unique_ptr<Pass> createPTOValidateVPTOIRPass();
std::unique_ptr<Pass> createPTOValidateVPTOEmissionIRPass();
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::unique_ptr<Pass> createFoldTileBufIntrinsicsPass();
Expand Down
52 changes: 46 additions & 6 deletions include/PTO/Transforms/Passes.td
Original file line number Diff line number Diff line change
Expand Up @@ -451,13 +451,44 @@ 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 TileLang DSL template functions";
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
TileLang Python DSL 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.).
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).
Expand All @@ -484,10 +515,19 @@ def ExpandTileOp : Pass<"pto-expand-tile-op", "ModuleOp"> {
"PYTHONPATH for tilelang_dsl package (added to env)">,
Option<"pythonExe", "python-exe", "std::string",
/*default=*/"\"python3\"",
"Python executable for tilelang DSL invocation">,
"Python executable for TileLib invocation">,
Option<"daemonSocketPath", "daemon-socket-path", "std::string",
/*default=*/"\"\"",
"Path to Unix domain socket for daemon RPC (if empty, uses subprocess)">
"Path to Unix domain socket for daemon RPC">,
Option<"tileLibBackend", "tile-lib-backend", "std::string",
/*default=*/"\"tilelang\"",
"TileLib backend: tilelang or 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=*/"\"tilelang_dsl.daemon_helper\"",
"Python module used for daemon helper RPC calls">
];
}

Expand Down
1 change: 1 addition & 0 deletions lib/PTO/Transforms/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -42,6 +42,7 @@ add_mlir_dialect_library(PTOTransforms
InsertSync/InsertSyncDebug.cpp
PTOViewToMemref.cpp
PTOValidateIntToPtrUses.cpp
InsertTemplateAttributes.cpp
ExpandTileOp.cpp
FoldTileBufIntrinsics.cpp
PTOLowerToOpLibCalls.cpp
Expand Down
Loading