clang format

tenstorrent · Jan 1, 2025 · b608099 · b608099
1 parent acec2a8
commit b608099
Show file tree

Hide file tree

Showing 5 changed files with 52 additions and 63 deletions.
diff --git a/include/ttmlir/Conversion/TTKernelToEmitC/TTKernelToEmitC.h b/include/ttmlir/Conversion/TTKernelToEmitC/TTKernelToEmitC.h
@@ -19,27 +19,24 @@ namespace mlir::tt {
 // Runs a conversion pass to EmitC dialect on a func op containing given
 // region's body. Also, it adds boilerplate code such as includes and namespace
 // declarations.
-LogicalResult convertTTKernelRegionToEmitC(
-    OpBuilder &builder, Region *region,
-    const ttkernel::KernelConfigInterface &kernelConfig);
+LogicalResult
+convertTTKernelRegionToEmitC(OpBuilder &builder, Region *region,
+                             const ttkernel::ThreadType &threadType);
 
 // Converts given region to EmitC dialect and translates it to C++ code.
-LogicalResult
-emitOpRegionAsCpp(Region *region, std::string &regionCpp,
-                          const ttkernel::ThreadType &threadType);
+LogicalResult emitOpRegionAsCpp(Region *region, std::string &regionCpp,
+                                const ttkernel::ThreadType &threadType);
 
-LogicalResult
-emitOpRegionAsCpp(Region *region, llvm::raw_ostream &os,
-                          const ttkernel::ThreadType &threadType);
+LogicalResult emitOpRegionAsCpp(Region *region, llvm::raw_ostream &os,
+                                const ttkernel::ThreadType &threadType);
 
 // Converts dispatch op's regions to C++ code.
 LogicalResult
 emitDispatchOpRegionsAsCpp(ttmetal::DispatchOp dispatchOp,
                            llvm::SmallVector<std::string> &cppStrings);
 
-
-LogicalResult
-emitKernelAsCpp( mlir::ModuleOp op, llvm::raw_ostream &os, const ttkernel::ThreadType &threadType);
+LogicalResult emitKernelAsCpp(mlir::ModuleOp op, llvm::raw_ostream &os,
+                              const ttkernel::ThreadType &threadType);
 
 } // namespace mlir::tt
 

diff --git a/include/ttmlir/Target/TTKernel/TTKernelToCpp.h b/include/ttmlir/Target/TTKernel/TTKernelToCpp.h
@@ -15,9 +15,8 @@ enum class ThreadType : uint32_t;
 // Translates a TTKernel operation to C++ and writes it to the given
 // stream.
 
-LogicalResult translateTTKernelToCpp(
-    Operation *op, llvm::raw_ostream &os, const ThreadType &threadType);
-
+LogicalResult translateTTKernelToCpp(Operation *op, llvm::raw_ostream &os,
+                                     const ThreadType &threadType);
 
 } // namespace mlir::tt::ttkernel
 

diff --git a/lib/Conversion/TTKernelToEmitC/TTKernelToEmitC.cpp b/lib/Conversion/TTKernelToEmitC/TTKernelToEmitC.cpp
@@ -510,7 +510,7 @@ class ThreadConfigHelper {
       builder->create<emitc::IncludeOp>(loc, "compute_kernel_api/reduce.h",
                                         /*isStandard=*/false);
       builder->create<emitc::VerbatimOp>(loc, "namespace NAMESPACE {");
-    }  
+    }
   }
 
   ~ThreadConfigHelper() {
@@ -522,17 +522,15 @@ class ThreadConfigHelper {
   }
 
 private:
-
   OpBuilder *builder;
   Location loc;
   ttkernel::ThreadType threadType;
 };
 
