Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Enzyme ops removal for stablehlo.while #262

Open
wants to merge 14 commits into
base: main
Choose a base branch
from
Open

Enzyme ops removal for stablehlo.while #262

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
339eead
ops removal for while op
Pangoraw Jan 18, 2025
a551585
more tests
Pangoraw Jan 22, 2025
3371475
rewriter
Pangoraw Jan 22, 2025
77ba894
fix formatting
Pangoraw Jan 22, 2025
d111e1f
tmp enzyme commit
Pangoraw Jan 22, 2025
f15da04
update
Pangoraw Jan 22, 2025
9b5498d
use new api
Pangoraw Jan 22, 2025
d063831
adapt to removeEnzymeOpsInBlock change
Pangoraw Jan 22, 2025
dabe2df
update upstream enzyme
Pangoraw Jan 28, 2025
8fe3e6b
comments
Pangoraw Jan 28, 2025
2d518c0
Update workspace.bzl
Pangoraw Jan 28, 2025
43df343
Merge branch 'main' into remove-ops-while
Pangoraw Jan 28, 2025
a8c3b0c
adapt to while dead branch removal
Pangoraw Jan 28, 2025
ce4bb7e
remove log
Pangoraw Feb 4, 2025
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
440 changes: 400 additions & 40 deletions src/enzyme_ad/jax/Implementations/StableHLOAutoDiffOpInterfaceImpl.cpp
View file
Open in desktop

Large diffs are not rendered by default.

5 changes: 5 additions & 0 deletions src/enzyme_ad/jax/Implementations/WhileLoopInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,6 +72,11 @@ std::optional<int64_t> WhileLoopInfo::getConstantLimit() {
return (*limitAttr.begin()).getSExtValue();
}

int64_t WhileLoopInfo::getConstantNumIters() {
return (getConstantLimit().value() - getConstantStart().value()) /
getConstantStep().value();
}

