diff --git a/include/ttmlir/Dialect/TTIR/IR/TTIROps.td b/include/ttmlir/Dialect/TTIR/IR/TTIROps.td index 075821fc8f..1b7189e6af 100644 --- a/include/ttmlir/Dialect/TTIR/IR/TTIROps.td +++ b/include/ttmlir/Dialect/TTIR/IR/TTIROps.td @@ -269,13 +269,6 @@ class TTIR_ElementwiseBinaryOp traits = []> : ]; } -def TTIR_SubtractOp : TTIR_ElementwiseBinaryOp<"subtract"> { - let summary = "Eltwise subtract."; - let description = [{ - Eltwise subtract operation. - }]; -} - def TTIR_EqualOp : TTIR_ElementwiseBinaryOp<"eq"> { let summary = "Eltwise equal to."; let description = [{ @@ -344,6 +337,26 @@ def TTIR_MaximumOp : TTIR_ElementwiseBinaryOp<"maximum"> { }]; } +def TTIR_MinimumOp : TTIR_ElementwiseBinaryOp<"minimum"> { + let summary = "Eltwise minimum OP."; + let description = [{ + Calculates minimum of input tensors' values element-wise and stores result + in output tensor. + + Example: + %lhs: [[3, 2, 7], [1, 4, 4]] + %rhs: [[1, 4, 2], [1, 2, 3]] + "ttir.minimum"(%lhs, %rhs, %out) -> %out: [[1, 2, 2], [1, 2, 3]] + }]; +} + +def TTIR_SubtractOp : TTIR_ElementwiseBinaryOp<"subtract"> { + let summary = "Eltwise subtract."; + let description = [{ + Eltwise subtract operation. + }]; +} + class TTIR_ReductionOp traits = []> : TTIR_DPSOp { diff --git a/include/ttmlir/Dialect/TTNN/IR/TTNNOps.td b/include/ttmlir/Dialect/TTNN/IR/TTNNOps.td index ea1ded85ea..9cd03ffe86 100644 --- a/include/ttmlir/Dialect/TTNN/IR/TTNNOps.td +++ b/include/ttmlir/Dialect/TTNN/IR/TTNNOps.td @@ -246,20 +246,6 @@ def TTNN_DivOp : TTNN_ElementwiseBinaryOp<"div"> { }]; } -def TTNN_SubtractOp : TTNN_ElementwiseBinaryOp<"subtract"> { - let summary = "Eltwise subtract."; - let description = [{ - Eltwise subtract operation. - }]; -} - -def TTNN_MultiplyOp : TTNN_ElementwiseBinaryOp<"multiply"> { - let summary = "Eltwise multiply."; - let description = [{ - Eltwise multiply operation. - }]; -} - def TTNN_EqualOp : TTNN_ElementwiseBinaryOp<"eq"> { let summary = "Eltwise equal to."; let description = [{ @@ -328,6 +314,33 @@ def TTNN_MaximumOp : TTNN_ElementwiseBinaryOp<"maximum"> { }]; } +def TTNN_MinimumOp : TTNN_ElementwiseBinaryOp<"minimum"> { + let summary = "Eltwise minimum OP."; + let description = [{ + Calculates minimum of input tensors' values element-wise and stores result + in output tensor. + + Example: + %lhs: [[3, 2, 7], [1, 4, 4]] + %rhs: [[1, 4, 2], [1, 2, 3]] + "ttnn.minimum"(%lhs, %rhs, %out) -> %out: [[1, 2, 2], [1, 2, 3]] + }]; +} + +def TTNN_MultiplyOp : TTNN_ElementwiseBinaryOp<"multiply"> { + let summary = "Eltwise multiply."; + let description = [{ + Eltwise multiply operation. + }]; +} + +def TTNN_SubtractOp : TTNN_ElementwiseBinaryOp<"subtract"> { + let summary = "Eltwise subtract."; + let description = [{ + Eltwise subtract operation. + }]; +} + class TTNN_ReductionOp traits = []> : TTNN_Op { let summary = "Reduction op."; let description = [{ diff --git a/include/ttmlir/Target/TTNN/program.fbs b/include/ttmlir/Target/TTNN/program.fbs index 652d7028ec..d1f8a06abe 100644 --- a/include/ttmlir/Target/TTNN/program.fbs +++ b/include/ttmlir/Target/TTNN/program.fbs @@ -82,6 +82,7 @@ enum EltwiseOpType: uint32 { LogicalOr = 21, LogicalNot = 22, Cbrt = 23, + Minimum = 24, } table EltwiseOp { diff --git a/lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp b/lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp index ddf6eca0cd..639f36c637 100644 --- a/lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp +++ b/lib/Conversion/StableHLOToTTIR/StableHLOToTTIRPatterns.cpp @@ -879,15 +879,17 @@ void addElementwiseBinaryOpsConversionPatterns(MLIRContext *ctx, patterns.add>(typeConverter, ctx); - patterns.add>(typeConverter, - ctx); - patterns.add>(typeConverter, ctx); patterns.add>(typeConverter, ctx); patterns.add>(typeConverter, ctx); + patterns.add>(typeConverter, ctx); + patterns.add>(typeConverter, ctx); + patterns.add>(typeConverter, + ctx); } void addReduceOpsConversionPatterns(MLIRContext *ctx, diff --git a/lib/Conversion/TTIRToTTNN/TTIRToTTNN.cpp b/lib/Conversion/TTIRToTTNN/TTIRToTTNN.cpp index 95834070c4..315051c0c3 100644 --- a/lib/Conversion/TTIRToTTNN/TTIRToTTNN.cpp +++ b/lib/Conversion/TTIRToTTNN/TTIRToTTNN.cpp @@ -820,6 +820,7 @@ void populateTTIRToTTNNPatterns(MLIRContext *ctx, RewritePatternSet &patterns, ElementwiseOpConversionPattern, ElementwiseOpConversionPattern, ElementwiseOpConversionPattern, + ElementwiseOpConversionPattern, ElementwiseOpConversionPattern, ElementwiseOpConversionPattern, ElementwiseOpConversionPattern, diff --git a/lib/Conversion/TTNNToEmitC/TTNNToEmitC.cpp b/lib/Conversion/TTNNToEmitC/TTNNToEmitC.cpp index 226efcf781..b7d4c30167 100644 --- a/lib/Conversion/TTNNToEmitC/TTNNToEmitC.cpp +++ b/lib/Conversion/TTNNToEmitC/TTNNToEmitC.cpp @@ -639,6 +639,7 @@ void populateTTNNToEmitCPatterns(mlir::MLIRContext *ctx, DefaultOpConversionPattern, DefaultOpConversionPattern, DefaultOpConversionPattern, + DefaultOpConversionPattern, DefaultOpConversionPattern>(typeConverter, ctx); // Tensor manipulation ops diff --git a/lib/Target/TTNN/TTNNToFlatbuffer.cpp b/lib/Target/TTNN/TTNNToFlatbuffer.cpp index 71d146f072..5cf0a0f02b 100644 --- a/lib/Target/TTNN/TTNNToFlatbuffer.cpp +++ b/lib/Target/TTNN/TTNNToFlatbuffer.cpp @@ -325,6 +325,8 @@ createEltwiseOp(FlatbufferObjectCache &cache, EltwiseOp op) { type = ::tt::target::ttnn::EltwiseOpType::LessThan; } else if constexpr (std::is_same_v) { type = ::tt::target::ttnn::EltwiseOpType::Maximum; + } else if constexpr (std::is_same_v) { + type = ::tt::target::ttnn::EltwiseOpType::Minimum; } else if constexpr (std::is_same_v) { type = ::tt::target::ttnn::EltwiseOpType::Relu; } else if constexpr (std::is_same_v) { @@ -568,6 +570,10 @@ emitTTNNOperation(FlatbufferObjectCache &cache, Operation *op, return createOperation(cache, createEltwiseOp(cache, maximumOp), debugString); } + if (auto minimumOp = dyn_cast(op); minimumOp) { + return createOperation(cache, createEltwiseOp(cache, minimumOp), + debugString); + } if (auto reluOp = dyn_cast(op); reluOp) { return createOperation(cache, createEltwiseOp(cache, reluOp), debugString); } diff --git a/runtime/lib/ttnn/operations/eltwise/binary.cpp b/runtime/lib/ttnn/operations/eltwise/binary.cpp index ce56e9ac68..2e35f20ed8 100644 --- a/runtime/lib/ttnn/operations/eltwise/binary.cpp +++ b/runtime/lib/ttnn/operations/eltwise/binary.cpp @@ -122,6 +122,10 @@ void run(const ::tt::target::ttnn::EltwiseOp *op, ProgramContext &context) { runEltwiseBinaryCompositeOP(op, tensorPool, ::ttnn::maximum); break; } + case ::tt::target::ttnn::EltwiseOpType::Minimum: { + runEltwiseBinaryCompositeOP(op, tensorPool, ::ttnn::minimum); + break; + } default: throw std::invalid_argument("Unsupported Eltwise Binary operation"); } diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/add_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/add_op.mlir similarity index 100% rename from test/ttmlir/Conversion/StableHLOToTTIR/add_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/binary/add_op.mlir diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/compare_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/compare_op.mlir similarity index 100% rename from test/ttmlir/Conversion/StableHLOToTTIR/compare_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/binary/compare_op.mlir diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/concat_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/concat_op.mlir similarity index 100% rename from test/ttmlir/Conversion/StableHLOToTTIR/concat_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/binary/concat_op.mlir diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/divide_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/divide_op.mlir similarity index 100% rename from test/ttmlir/Conversion/StableHLOToTTIR/divide_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/binary/divide_op.mlir diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/logical_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/logical_op.mlir similarity index 100% rename from test/ttmlir/Conversion/StableHLOToTTIR/logical_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/binary/logical_op.mlir diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/maximum_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/maximum_op.mlir similarity index 100% rename from test/ttmlir/Conversion/StableHLOToTTIR/maximum_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/binary/maximum_op.mlir diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/binary/minimum_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/minimum_op.mlir new file mode 100644 index 0000000000..bd9c241847 --- /dev/null +++ b/test/ttmlir/Conversion/StableHLOToTTIR/binary/minimum_op.mlir @@ -0,0 +1,15 @@ +// REQUIRES: stablehlo +// RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | FileCheck %s +module @jit_eltwise_minimum attributes {} { + func.func public @test_minimum(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> { + // CHECK: %[[C:.*]] = tensor.empty() + // CHECK-SAME: [[TENSOR:tensor<13x21x3xf32>]] + // CHECK: %[[C:.*]] = "ttir.minimum" + // CHECK-SAME: [[TENSOR]] + // CHECK-SAME: [[TENSOR]] + // CHECK-SAME: [[TENSOR]] + // CHECK-SAME: -> [[TENSOR]] + %0 = stablehlo.minimum %arg0, %arg1 : tensor<13x21x3xf32> + return %0 : tensor<13x21x3xf32> + } +} diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/multiply_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/multiply_op.mlir similarity index 100% rename from test/ttmlir/Conversion/StableHLOToTTIR/multiply_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/binary/multiply_op.mlir diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/subtract_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/binary/subtract_op.mlir similarity index 100% rename from test/ttmlir/Conversion/StableHLOToTTIR/subtract_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/binary/subtract_op.mlir diff --git a/test/ttmlir/Dialect/TTNN/eltwise/binary/minimum/simple_minimum.mlir b/test/ttmlir/Dialect/TTNN/eltwise/binary/minimum/simple_minimum.mlir new file mode 100644 index 0000000000..8ebdfe0a47 --- /dev/null +++ b/test/ttmlir/Dialect/TTNN/eltwise/binary/minimum/simple_minimum.mlir @@ -0,0 +1,11 @@ +// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline %s | FileCheck %s +#any_device = #tt.operand_constraint +module attributes {} { + func.func @forward(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] + %0 = tensor.empty() : tensor<64x128xf32> + // CHECK: %[[C:.*]] = "ttnn.minimum"[[C:.*]] + %1 = "ttir.minimum"(%arg0, %arg1, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> + return %1 : tensor<64x128xf32> + } +} diff --git a/test/ttmlir/Silicon/TTNN/simple_eltwise.mlir b/test/ttmlir/Silicon/TTNN/simple_eltwise.mlir index 93a706086c..7a8736b9d3 100644 --- a/test/ttmlir/Silicon/TTNN/simple_eltwise.mlir +++ b/test/ttmlir/Silicon/TTNN/simple_eltwise.mlir @@ -20,6 +20,19 @@ func.func @div(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<6 return %1 : tensor<64x128xf32> } +func.func @minimum(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: %[[C:.*]] = "ttnn.empty" + // CHECK-SAME: [[TENSOR:tensor<64x128xf32,]] + %0 = tensor.empty() : tensor<64x128xf32> + // CHECK: %[[C:.*]] = "ttnn.minimum" + // CHECK-SAME: [[TENSOR]] + // CHECK-SAME: [[TENSOR]] + // CHECK-SAME: [[TENSOR]] + // CHECK-SAME: -> [[TENSOR]] + %1 = "ttir.minimum"(%arg0, %arg1, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> + return %1 : tensor<64x128xf32> +} + func.func @multiply(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] %0 = tensor.empty() : tensor<64x128xf32>