TTKernel to emitc pipeline

include/ttmlir/Dialect/TTKernel/Pipelines/TTKernelPipelines.h

@@ -0,0 +1,34 @@
+// SPDX-FileCopyrightText: (c) 2024 Tenstorrent AI ULC
+//
+// SPDX-License-Identifier: Apache-2.0
+
+#ifndef TTMLIR_DIALECT_TTKERNEL_PIPELINES_TTKERNELPIPELINES_H
+#define TTMLIR_DIALECT_TTKERNEL_PIPELINES_TTKERNELPIPELINES_H
+
+#include "mlir/Pass/PassOptions.h"
+
+namespace mlir::tt::ttkernel {
+// Options for the TTKernel to EmitC pipeline.
+//
+// struct TTIRToTTMetalBackendPipelineOptions
+//     : public PassPipelineOptions<TTIRToTTMetalBackendPipelineOptions> {
+//   ListOption<int64_t> meshShape{
+//       *this, "mesh-shape", llvm::cl::desc("Set the multi-device mesh
+//       shape.")};
+
+//   // Option to provide a system descriptor flatbuffer file to compile
+//   // against.
+//   //
+//   Option<std::string> systemDescPath{
+//       *this, "system-desc-path",
+//       llvm::cl::desc(
+//           "Pass in a system descriptor flatbuffer to compile against."),
+//       llvm::cl::init("")};
+// };
+
+void createTTKernelToEmitCPipeline(OpPassManager &pm);
+
+void registerTTKernelPipelines();
+} // namespace mlir::tt::ttkernel
+
+#endif

lib/CMakeLists.txt

@@ -27,6 +27,7 @@ MLIRTTMetalDialect
 MLIRTTIRPipelines
 MLIRTTNNPipelines
 MLIRTTMetalPipelines
+MLIRTTKernelPipelines
 )
 
 if (TTMLIR_ENABLE_STABLEHLO)

lib/Dialect/TTKernel/CMakeLists.txt

@@ -1 +1,2 @@
 add_subdirectory(IR)
+add_subdirectory(Pipelines)

lib/Dialect/TTKernel/Pipelines/CMakeLists.txt

@@ -0,0 +1,11 @@
+add_mlir_dialect_library(MLIRTTKernelPipelines
+  TTKernelPipelines.cpp
+
+  ADDITIONAL_HEADER_DIRS
+  ${PROJECT_SOURCE_DIR}/include/ttmlir
+
+  LINK_LIBS PUBLIC
+  MLIRTTIRDialect
+  MLIRTTKernelDialect
+  MLIRTransforms
+)

lib/Dialect/TTKernel/Pipelines/TTKernelPipelines.cpp

@@ -0,0 +1,30 @@
+// SPDX-FileCopyrightText: (c) 2024 Tenstorrent AI ULC
+//
+// SPDX-License-Identifier: Apache-2.0
+
+#include "ttmlir/Dialect/TTKernel/Pipelines/TTKernelPipelines.h"
+
+#include "mlir/Pass/PassManager.h"
+
+#include "ttmlir/Conversion/Passes.h"
+#include "ttmlir/Dialect/TTIR/Transforms/Passes.h"
+
+namespace mlir::tt::ttkernel {
+//===----------------------------------------------------------------------===//
+// Pipeline implementation.
+//===----------------------------------------------------------------------===//
+
+void createTTKernelToEmitCPipeline(OpPassManager &pm) {
+  pm.addPass(createConvertTTKernelToEmitC());
+}
+
+//===----------------------------------------------------------------------===//
+// Pipeline registration.
+//===----------------------------------------------------------------------===//
+
+void registerTTKernelPipelines() {
+  mlir::PassPipelineRegistration<>(
+      "ttkernel-to-emitc-pipeline", "Pipeline lowering ttkernel to emitc.",
+      mlir::tt::ttkernel::createTTKernelToEmitCPipeline);
+}
+} // namespace mlir::tt::ttkernel

lib/RegisterAll.cpp

@@ -14,6 +14,7 @@
 #include "ttmlir/Dialect/TTIR/Pipelines/TTIRPipelines.h"
 #include "ttmlir/Dialect/TTIR/Transforms/Passes.h"
 #include "ttmlir/Dialect/TTKernel/IR/TTKernel.h"
+#include "ttmlir/Dialect/TTKernel/Pipelines/TTKernelPipelines.h"
 #include "ttmlir/Dialect/TTMetal/Pipelines/TTMetalPipelines.h"
 #include "ttmlir/Dialect/TTMetal/Transforms/Passes.h"
 #include "ttmlir/Dialect/TTNN/IR/TTNN.h"
@@ -76,4 +77,5 @@ void mlir::tt::registerAllPasses() {
   mlir::tt::ttir::registerTTIRPipelines();
   mlir::tt::ttnn::registerTTNNPipelines();
   mlir::tt::ttmetal::registerTTMetalPipelines();
+  mlir::tt::ttkernel::registerTTKernelPipelines();
 }

lib/SharedLib/CMakeLists.txt

@@ -24,6 +24,7 @@ set(TTMLIR_LIBS
     MLIRTTNNAnalysis
     MLIRTTNNPipelines
     MLIRTTMetalPipelines
+    MLIRTTKernelPipelines
     TTMLIRTTNNToEmitC
 )