-LogicalResult convertTTKernelRegionToEmitC(
-    OpBuilder &builder, Region *region,
-    const ttkernel::ThreadType &threadType) {
-  ThreadConfigHelper threadConfigHelper(&builder, region->getLoc(),
-                                        threadType);
+LogicalResult
+convertTTKernelRegionToEmitC(OpBuilder &builder, Region *region,
+                             const ttkernel::ThreadType &threadType) {
+  ThreadConfigHelper threadConfigHelper(&builder, region->getLoc(), threadType);
 
   auto funcOp = builder.create<func::FuncOp>(
       region->getLoc(), "kernel_main",
@@ -551,17 +549,15 @@ LogicalResult convertTTKernelRegionToEmitC(
   return success();
 }
 
-LogicalResult
-emitOpRegionAsCpp(Region *region, std::string &regionCpp,
-                          const ttkernel::ThreadType &threadType) {
+LogicalResult emitOpRegionAsCpp(Region *region, std::string &regionCpp,
+                                const ttkernel::ThreadType &threadType) {
 
   llvm::raw_string_ostream os(regionCpp);
   return emitOpRegionAsCpp(region, os, threadType);
 }
 
-LogicalResult
-emitOpRegionAsCpp(Region *region, llvm::raw_ostream &os,
-                          const ttkernel::ThreadType &threadType) {
+LogicalResult emitOpRegionAsCpp(Region *region, llvm::raw_ostream &os,
+                                const ttkernel::ThreadType &threadType) {
 
   // We must load the EmitC dialect before we can emit any EmitC code. This
   // dialect won't be loaded by MLIR until pass manager starts a pass that
@@ -593,13 +589,11 @@ emitDispatchOpRegionsAsCpp(ttmetal::DispatchOp dispatchOp,
   assert(cppStrings.size() == dispatchOp.getNumRegions() &&
          "cppStrings size must match number of regions");
 
-
-
   for (auto &reg : dispatchOp->getRegions()) {
     auto kernelConfig = mlir::cast<ttkernel::KernelConfigInterface>(
         dispatchOp.getKernelConfigs()[reg.getRegionNumber()]);
     if (emitOpRegionAsCpp(&reg, cppStrings[reg.getRegionNumber()],
-                                  kernelConfig.getThreadType())
+                          kernelConfig.getThreadType())
             .failed()) {
       return llvm::failure();
     }
@@ -608,19 +602,14 @@ emitDispatchOpRegionsAsCpp(ttmetal::DispatchOp dispatchOp,
   return success();
 }
 
-LogicalResult
-emitKernelAsCpp(mlir::ModuleOp op, llvm::raw_ostream &os, const ttkernel::ThreadType &threadType )
-{
+LogicalResult emitKernelAsCpp(mlir::ModuleOp op, llvm::raw_ostream &os,
+                              const ttkernel::ThreadType &threadType) {
   std::vector<func::FuncOp> ops;
-  op->walk([&](func::FuncOp entry) {
-    ops.push_back(entry);
-  });
+  op->walk([&](func::FuncOp entry) { ops.push_back(entry); });
 
-  for (const auto &op : ops){
+  for (const auto &op : ops) {
     for (auto &reg : op->getRegions()) {
-      if (emitOpRegionAsCpp(&reg, os,
-                              threadType)
-        .failed()) {
+      if (emitOpRegionAsCpp(&reg, os, threadType).failed()) {
         return llvm::failure();
       }
     }

diff --git a/lib/Target/TTKernel/TTKernelToCpp.cpp b/lib/Target/TTKernel/TTKernelToCpp.cpp
@@ -6,35 +6,35 @@
 #include <cstddef>
 #include <memory>
 
+#include "mlir/Conversion/ArithToEmitC/ArithToEmitCPass.h"
+#include "mlir/Conversion/FuncToEmitC/FuncToEmitCPass.h"
+#include "mlir/Conversion/SCFToEmitC/SCFToEmitC.h"
 #include "mlir/Dialect/EmitC/IR/EmitC.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/Pass/PassManager.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Support/LogicalResult.h"
+#include "mlir/Target/Cpp/CppEmitter.h"
+#include "ttmlir/Conversion/TTKernelToEmitC/TTKernelToEmitC.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/Support/LogicalResult.h"
 #include "llvm/Support/raw_ostream.h"
-#include "mlir/Target/Cpp/CppEmitter.h"
-#include "mlir/Pass/PassManager.h"
-#include "ttmlir/Conversion/TTKernelToEmitC/TTKernelToEmitC.h"
-#include "mlir/Conversion/ArithToEmitC/ArithToEmitCPass.h"
-#include "mlir/Conversion/SCFToEmitC/SCFToEmitC.h"
-#include "mlir/Conversion/FuncToEmitC/FuncToEmitCPass.h"
 #include <llvm/Support/raw_ostream.h>
 
 namespace mlir::tt::ttkernel {
 
-static llvm::LogicalResult translateModuleToCpp(
-    Operation *op, llvm::raw_ostream &os, const ttkernel::ThreadType &threadType) {
-
-    ModuleOp module = dyn_cast<ModuleOp>(op);
-    assert(module && "Expected ModuleOp as top level operation");
-    return mlir::tt::emitKernelAsCpp(module, os, threadType);
+static llvm::LogicalResult
+translateModuleToCpp(Operation *op, llvm::raw_ostream &os,
+                     const ttkernel::ThreadType &threadType) {
 
+  ModuleOp module = dyn_cast<ModuleOp>(op);
+  assert(module && "Expected ModuleOp as top level operation");
+  return mlir::tt::emitKernelAsCpp(module, os, threadType);
 }
 
-LogicalResult translateTTKernelToCpp(
-    Operation *op, llvm::raw_ostream &os, const ttkernel::ThreadType &threadType) {
+LogicalResult translateTTKernelToCpp(Operation *op, llvm::raw_ostream &os,
+                                     const ttkernel::ThreadType &threadType) {
   return translateModuleToCpp(op, os, threadType);
 }
 

diff --git a/lib/Target/TTKernel/TTKernelToCppRegistration.cpp b/lib/Target/TTKernel/TTKernelToCppRegistration.cpp
@@ -6,27 +6,29 @@
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Tools/mlir-translate/Translation.h"
 
+#include "mlir/Dialect/SCF/IR/SCF.h"
 #include "ttmlir/Dialect/TT/IR/TT.h"
 #include "ttmlir/Dialect/TTKernel/IR/TTKernel.h"
 #include "ttmlir/Dialect/TTKernel/IR/TTKernelOpsTypes.h"
 #include "ttmlir/Target/TTKernel/TTKernelToCpp.h"
-#include "mlir/Dialect/SCF/IR/SCF.h"
 #include <mlir/Dialect/MemRef/IR/MemRef.h>
 using namespace mlir;
 
 namespace mlir::tt::ttkernel {
 
-//TODO: Should generalize this to read kernel type from Attribute?
+// TODO: Should generalize this to read kernel type from Attribute?
 void registerTensixKernelToCpp() {
   TranslateFromMLIRRegistration reg(
       "tensixkernel-to-cpp", "translate tensix kernel to C++",
       [](Operation *op, llvm::raw_ostream &os) -> LogicalResult {
         return translateTTKernelToCpp(op, os, tt::ttkernel::ThreadType::Tensix);
       },
       [](DialectRegistry &registry) {
-        registry.insert<mlir::scf::SCFDialect,
-                        mlir::tt::ttkernel::TTKernelDialect, mlir::arith::ArithDialect,
-                        mlir::emitc::EmitCDialect, mlir::func::FuncDialect, mlir::tt::TTDialect, mlir::memref::MemRefDialect>();
+        registry
+            .insert<mlir::scf::SCFDialect, mlir::tt::ttkernel::TTKernelDialect,
+                    mlir::arith::ArithDialect, mlir::emitc::EmitCDialect,
+                    mlir::func::FuncDialect, mlir::tt::TTDialect,
+                    mlir::memref::MemRefDialect>();
       });
 }
 
@@ -37,9 +39,11 @@ void registerNocKernelToCpp() {
         return translateTTKernelToCpp(op, os, tt::ttkernel::ThreadType::Noc);
       },
       [](DialectRegistry &registry) {
-        registry.insert<mlir::scf::SCFDialect,
-                        mlir::tt::ttkernel::TTKernelDialect, mlir::arith::ArithDialect,
-                        mlir::emitc::EmitCDialect, mlir::func::FuncDialect, mlir::tt::TTDialect, mlir::memref::MemRefDialect>();
+        registry
+            .insert<mlir::scf::SCFDialect, mlir::tt::ttkernel::TTKernelDialect,
+                    mlir::arith::ArithDialect, mlir::emitc::EmitCDialect,
+                    mlir::func::FuncDialect, mlir::tt::TTDialect,
+                    mlir::memref::MemRefDialect>();
       });
 }