From b6f73a3efcbfed0e57e0cde9669385f97401583c Mon Sep 17 00:00:00 2001 From: uazizTT Date: Wed, 30 Oct 2024 21:20:23 -0400 Subject: [PATCH] Refactoring. --- include/ttmlir/Dialect/TTIR/IR/TTIROps.td | 2 +- include/ttmlir/Target/TTNN/program.fbs | 2 +- lib/Target/TTNN/TTNNToFlatbuffer.cpp | 2 +- runtime/include/tt/runtime/detail/ttnn.h | 1 - runtime/lib/ttnn/operations/eltwise/unary.cpp | 2 +- .../StableHLOToTTIR/unary/{logistic_op.mlir => logit_op.mlir} | 4 ++-- test/ttmlir/Silicon/TTNN/perf_unit/test_perf_ceil.mlir | 4 ++-- test/ttmlir/Silicon/TTNN/perf_unit/test_perf_cosine.mlir | 4 ++-- test/ttmlir/Silicon/TTNN/perf_unit/test_perf_log.mlir | 4 ++-- test/ttmlir/Silicon/TTNN/perf_unit/test_perf_logit.mlir | 4 ++-- test/ttmlir/Silicon/TTNN/perf_unit/test_perf_sine.mlir | 4 ++-- test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tan.mlir | 4 ++-- test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tanh.mlir | 4 ++-- 13 files changed, 20 insertions(+), 21 deletions(-) rename test/ttmlir/Conversion/StableHLOToTTIR/unary/{logistic_op.mlir => logit_op.mlir} (79%) diff --git a/include/ttmlir/Dialect/TTIR/IR/TTIROps.td b/include/ttmlir/Dialect/TTIR/IR/TTIROps.td index 52b5c7a44..c4bf1b395 100644 --- a/include/ttmlir/Dialect/TTIR/IR/TTIROps.td +++ b/include/ttmlir/Dialect/TTIR/IR/TTIROps.td @@ -232,7 +232,7 @@ def TTIR_LogitOp: TTIR_ElementwiseUnaryOp<"logit"> { } def TTIR_TanOp: TTIR_ElementwiseUnaryOp<"tan"> { - let summary = "Eltwise logistic op."; + let summary = "Eltwise tan op."; let description = [{ Eltwise tan operation. }]; diff --git a/include/ttmlir/Target/TTNN/program.fbs b/include/ttmlir/Target/TTNN/program.fbs index 2b84bddcb..45179f716 100644 --- a/include/ttmlir/Target/TTNN/program.fbs +++ b/include/ttmlir/Target/TTNN/program.fbs @@ -86,7 +86,7 @@ enum EltwiseOpType: uint32 { Ceil = 25, Sin = 26, Cos = 27, - Logistic = 28, + Logit = 28, Tan = 29, Tanh = 30, Log = 31 diff --git a/lib/Target/TTNN/TTNNToFlatbuffer.cpp b/lib/Target/TTNN/TTNNToFlatbuffer.cpp index 5dece94c2..0a1a92c1e 100644 --- a/lib/Target/TTNN/TTNNToFlatbuffer.cpp +++ b/lib/Target/TTNN/TTNNToFlatbuffer.cpp @@ -349,7 +349,7 @@ createEltwiseOp(FlatbufferObjectCache &cache, EltwiseOp op) { } else if constexpr (std::is_same_v) { type = ::tt::target::ttnn::EltwiseOpType::Sin; } else if constexpr (std::is_same_v) { - type = ::tt::target::ttnn::EltwiseOpType::Logistic; + type = ::tt::target::ttnn::EltwiseOpType::Logit; } else if constexpr (std::is_same_v) { type = ::tt::target::ttnn::EltwiseOpType::Tan; } else if constexpr (std::is_same_v) { diff --git a/runtime/include/tt/runtime/detail/ttnn.h b/runtime/include/tt/runtime/detail/ttnn.h index 10f133636..0fdfdbddf 100644 --- a/runtime/include/tt/runtime/detail/ttnn.h +++ b/runtime/include/tt/runtime/detail/ttnn.h @@ -57,7 +57,6 @@ #include "ttnn/operations/eltwise/binary/binary.hpp" #include "ttnn/operations/eltwise/binary/binary_composite.hpp" #include "ttnn/operations/eltwise/unary/unary.hpp" -#include "ttnn/operations/eltwise/unary/unary_composite.hpp" #include "ttnn/operations/embedding/embedding.hpp" #include "ttnn/operations/matmul/matmul.hpp" #include "ttnn/operations/normalization/softmax/softmax.hpp" diff --git a/runtime/lib/ttnn/operations/eltwise/unary.cpp b/runtime/lib/ttnn/operations/eltwise/unary.cpp index 39e07ad40..c7425619f 100644 --- a/runtime/lib/ttnn/operations/eltwise/unary.cpp +++ b/runtime/lib/ttnn/operations/eltwise/unary.cpp @@ -124,7 +124,7 @@ void run(const ::tt::target::ttnn::EltwiseOp *op, ProgramContext &context) { runEltwiseUnaryOP(op, tensorPool, ::ttnn::reciprocal); break; } - case ::tt::target::ttnn::EltwiseOpType::Logistic: { + case ::tt::target::ttnn::EltwiseOpType::Logit: { runEltwiseUnaryOP(op, tensorPool, ::ttnn::sigmoid); break; } diff --git a/test/ttmlir/Conversion/StableHLOToTTIR/unary/logistic_op.mlir b/test/ttmlir/Conversion/StableHLOToTTIR/unary/logit_op.mlir similarity index 79% rename from test/ttmlir/Conversion/StableHLOToTTIR/unary/logistic_op.mlir rename to test/ttmlir/Conversion/StableHLOToTTIR/unary/logit_op.mlir index 66bed61c6..dc7003e7c 100644 --- a/test/ttmlir/Conversion/StableHLOToTTIR/unary/logistic_op.mlir +++ b/test/ttmlir/Conversion/StableHLOToTTIR/unary/logit_op.mlir @@ -1,7 +1,7 @@ // REQUIRES: stablehlo // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | FileCheck %s -module @jit_eltwise_logistic attributes {} { - func.func public @test_logistic(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> { +module @jit_eltwise_logit attributes {} { + func.func public @test_logit(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> { %0 = stablehlo.logistic %arg0 : tensor<13x21x3xf32> // CHECK: [[VAL0:%[0-9]+]] = tensor.empty() : [[TENSOR_SIZE:tensor<[0-9]+x[0-9]+x[0-9]+xf[0-9]+>]] // CHECK: [[VAL1:%[0-9]+]] = "ttir.logit"(%arg0, [[VAL0]]) <{operandSegmentSizes = array, operand_constraints = [#any_device_tile, #any_device_tile]}> : ([[TENSOR_SIZE]], [[TENSOR_SIZE]]) -> [[TENSOR_SIZE]] diff --git a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_ceil.mlir b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_ceil.mlir index c31c789f4..2e7f55428 100644 --- a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_ceil.mlir +++ b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_ceil.mlir @@ -5,9 +5,9 @@ #any_device_tile = #tt.operand_constraint func.func @ceil(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { - // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] %0 = tensor.empty() : tensor<64x128xf32> - // CHECK: %[[C:.*]] = "ttnn.ceil"[[C:.*]] + // CHECK: [[VAL0:%[0-9]+]] = "ttnn.empty"(%{{[0-9]+}}) + // CHECK: %{{[0-9]+}} = "ttnn.ceil"(%{{[0-9]+}}, [[VAL0]]) %1 = "ttir.ceil"(%arg0, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> return %1 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_cosine.mlir b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_cosine.mlir index 91a7fea47..ede823439 100644 --- a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_cosine.mlir +++ b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_cosine.mlir @@ -5,9 +5,9 @@ #any_device_tile = #tt.operand_constraint func.func @cosine(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { - // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] %0 = tensor.empty() : tensor<64x128xf32> - // CHECK: %[[C:.*]] = "ttnn.cos"[[C:.*]] + // CHECK: [[VAL0:%[0-9]+]] = "ttnn.empty"(%{{[0-9]+}}) + // CHECK: %{{[0-9]+}} = "ttnn.cos"(%{{[0-9]+}}, [[VAL0]]) %1 = "ttir.cos"(%arg0, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> return %1 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_log.mlir b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_log.mlir index 9656ba567..b3de1bba4 100644 --- a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_log.mlir +++ b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_log.mlir @@ -5,9 +5,9 @@ #any_device_tile = #tt.operand_constraint func.func @log(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { - // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] %0 = tensor.empty() : tensor<64x128xf32> - // CHECK: %[[C:.*]] = "ttnn.log"[[C:.*]] + // CHECK: [[VAL0:%[0-9]+]] = "ttnn.empty"(%{{[0-9]+}}) + // CHECK: %{{[0-9]+}} = "ttnn.log"(%{{[0-9]+}}, [[VAL0]]) %1 = "ttir.log"(%arg0, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> return %1 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_logit.mlir b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_logit.mlir index 554a483db..9ba7f87a2 100644 --- a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_logit.mlir +++ b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_logit.mlir @@ -5,9 +5,9 @@ #any_device_tile = #tt.operand_constraint func.func @logit(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { - // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] %0 = tensor.empty() : tensor<64x128xf32> - // CHECK: %[[C:.*]] = "ttnn.logit"[[C:.*]] + // CHECK: [[VAL0:%[0-9]+]] = "ttnn.empty"(%{{[0-9]+}}) + // CHECK: %{{[0-9]+}} = "ttnn.logit"(%{{[0-9]+}}, [[VAL0]]) %1 = "ttir.logit"(%arg0, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> return %1 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_sine.mlir b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_sine.mlir index e72d57ffa..36f71d8e6 100644 --- a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_sine.mlir +++ b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_sine.mlir @@ -5,9 +5,9 @@ #any_device_tile = #tt.operand_constraint func.func @sine(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { - // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] %0 = tensor.empty() : tensor<64x128xf32> - // CHECK: %[[C:.*]] = "ttnn.sin"[[C:.*]] + // CHECK: [[VAL0:%[0-9]+]] = "ttnn.empty"(%{{[0-9]+}}) + // CHECK: %{{[0-9]+}} = "ttnn.sin"(%{{[0-9]+}}, [[VAL0]]) %1 = "ttir.sin"(%arg0, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> return %1 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tan.mlir b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tan.mlir index 0cea8df02..aa7b97298 100644 --- a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tan.mlir +++ b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tan.mlir @@ -5,9 +5,9 @@ #any_device_tile = #tt.operand_constraint func.func @tan(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { - // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] %0 = tensor.empty() : tensor<64x128xf32> - // CHECK: %[[C:.*]] = "ttnn.tan"[[C:.*]] + // CHECK: [[VAL0:%[0-9]+]] = "ttnn.empty"(%{{[0-9]+}}) + // CHECK: %{{[0-9]+}} = "ttnn.tan"(%{{[0-9]+}}, [[VAL0]]) %1 = "ttir.tan"(%arg0, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> return %1 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tanh.mlir b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tanh.mlir index 34de25c50..ecb7266c9 100644 --- a/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tanh.mlir +++ b/test/ttmlir/Silicon/TTNN/perf_unit/test_perf_tanh.mlir @@ -5,9 +5,9 @@ #any_device_tile = #tt.operand_constraint func.func @tanh(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { - // CHECK: %[[C:.*]] = "ttnn.empty"[[C:.*]] %0 = tensor.empty() : tensor<64x128xf32> - // CHECK: %[[C:.*]] = "ttnn.tanh"[[C:.*]] + // CHECK: [[VAL0:%[0-9]+]] = "ttnn.empty"(%{{[0-9]+}}) + // CHECK: %{{[0-9]+}} = "ttnn.tanh"(%{{[0-9]+}}, [[VAL0]]) %1 = "ttir.tanh"(%arg0, %0) <{operandSegmentSizes = array, operand_constraints = [#any_device, #any_device]}> : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> return %1 : tensor<64x128xf32> }