Convert stablehlo.composite ops to func.call ops before converting to…

… TTIR so that the TTIR inliner can inline them

lib/Conversion/StableHLOToTTIR/CMakeLists.txt

@@ -17,4 +17,5 @@ add_mlir_library(TTMLIRStableHLOToTTIR
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRPass
+  StablehloPasses
 )

lib/Dialect/TTIR/Pipelines/TTIRPipelines.cpp

@@ -7,6 +7,7 @@
 #include "mlir/Pass/PassManager.h"
 #include "mlir/Transforms/Passes.h"
 
+#include "stablehlo/transforms/Passes.h"
 #include "ttmlir/Conversion/Passes.h"
 
 namespace mlir::tt::ttir {
@@ -20,6 +21,7 @@ void createStableHLOToTTIRPipeline(
   if (options.arithDialectConversionsEnabled) {
     pm.addPass(createConvertArithToStableHLOPass());
   }
+  pm.addPass(stablehlo::createStablehloLegalizeCompositeToCallPass());
   pm.addPass(createConvertStableHLOToTTIRPass());
   if (options.removeDeadValuesEnabled) {
     pm.addPass(mlir::createRemoveDeadValuesPass());

test/ttmlir/Conversion/StableHLOToTTIR/composite_op.mlir

@@ -0,0 +1,16 @@
+// REQUIRES: stablehlo
+// RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | FileCheck %s
+module @jit_eltwise_add attributes {} {
+  func.func private @add_impl(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
+    %0 = stablehlo.add %arg0, %arg1 : tensor<13x21x3xf32>
+    return %0 : tensor<13x21x3xf32>
+  }
+
+  func.func public @main(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
+    %results = stablehlo.composite "jit_eltwise_add.my_add" %arg0, %arg1 {
+        decomposition = @add_impl
+    } : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+    // CHECK: %[[C:.*]] = call @add_impl [[C:.*]]
+    return %results : tensor<13x21x3xf32>
+  }
+}