Value WhileLoopInfo::getNumIters(mlir::OpBuilder &builder) {
auto opReg = op->getParentRegion();
if (!opReg->isAncestor(limit.getParentRegion()) ||
Expand Down
1 change: 1 addition & 0 deletions src/enzyme_ad/jax/Implementations/WhileLoopInfo.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,6 +37,7 @@ struct WhileLoopInfo {
std::optional<int64_t> getConstantStart();
std::optional<int64_t> getConstantLimit();

int64_t getConstantNumIters();
Value getNumIters(OpBuilder &builder);
};

Expand Down
11 changes: 5 additions & 6 deletions src/enzyme_ad/jax/Passes/EnzymeHLOUnroll.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,13 +1,14 @@
//===- EnzymeWrapPass.cpp - Replace calls with their derivatives ------------ //
//===- EnzymeHLOUnroll.cpp - Unroll stablehlo.while loops -----------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
// This file implements a pass to create wrapper functions which differentiate
// ops.
// This file implements a pass to unroll stablehlo.while ops with known number
// of iterations.
//
//===----------------------------------------------------------------------===//

#include "mlir/IR/Builders.h"
Expand Down Expand Up @@ -55,9 +56,7 @@ struct WhileUnroll : public OpRewritePattern<mlir::stablehlo::WhileOp> {
auto bodyTerm = cast<stablehlo::ReturnOp>(&op.getBody().front().back());
auto loopBodyBlock = &op.getBody().front();

auto iters =
(info.getConstantLimit().value() - info.getConstantStart().value()) /
info.getConstantStep().value();
auto iters = info.getConstantNumIters();

SmallVector<Value> results(op.getOperands().begin(),
op.getOperands().end());
Expand Down
61 changes: 31 additions & 30 deletions test/lit_tests/diffrules/stablehlo/while3.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,40 +57,41 @@ module {
// FORWARD-NEXT: }

// REVERSE: func.func private @diffef(%arg0: tensor<f64>, %arg1: tensor<f64>) -> (tensor<f64>, tensor<f64>) {
// REVERSE-NEXT: %c = stablehlo.constant dense<3> : tensor<i64>
// REVERSE-NEXT: %c_0 = stablehlo.constant dense<1> : tensor<i64>
// REVERSE-NEXT: %c_1 = stablehlo.constant dense<0> : tensor<i64>
// REVERSE-NEXT: %cst = arith.constant dense<0.000000e+00> : tensor<f64>
// REVERSE-NEXT: %0 = "enzyme.init"() : () -> !enzyme.Cache<tensor<f64>>
// REVERSE-NEXT: %1 = "enzyme.init"() : () -> !enzyme.Cache<tensor<f64>>
// REVERSE-NEXT: %2 = stablehlo.subtract %c, %c_1 : tensor<i64>
// REVERSE-NEXT: %3 = stablehlo.divide %2, %c_0 : tensor<i64>
// REVERSE-NEXT: %4:2 = stablehlo.while(%iterArg = %c_1, %iterArg_2 = %arg0) : tensor<i64>, tensor<f64>
// REVERSE-NEXT: %c = stablehlo.constant dense<2> : tensor<i64>
// REVERSE-NEXT: %cst = arith.constant dense<0.000000e+00> : tensor<3xf64>
// REVERSE-NEXT: %c_0 = stablehlo.constant dense<3> : tensor<i64>
// REVERSE-NEXT: %c_1 = stablehlo.constant dense<1> : tensor<i64>
// REVERSE-NEXT: %c_2 = stablehlo.constant dense<0> : tensor<i64>
// REVERSE-NEXT: %cst_3 = arith.constant dense<0.000000e+00> : tensor<f64>
// REVERSE-NEXT: %0 = stablehlo.subtract %c_0, %c_2 : tensor<i64>
// REVERSE-NEXT: %1 = stablehlo.divide %0, %c_1 : tensor<i64>
// REVERSE-NEXT: %2:3 = stablehlo.while(%iterArg = %c_2, %iterArg_4 = %arg0, %iterArg_5 = %cst) : tensor<i64>, tensor<f64>, tensor<3xf64>
// REVERSE-NEXT: cond {
// REVERSE-NEXT: %8 = stablehlo.compare LT, %iterArg, %c : (tensor<i64>, tensor<i64>) -> tensor<i1>
// REVERSE-NEXT: stablehlo.return %8 : tensor<i1>
// REVERSE-NEXT: %6 = stablehlo.compare LT, %iterArg, %c_0 : (tensor<i64>, tensor<i64>) -> tensor<i1>
// REVERSE-NEXT: stablehlo.return %6 : tensor<i1>
// REVERSE-NEXT: } do {
// REVERSE-NEXT: "enzyme.push"(%1, %iterArg_2) : (!enzyme.Cache<tensor<f64>>, tensor<f64>) -> ()
// REVERSE-NEXT: "enzyme.push"(%0, %iterArg_2) : (!enzyme.Cache<tensor<f64>>, tensor<f64>) -> ()
// REVERSE-NEXT: %8 = stablehlo.multiply %iterArg_2, %iterArg_2 : tensor<f64>
// REVERSE-NEXT: %9 = stablehlo.add %iterArg, %c_0 : tensor<i64>
// REVERSE-NEXT: stablehlo.return %9, %8 : tensor<i64>, tensor<f64>
// REVERSE-NEXT: %6 = stablehlo.reshape %iterArg_4 : (tensor<f64>) -> tensor<1xf64>
// REVERSE-NEXT: %7 = stablehlo.dynamic_update_slice %iterArg_5, %6, %iterArg : (tensor<3xf64>, tensor<1xf64>, tensor<i64>) -> tensor<3xf64>
// REVERSE-NEXT: %8 = stablehlo.multiply %iterArg_4, %iterArg_4 : tensor<f64>
// REVERSE-NEXT: %9 = stablehlo.add %iterArg, %c_1 : tensor<i64>
// REVERSE-NEXT: stablehlo.return %9, %8, %7 : tensor<i64>, tensor<f64>, tensor<3xf64>
// REVERSE-NEXT: }
// REVERSE-NEXT: %5 = arith.addf %arg1, %cst : tensor<f64>
// REVERSE-NEXT: %6:2 = stablehlo.while(%iterArg = %c_1, %iterArg_2 = %5) : tensor<i64>, tensor<f64>
// REVERSE-NEXT: %3 = arith.addf %arg1, %cst_3 : tensor<f64>
// REVERSE-NEXT: %4:3 = stablehlo.while(%iterArg = %c_2, %iterArg_4 = %3, %iterArg_5 = %c) : tensor<i64>, tensor<f64>, tensor<i64>
// REVERSE-NEXT: cond {
// REVERSE-NEXT: %8 = stablehlo.compare LT, %iterArg, %3 : (tensor<i64>, tensor<i64>) -> tensor<i1>
// REVERSE-NEXT: stablehlo.return %8 : tensor<i1>
// REVERSE-NEXT: %6 = stablehlo.compare LT, %iterArg, %1 : (tensor<i64>, tensor<i64>) -> tensor<i1>
// REVERSE-NEXT: stablehlo.return %6 : tensor<i1>
// REVERSE-NEXT: } do {
// REVERSE-NEXT: %8 = stablehlo.add %iterArg, %c_0 : tensor<i64>
// REVERSE-NEXT: %9 = "enzyme.pop"(%1) : (!enzyme.Cache<tensor<f64>>) -> tensor<f64>
// REVERSE-NEXT: %10 = "enzyme.pop"(%0) : (!enzyme.Cache<tensor<f64>>) -> tensor<f64>
// REVERSE-NEXT: %11 = stablehlo.multiply %iterArg_2, %10 : tensor<f64>
// REVERSE-NEXT: %12 = arith.addf %11, %cst : tensor<f64>
// REVERSE-NEXT: %13 = stablehlo.multiply %iterArg_2, %9 : tensor<f64>
// REVERSE-NEXT: %14 = arith.addf %12, %13 : tensor<f64>
// REVERSE-NEXT: stablehlo.return %8, %14 : tensor<i64>, tensor<f64>
// REVERSE-NEXT: %6 = stablehlo.add %iterArg, %c_1 : tensor<i64>
// REVERSE-NEXT: %7 = stablehlo.dynamic_slice %2#2, %iterArg_5, sizes = [1] : (tensor<3xf64>, tensor<i64>) -> tensor<1xf64>
// REVERSE-NEXT: %8 = stablehlo.reshape %7 : (tensor<1xf64>) -> tensor<f64>
// REVERSE-NEXT: %9 = stablehlo.multiply %iterArg_4, %8 : tensor<f64>
// REVERSE-NEXT: %10 = arith.addf %9, %cst_3 : tensor<f64>
// REVERSE-NEXT: %11 = stablehlo.multiply %iterArg_4, %8 : tensor<f64>
// REVERSE-NEXT: %12 = arith.addf %10, %11 : tensor<f64>
// REVERSE-NEXT: %13 = stablehlo.subtract %iterArg_5, %c_1 : tensor<i64>
// REVERSE-NEXT: stablehlo.return %6, %12, %13 : tensor<i64>, tensor<f64>, tensor<i64>
// REVERSE-NEXT: }
// REVERSE-NEXT: %7 = arith.addf %6#1, %cst : tensor<f64>
// REVERSE-NEXT: return %4#1, %7 : tensor<f64>, tensor<f64>
// REVERSE-NEXT: %5 = arith.addf %4#1, %cst_3 : tensor<f64>
// REVERSE-NEXT: return %2#1, %5 : tensor<f64>, tensor<f64>
// REVERSE-NEXT: }
Loading
Loading