Skip to content
Merged
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
11 changes: 6 additions & 5 deletions shardy/dialect/sdy/transforms/import/add_data_flow_edges.cc
Original file line number Diff line number Diff line change
Expand Up @@ -38,9 +38,10 @@ using func::FuncOp;

// Adds func input and output data flow edges. Adds func input data flow edges
// only for non-main funcs.
void addFuncDataFlowEdges(ModuleOp moduleOp, const SymbolTable& symbolTable,
IRRewriter& rewriter) {
FuncOp mainFuncOp = getMainFuncOrDie(moduleOp, symbolTable);
void addFuncDataFlowEdgeOps(ModuleOp moduleOp, const SymbolTable& symbolTable,
IRRewriter& rewriter) {
FuncOp mainFuncOp =
getMainFuncOrDie(moduleOp, symbolTable, /*useSingleFunc=*/true);
moduleOp.walk([&](FuncOp funcOp) {
if (funcOp == mainFuncOp) {
return;
Expand Down Expand Up @@ -84,8 +85,8 @@ struct AddDataFlowEdgesPass
addDataFlowEdges(op.getBlockArgumentEdgeOwners(), rewriter);
addDataFlowEdges(op.getOpResultEdgeOwners(), rewriter);
});
if (enableNativeNonFlatSupport) {
addFuncDataFlowEdges(moduleOp, symbolTable, rewriter);
if (addFuncDataFlowEdges) {
addFuncDataFlowEdgeOps(moduleOp, symbolTable, rewriter);
}
}
};
Expand Down
2 changes: 1 addition & 1 deletion shardy/dialect/sdy/transforms/import/import_pipeline.cc
Original file line number Diff line number Diff line change
Expand Up @@ -43,7 +43,7 @@ void addImportPipeline(OpPassManager& pm, int& dumpIndex,
options.dumpDirectory, "before_propagation", dumpIndex++));

pm.addPass(createAddDataFlowEdgesPass(
AddDataFlowEdgesPassOptions{options.enableNativeNonFlatSupport}));
AddDataFlowEdgesPassOptions{/*addFuncDataFlowEdges=*/true}));
pm.addPass(
createApplyShardingConstraintsPass(ApplyShardingConstraintsPassOptions{
options.debugShardingOrigins, options.debugPropagationEdgeSharding}));
Expand Down
6 changes: 2 additions & 4 deletions shardy/dialect/sdy/transforms/import/passes.td
Original file line number Diff line number Diff line change
Expand Up @@ -92,11 +92,9 @@ def AddDataFlowEdgesPass : Pass<"sdy-add-data-flow-edges", "ModuleOp"> {
let dependentDialects = ["mlir::sdy::SdyDialect"];

let options = [
Option<"enableNativeNonFlatSupport", "enable-native-non-flat-support", "bool",
Option<"addFuncDataFlowEdges", "add-func-data-flow-edges", "bool",
/*default=*/"false",
"Whether to propagate shardings directly on a non-flat graph without "
"flattening it. The default is false, meaning it will flatten the "
"graph and then propagate.">
"Whether to add func data flow edges.">
];
}

Expand Down
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: sdy_opt %s -split-input-file -sdy-add-data-flow-edges='enable-native-non-flat-support=true' | FileCheck %s
// RUN: sdy_opt %s -split-input-file -sdy-add-data-flow-edges='add-func-data-flow-edges=true' | FileCheck %s

