Stable HLO runtime tests

test/ttmlir/Silicon/StableHLO/binary/2add_op.mlir

@@ -0,0 +1,17 @@
+// REQUIRES: stablehlo
+// RUN: rm -rf %t.ttnn
+// RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \
+// RUN:     ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" > %t.mlir
+// RUN:     ttmlir-translate --ttnn-to-flatbuffer %t.mlir > %t.ttnn
+
+#any_device = #tt.operand_constraint<dram|l1|scalar|tile|any_device|any_device_tile>
+#any_device_tile = #tt.operand_constraint<dram|l1|tile|any_device_tile>
+
+module @jit_eltwise_add attributes {} {
+  func.func public @test_add(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> {
+    %0 = stablehlo.add %arg0, %arg1 : tensor<64x128xf32>
+    // CHECK: %[[C:.*]] = tensor.empty[[C:.*]]
+    // CHECK: %[[C:.*]] = "ttir.add"[[C:.*]]
+    return %0 : tensor<64x128xf32>
+  }
+}

test/ttmlir/Silicon/StableHLO/binary/3add_op.mlir

@@ -0,0 +1,17 @@
+// REQUIRES: stablehlo
+// RUN: rm -rf %t.ttnn
+// RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \
+// RUN:     ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" > %t.mlir | \
+// RUN:     ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn
+
+#any_device = #tt.operand_constraint<dram|l1|scalar|tile|any_device|any_device_tile>
+#any_device_tile = #tt.operand_constraint<dram|l1|tile|any_device_tile>
+
+module @jit_eltwise_add attributes {} {
+  func.func public @test_add(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> {
+    %0 = stablehlo.add %arg0, %arg1 : tensor<64x128xf32>
+    // CHECK: %[[C:.*]] = tensor.empty[[C:.*]]
+    // CHECK: %[[C:.*]] = "ttir.add"[[C:.*]]
+    return %0 : tensor<64x128xf32>
+  }
+}

test/ttmlir/Silicon/StableHLO/binary/4add_op.mlir

@@ -0,0 +1,16 @@
+// REQUIRES: stablehlo
+// RUN: rm -rf %t.ttnn
+// RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" > %t.mlir
+// RUN:     ttmlir-translate --ttnn-to-flatbuffer %t.mlir > %t.ttnn
+
+#any_device = #tt.operand_constraint<dram|l1|scalar|tile|any_device|any_device_tile>
+#any_device_tile = #tt.operand_constraint<dram|l1|tile|any_device_tile>
+
+module @jit_eltwise_add attributes {} {
+  func.func public @test_add(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> {
+    %0 = stablehlo.add %arg0, %arg1 : tensor<64x128xf32>
+    // CHECK: %[[C:.*]] = tensor.empty[[C:.*]]
+    // CHECK: %[[C:.*]] = "ttir.add"[[C:.*]]
+    return %0 : tensor<64x128xf32>
+  }
+}

test/ttmlir/Silicon/StableHLO/binary/add_op.mlir

@@ -0,0 +1,17 @@
+// REQUIRES: stablehlo
+// RUN: rm -rf %t.ttnn
+// RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s > %t.ttir
+// RUN:     ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" %t.ttir > %t.mlir
+// RUN:     ttmlir-translate --ttnn-to-flatbuffer %t.mlir > %t.ttnn
+
+#any_device = #tt.operand_constraint<dram|l1|scalar|tile|any_device|any_device_tile>
+#any_device_tile = #tt.operand_constraint<dram|l1|tile|any_device_tile>
+
+module @jit_eltwise_add attributes {} {
+  func.func public @test_add(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> {
+    %0 = stablehlo.add %arg0, %arg1 : tensor<64x128xf32>
+    // CHECK: %[[C:.*]] = tensor.empty[[C:.*]]
+    // CHECK: %[[C:.*]] = "ttir.add"[[C:.*]]
+    return %0 : tensor<64x128xf32>
+  }
+}

test/ttmlir/Silicon/StableHLO/binary/multiply_op.mlir

@@ -0,0 +1,15 @@
+// REQUIRES: stablehlo
+// RUN: ttmlir-opt --stablehlo-to-ttir-pipeline --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" %s | \
+// RUN:     ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn
+
+#any_device = #tt.operand_constraint<dram|l1|scalar|tile|any_device|any_device_tile>
+#any_device_tile = #tt.operand_constraint<dram|l1|tile|any_device_tile>
+
+module @jit_eltwise_multiply attributes {} {
+  func.func public @test_multiply(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
+    %0 = stablehlo.multiply %arg0, %arg1 : tensor<13x21x3xf32>
+    // CHECK: %[[C:.*]] = tensor.empty[[C:.*]]
+    // CHECK: %[[C:.*]] = "ttir.multiply"[[C:.*]]
+    return %0 : tensor<13x21x3xf32>
+  }
+}