Skip to content

Commit df0b37e

Browse files
committed
[Nvidia][TritonGPU] Add phases to RemoveLayoutConversions
1 parent d6e3640 commit df0b37e

4 files changed

Lines changed: 2797 additions & 84 deletions

File tree

include/triton/Dialect/TritonGPU/Transforms/Passes.td

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -273,6 +273,26 @@ def TritonGPURemoveLayoutConversions : Pass<"tritongpu-remove-layout-conversions
273273
let dependentDialects = ["mlir::triton::gpu::TritonGPUDialect",
274274
"mlir::triton::TritonDialect"];
275275

276+
// These enhancement phases are opt-in (default off) so the pass keeps its
277+
// original behavior for every caller; the NVIDIA backend enables them
278+
// explicitly (see third_party/nvidia/backend/compiler.py).
279+
let options = [
280+
Option<"enableCostBasedResolution", "enable-cost-based-resolution",
281+
"bool", /*default=*/"false",
282+
"Use cost-based conflict resolution instead of heuristic rules">,
283+
Option<"enableBackwardPropagation", "enable-backward-propagation",
284+
"bool", /*default=*/"false",
285+
"Enable backward layout propagation from convert_layout ops">,
286+
Option<"enableSmallComponentSolving", "enable-small-component-solving",
287+
"bool", /*default=*/"false",
288+
"Enable small local layout component solving">,
289+
Option<"enableStoreLayoutRematerialization",
290+
"enable-store-layout-rematerialization",
291+
"bool", /*default=*/"false",
292+
"Drop store/atomic writeback convert_layout ops by rematerializing "
293+
"their address/mask chains in the value's layout">
294+
];
295+
276296
}
277297

278298
def TritonGPUOptimizeThreadLocality : Pass<"tritongpu-optimize-thread-locality", "mlir::ModuleOp"> {

0 commit comments

Comments
 (0)