// CHECK-LABEL: @bar(%arg0: tensor<8xf32>)
func.func @bar(%arg0: tensor<8xf32>) -> tensor<8xf32> {
Expand Down
45 changes: 34 additions & 11 deletions shardy/dialect/sdy/transforms/import/test/import_pipeline.mlir
Original file line number Diff line number Diff line change
Expand Up @@ -32,12 +32,17 @@ func.func @main(%arg0: tensor<8x8xf32>, %arg1: tensor<8x8xf32>) {

sdy.mesh @mesh = <["c"=2, "a"=2, "b"=2]>

// CHECK-LABEL: @add_manual_axes_to_replicated
func.func @add_manual_axes_to_replicated(%arg0: tensor<8xf32>) -> tensor<8xf32> {
// CHECK-NEXT: sdy.manual_computation(%arg0)
// CHECK-LABEL: func @main
func.func @main(%arg0: tensor<8xf32>) -> tensor<8xf32> {
// CHECK-NEXT: %0 = sdy.manual_computation(%arg0)
// CHECK-SAME{LITERAL}: in_shardings=[<@mesh, [{"c", ?}], replicated={"a"}>]
// CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a"}>]
// CHECK-SAME{LITERAL}: manual_axes={"c", "a"} (%arg1: tensor<4xf32>) {
// CHECK-NEXT: %2 = sdy.data_flow_edge %arg1 sharding=<@mesh, [{?}]>
// CHECK-NEXT: sdy.return %2
// CHECK-NEXT: }
// CHECK-NEXT: %1 = sdy.data_flow_edge %0 sharding=<@mesh, [{"c", ?}], replicated={"a"}>
// CHECK-NEXT: return %1
%0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{"c", ?}]>] out_shardings=[<@mesh, [{"c", ?}]>] manual_axes={"c", "a"} (%arg1: tensor<4xf32>) {
sdy.return %arg1 : tensor<4xf32>
} : (tensor<8xf32>) -> tensor<8xf32>
Expand All @@ -50,10 +55,17 @@ sdy.mesh @mesh = <["c"=2, "a"=2, "b"=2]>

// Due to the in_sharding being fully closed, the in_sharding is added to the
// func arg but with the manual axis added as replicated.
// CHECK-LABEL: @add_manual_axes_to_replicated_applied_constraint
// CHECK-LABEL: func @main
// CHECK-SAME %arg0: tensor<16x16xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"c"}, {}], replicated={"a"}>}
// CHECK-SAME -> tensor<16x16xf32> {
func.func @add_manual_axes_to_replicated_applied_constraint(%arg0: tensor<8xf32>) -> tensor<8xf32> {
func.func @main(%arg0: tensor<8xf32>) -> tensor<8xf32> {
// CHECK-NEXT: %0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{"c"}], replicated={"a"}>] out_shardings=[<@mesh, [{"c"}], replicated={"a"}>] manual_axes={"c", "a"} (%arg1: tensor<4xf32>) {
// CHECK-NEXT: %2 = sdy.data_flow_edge %arg1 sharding=<@mesh, [{}]>
// CHECK-NEXT: %3 = stablehlo.add %2, %2
// CHECK-NEXT: sdy.return %3
// CHECK-NEXT: }
// CHECK-NEXT: %1 = sdy.data_flow_edge %0 sharding=<@mesh, [{"c"}], replicated={"a"}>
// CHECK-NEXT: return %1
%0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{"c"}]>] out_shardings=[<@mesh, [{"c"}]>] manual_axes={"c", "a"} (%arg1: tensor<4xf32>) {
%1 = stablehlo.add %arg1, %arg1 : tensor<4xf32>
sdy.return %1 : tensor<4xf32>
Expand All @@ -67,28 +79,39 @@ sdy.mesh @mesh = <["a"=2]>

// This test verifies that the manual axes are cleaned up before adding data
// flow edges.
func.func @manual_axes_cleanup_before_adding_data_flow_edges(%arg0: tensor<8xf32>) -> tensor<8xf32> {
// CHECK-LABEL: func @main
func.func @main(%arg0: tensor<8xf32>) -> tensor<8xf32> {
// CHECK-NEXT: %0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{?}], replicated={"a"}>] out_shardings=[<@mesh, [{?}], replicated={"a"}>] manual_axes={"a"} (%arg1: tensor<8xf32>) {
// CHECK-NEXT: %2 = sdy.data_flow_edge %arg1 sharding=<@mesh, [{?}]>
// CHECK-NEXT: %3 = stablehlo.add %2, %2
// CHECK-NEXT: sdy.return %3
// CHECK-NEXT: }
// CHECK-NEXT: %1 = sdy.data_flow_edge %0 sharding=<@mesh, [{?}], replicated={"a"}>
// CHECK-NEXT: return %1
%0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{?}]>] out_shardings=[<@mesh, [{?}]>] manual_axes={"a"} (%arg1: tensor<8xf32>) {
%1 = stablehlo.add %arg1, %arg1 : tensor<8xf32>
sdy.return %1 : tensor<8xf32>
} : (tensor<8xf32>) -> tensor<8xf32>
// CHECK: sdy.data_flow_edge %0 sharding=<@mesh, [{?}], replicated={"a"}> : tensor<8xf32>
return %0 : tensor<8xf32>
}

// -----

sdy.mesh @mesh = <["a"=2]>

// CHECK-LABEL: func @single_call
func.func @single_call(%arg0: tensor<8xf32>) -> tensor<8xf32> {
// CHECK-NEXT: %0 = call @foo(%arg0) : (tensor<8xf32>) -> tensor<8xf32>
// CHECK-NEXT: return %0 : tensor<8xf32>
// test: single_call
// CHECK-LABEL: func @main
func.func @main(%arg0: tensor<8xf32>) -> tensor<8xf32> {
// CHECK-NEXT: %[[CALL:.*]] = call @foo(%arg0) : (tensor<8xf32>) -> tensor<8xf32>
// CHECK-NEXT: %[[EDGE:.*]] = sdy.func_data_flow_edge %[[CALL]] : tensor<8xf32>
// CHECK-NEXT: return %[[EDGE]] : tensor<8xf32>
%0 = call @foo(%arg0) : (tensor<8xf32>) -> tensor<8xf32>
return %0 : tensor<8xf32>
}

// CHECK-LABEL: func private @foo
func.func private @foo(%arg0: tensor<8xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"a"}]>}) -> tensor<8xf32> {
// CHECK-NEXT: %0 = sdy.func_data_flow_edge %arg0 {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"a"}]>]>}
// CHECK-NEXT: return %0
return %arg0 : tensor<8xf32>
}
1 change: 1 addition & 0 deletions shardy/dialect/sdy/transforms/propagation/BUILD
Original file line number Diff line number Diff line change
Expand Up @@ -57,6 +57,7 @@ cc_library(
":sharding_projection",
":utils",
"//shardy/common:file_utils",
"//shardy/common:logging",
"//shardy/dialect/sdy/ir:dialect",
"//shardy/dialect/sdy/transforms/common:op_properties",
"//shardy/dialect/sdy/transforms/common:propagation_options",
Expand Down
Loading
Loading