5 runs of PO2=19 completed in 8.994s, avg=1.799s, 291481.380 cycles/s…

…ec, 1457.162 keccak/sec __device__ static constexpr size_t getFpBufferSize() { return 1554; }
risc0 · Feb 4, 2025 · 2505128 · 2505128
1 parent 9937c00
commit 2505128
Show file tree

Hide file tree

Showing 6 changed files with 17 additions and 4 deletions.
diff --git a/zirgen/Dialect/ByteCode/Transforms/Schedule.cpp b/zirgen/Dialect/ByteCode/Transforms/Schedule.cpp
@@ -327,7 +327,7 @@ const OpInfo& Scheduler::getOpInfo(Operation* topOp) {
       return lhs.second->queuePriority() < rhs.second->queuePriority();
     });
 
-    // Registers that have mee saved so far by operands making use of them.
+    // Registers that have been saved so far by operands making use of them.
     ssize_t adjustedRegs = 0;
 
     size_t maxLive = 0;
@@ -404,3 +404,4 @@ void scheduleBlock(Block* block, ScheduleInterface& scheduleInterface) {
 }
 
 } // namespace zirgen::ByteCode
+
diff --git a/zirgen/Dialect/ByteCode/Transforms/ScheduleZll.cpp b/zirgen/Dialect/ByteCode/Transforms/ScheduleZll.cpp
@@ -484,9 +484,9 @@ struct CloneActiveZllPass : public impl::CloneActiveZllBase<CloneActiveZllPass>
             activeVal.print(llvm::dbgs(), asmState);
             llvm::dbgs() << "\n";
           });
-/*          llvm::errs() << "Spilling ";
+          llvm::errs() << "Spilling ";
           activeVal.print(llvm::errs(), asmState);
-          llvm::errs() << "\n";*/
+          llvm::errs() << "\n";
           active.update(it, ActiveState::SpilledDoNotClone);
           ++opSpillVals;
         } else if (state == ActiveState::Used) {

diff --git a/zirgen/Dialect/Zll/Transforms/BUILD.bazel b/zirgen/Dialect/Zll/Transforms/BUILD.bazel
@@ -42,6 +42,7 @@ cc_library(
         "MakePolynomial.cpp",
         "MultiplyToIf.cpp",
         "PassDetail.h",
+        "ReorderConstraints.cpp",
         "SortForReproducibility.cpp",
         "SplitStage.cpp",
     ],

diff --git a/zirgen/Dialect/Zll/Transforms/Passes.h b/zirgen/Dialect/Zll/Transforms/Passes.h
@@ -34,6 +34,7 @@ std::unique_ptr<mlir::Pass> createBalancedSplitPass(size_t maxOps = 1000);
 std::unique_ptr<mlir::Pass> createSortForReproducibilityPass();
 std::unique_ptr<mlir::OperationPass<mlir::func::FuncOp>> createExtractPolyMixPass();
 std::unique_ptr<mlir::OperationPass<mlir::func::FuncOp>> createAnnotatePolyMixPass();
+std::unique_ptr<mlir::Pass> createReorderConstraintsPass();
 
 // Generate the code for registering passes.
 #define GEN_PASS_REGISTRATION

diff --git a/zirgen/Dialect/Zll/Transforms/Passes.td b/zirgen/Dialect/Zll/Transforms/Passes.td
@@ -106,4 +106,13 @@ def AnnotatePolyMix : Pass<"annotate-poly-mix", "mlir::func::FuncOp"> {
   let constructor = "zirgen::Zll::createAnnotatePolyMixPass()";
 }
 
+def ReorderConstraints : Pass<"reorder-constraints"> {
+  let summary = "Optimize by reordering zll.eqz constraints";
+  let description = [{
+     ReorderConstraints reorders constraints, and puts groups of constraints
+     inside a `if(1)` such that they can be better optimized by CSE.
+  }];
+  let constructor = "zirgen::Zll::createReorderConstraintsPass()";
+}
+
 #endif // ZLL_TRANSFORM_PASSES
diff --git a/zirgen/Main/gen_zirgen.cpp b/zirgen/Main/gen_zirgen.cpp
@@ -1,4 +1,4 @@
-b// Copyright 2025 RISC Zero, Inc.
+// Copyright 2025 RISC Zero, Inc.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -406,6 +406,7 @@ int main(int argc, char* argv[]) {
     checkPasses.addPass(mlir::createCanonicalizerPass());
     checkPasses.addPass(zirgen::dsl::createTopologicalShufflePass());
   }
+  checkPasses.addPass(zirgen::Zll::createReorderConstraintsPass());
 
   if (failed(pm.run(typedModule.value()))) {
     llvm::errs() << "an internal compiler error occurred while lowering this module:\n";