From ff509d666ea7fad3f9c7dbf9120b29903a4348e5 Mon Sep 17 00:00:00 2001 From: Muhammad Asif Manzoor Date: Thu, 14 Nov 2024 17:52:35 +0000 Subject: [PATCH] h --- .../Silicon/StableHLO/Binary/add_op.mlir | 12 +++++- .../Silicon/StableHLO/Binary/compare_op.mlir | 42 ++++++++++++++++++- .../Silicon/StableHLO/Binary/concat_op.mlir | 41 +++++++++++++++++- .../Silicon/StableHLO/Binary/divide_op.mlir | 12 +++++- .../Silicon/StableHLO/Binary/logical_op.mlir | 29 ++++++++++++- .../Silicon/StableHLO/Binary/maximum_op.mlir | 12 +++++- .../Silicon/StableHLO/Binary/minimum_op.mlir | 12 +++++- .../Silicon/StableHLO/Binary/multiply_op.mlir | 18 +++++--- .../StableHLO/Binary/remainder_op.mlir | 12 +++++- .../Silicon/StableHLO/Binary/subtract_op.mlir | 12 +++++- .../StableHLO/Constant/constant_bf16.mlir | 21 +++++++++- .../StableHLO/Constant/constant_bool.mlir | 32 +++++++++++--- .../StableHLO/Constant/constant_f32.mlir | 35 ++++++++++++---- .../StableHLO/Constant/constant_i16.mlir | 27 +++++++++--- .../StableHLO/Constant/constant_i32.mlir | 21 +++++++++- .../StableHLO/Constant/constant_i64.mlir | 21 +++++++++- .../Silicon/StableHLO/Unary/absolute_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/cbrt_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/ceil_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/cosine_op.mlir | 11 ++++- .../Unary/exponential_minus_one_op.mlir | 11 ++++- .../StableHLO/Unary/exponential_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/floor_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/isfinite_op.mlir | 11 ++++- .../StableHLO/Unary/log_plus_one_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/logical_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/negate_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/rsqrt_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/sign_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/sine_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/sqrt_op.mlir | 11 ++++- .../Silicon/StableHLO/Unary/tranpose_op.mlir | 10 ++++- .../Silicon/StableHLO/broadcast_op.mlir | 9 +++- .../Silicon/StableHLO/composite_op.mlir | 10 ++++- test/ttmlir/Silicon/StableHLO/conv2d_op.mlir | 6 ++- test/ttmlir/Silicon/StableHLO/convert_op.mlir | 18 +++++++- .../Silicon/StableHLO/dot_general_op.mlir | 12 +++++- .../StableHLO/get_dimension_size_op.mlir | 9 +++- .../Silicon/StableHLO/maxpool2d_op.mlir | 6 ++- .../Silicon/StableHLO/mnist_inference.mlir | 7 +++- .../Silicon/StableHLO/reduce_add_op.mlir | 13 +++++- .../Silicon/StableHLO/reduce_maximum_op.mlir | 13 +++++- test/ttmlir/Silicon/StableHLO/rehsape_op.mlir | 10 ++++- .../Silicon/StableHLO/scalar_add_op.mlir | 12 +++++- test/ttmlir/Silicon/StableHLO/slice_op.mlir | 14 ++++++- 45 files changed, 567 insertions(+), 106 deletions(-) diff --git a/test/ttmlir/Silicon/StableHLO/Binary/add_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/add_op.mlir index e31d97eb65..94a36f77c8 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/add_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/add_op.mlir @@ -1,11 +1,19 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_add attributes {} { func.func public @test_add(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.add + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.add %arg0, %arg1 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Binary/compare_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/compare_op.mlir index e9716401dd..da30dd301c 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/compare_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/compare_op.mlir @@ -1,36 +1,74 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_compare attributes {} { func.func public @test_eq(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.eq + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.compare EQ, %arg0, %arg1 : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xi1> return %0 : tensor<64x128xi1> } func.func public @test_ne(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.ne + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.compare NE, %arg0, %arg1 : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xi1> return %0 : tensor<64x128xi1> } func.func public @test_ge(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.ge + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.compare GE, %arg0, %arg1 : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xi1> return %0 : tensor<64x128xi1> } func.func public @test_gt(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.gt + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.compare GT, %arg0, %arg1 : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xi1> return %0 : tensor<64x128xi1> } func.func public @test_le(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.le + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.compare LE, %arg0, %arg1 : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xi1> return %0 : tensor<64x128xi1> } func.func public @test_lt(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.lt + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.compare LT, %arg0, %arg1 : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xi1> return %0 : tensor<64x128xi1> } diff --git a/test/ttmlir/Silicon/StableHLO/Binary/concat_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/concat_op.mlir index c7c17fbac1..cad2b894b0 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/concat_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/concat_op.mlir @@ -1,11 +1,20 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_concat attributes {} { func.func public @test_concat_0(%arg0: tensor<32x32xf32>, %arg1: tensor<64x32xf32>) -> tensor<96x32xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.concat + // CHECK-SAME: dim = 0 + // CHECK-SAME: tensor<32x32xf32, + // CHECK-SAME: tensor<64x32xf32, + // CHECK-SAME: tensor<96x32xf32, + // CHECK-SAME: -> tensor<96x32xf32, %0 = "stablehlo.concatenate"(%arg0, %arg1) { dimension = 0 : i64 } : (tensor<32x32xf32>, tensor<64x32xf32>) -> tensor<96x32xf32> @@ -13,6 +22,13 @@ module @jit_concat attributes {} { } func.func public @test_concat_1(%arg0: tensor<32x32xf32>, %arg1: tensor<32x64xf32>) -> tensor<32x96xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.concat + // CHECK-SAME: dim = 1 + // CHECK-SAME: tensor<32x32xf32, + // CHECK-SAME: tensor<32x64xf32, + // CHECK-SAME: tensor<32x96xf32, + // CHECK-SAME: -> tensor<32x96xf32, %0 = "stablehlo.concatenate"(%arg0, %arg1) { dimension = 1 : i64 } : (tensor<32x32xf32>, tensor<32x64xf32>) -> tensor<32x96xf32> @@ -21,6 +37,13 @@ module @jit_concat attributes {} { func.func public @test_concat_2(%arg0: tensor<128x64xf32>, %arg1: tensor<128x96xf32>) -> tensor<128x160xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.concat + // CHECK-SAME: dim = 1 + // CHECK-SAME: tensor<128x64xf32, + // CHECK-SAME: tensor<128x96xf32, + // CHECK-SAME: tensor<128x160xf32, + // CHECK-SAME: -> tensor<128x160xf32, %0 = "stablehlo.concatenate"(%arg0, %arg1) { dimension = 1 : i64 } : (tensor<128x64xf32>, tensor<128x96xf32>) -> tensor<128x160xf32> @@ -28,6 +51,13 @@ module @jit_concat attributes {} { } func.func public @test_concat_3(%arg0: tensor<64x32xf32>, %arg1: tensor<64x64xf32>) -> tensor<64x96xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.concat + // CHECK-SAME: dim = 1 + // CHECK-SAME: tensor<64x32xf32, + // CHECK-SAME: tensor<64x64xf32, + // CHECK-SAME: tensor<64x96xf32, + // CHECK-SAME: -> tensor<64x96xf32, %0 = "stablehlo.concatenate"(%arg0, %arg1) { dimension = 1 : i64 } : (tensor<64x32xf32>, tensor<64x64xf32>) -> tensor<64x96xf32> @@ -35,6 +65,13 @@ module @jit_concat attributes {} { } func.func public @test_concat_4(%arg0: tensor<32x32x32x32xf32>, %arg1: tensor<32x32x32x64xf32>) -> tensor<32x32x32x96xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.concat + // CHECK-SAME: dim = 3 + // CHECK-SAME: tensor<32x32x32x32xf32, + // CHECK-SAME: tensor<32x32x32x64xf32, + // CHECK-SAME: tensor<32x32x32x96xf32, + // CHECK-SAME: -> tensor<32x32x32x96xf32, %0 = "stablehlo.concatenate"(%arg0, %arg1) { dimension = 3 : i64 } : (tensor<32x32x32x32xf32>, tensor<32x32x32x64xf32>) -> tensor<32x32x32x96xf32> diff --git a/test/ttmlir/Silicon/StableHLO/Binary/divide_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/divide_op.mlir index c85cbf60e6..ad72cb1fff 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/divide_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/divide_op.mlir @@ -1,11 +1,19 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_divice attributes {} { func.func public @test_divide(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.div + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.divide %arg0, %arg1 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Binary/logical_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/logical_op.mlir index 82e94d687f..879ede5040 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/logical_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/logical_op.mlir @@ -1,17 +1,42 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_compare attributes {} { func.func public @logical_and(%arg0: tensor<64x128xi1>, %arg1: tensor<64x128xi1>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.logical_and + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.and %arg0, %arg1 : tensor<64x128xi1> return %0 : tensor<64x128xi1> } func.func public @logical_or(%arg0: tensor<64x128xi1>, %arg1: tensor<64x128xi1>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.logical_or + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.or %arg0, %arg1 : tensor<64x128xi1> return %0 : tensor<64x128xi1> } + + func.func public @logical_xor(%arg0: tensor<64x128xi1>, %arg1: tensor<64x128xi1>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.logical_xor + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, + %0 = stablehlo.xor %arg0, %arg1 : tensor<64x128xi1> + return %0 : tensor<64x128xi1> + } } diff --git a/test/ttmlir/Silicon/StableHLO/Binary/maximum_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/maximum_op.mlir index 72e0bd14c7..6c595d2efb 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/maximum_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/maximum_op.mlir @@ -1,11 +1,19 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_maximum attributes {} { func.func public @test_maximum(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.maximum + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.maximum %arg0, %arg1 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Binary/minimum_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/minimum_op.mlir index 7d7f3cac46..1dfce648bc 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/minimum_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/minimum_op.mlir @@ -1,11 +1,19 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_minimum attributes {} { func.func public @test_minimum(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.minimum + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.minimum %arg0, %arg1 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Binary/multiply_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/multiply_op.mlir index 95460ea6a4..0c491fa55f 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/multiply_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/multiply_op.mlir @@ -1,12 +1,20 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s 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> - return %0 : tensor<13x21x3xf32> + func.func public @test_multiply(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.multiply + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, + %0 = stablehlo.multiply %arg0, %arg1 : tensor<64x128xf32> + return %0 : tensor<64x128xf32> } } diff --git a/test/ttmlir/Silicon/StableHLO/Binary/remainder_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/remainder_op.mlir index 67a60c659a..f4e64095b6 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/remainder_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/remainder_op.mlir @@ -1,11 +1,19 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_remainder attributes {} { func.func public @test_remainder(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.remainder + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.remainder %arg0, %arg1 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Binary/subtract_op.mlir b/test/ttmlir/Silicon/StableHLO/Binary/subtract_op.mlir index 003c11dfb4..84804023a2 100644 --- a/test/ttmlir/Silicon/StableHLO/Binary/subtract_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Binary/subtract_op.mlir @@ -1,11 +1,19 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_subtract attributes {} { func.func public @test_subtract(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.subtract + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.subtract %arg0, %arg1 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Constant/constant_bf16.mlir b/test/ttmlir/Silicon/StableHLO/Constant/constant_bf16.mlir index bd438963c4..8331f766e7 100644 --- a/test/ttmlir/Silicon/StableHLO/Constant/constant_bf16.mlir +++ b/test/ttmlir/Silicon/StableHLO/Constant/constant_bf16.mlir @@ -1,21 +1,38 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_constant attributes {} { func.func public @test_bfloat16_scalar() -> tensor { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<1xbf16 %0 = stablehlo.constant dense<3.0> : tensor return %0 : tensor } + func.func public @test_bfloat16_scalar_empty() -> tensor { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<1xbf16 + %0 = stablehlo.constant dense<0.0> : tensor + return %0 : tensor + } + func.func public @test_bfloat16_empty() -> tensor<64x128xbf16> { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<64x128xbf16 %0 = stablehlo.constant dense<0.0> : tensor<64x128xbf16> return %0 : tensor<64x128xbf16> } func.func public @test_bfloat16_splat() -> tensor<64x128xbf16> { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<64x128xbf16 %0 = stablehlo.constant dense<3.0> : tensor<64x128xbf16> return %0 : tensor<64x128xbf16> } diff --git a/test/ttmlir/Silicon/StableHLO/Constant/constant_bool.mlir b/test/ttmlir/Silicon/StableHLO/Constant/constant_bool.mlir index 8fe2ceefc3..6d87380624 100644 --- a/test/ttmlir/Silicon/StableHLO/Constant/constant_bool.mlir +++ b/test/ttmlir/Silicon/StableHLO/Constant/constant_bool.mlir @@ -1,17 +1,39 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_constant attributes {} { func.func public @test_boolean_scalar() -> tensor { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 1.000000e+00 : f32 + // CHECK-SAME: -> tensor<1xbf16 %0 = stablehlo.constant dense : tensor return %0 : tensor } - func.func public @test_boolean_splat() -> tensor<64xi1> { - %0 = stablehlo.constant dense : tensor<64xi1> - return %0 : tensor<64xi1> + func.func public @test_boolean_scalar_empty() -> tensor { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<1xbf16 + %0 = stablehlo.constant dense : tensor + return %0 : tensor + } + + func.func public @test_boolean_empty() -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<64x128xbf16 + %0 = stablehlo.constant dense : tensor<64x128xi1> + return %0 : tensor<64x128xi1> + } + + func.func public @test_boolean_splat() -> tensor<64x128xi1> { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 1.000000e+00 : f32 + // CHECK-SAME: -> tensor<64x128xbf16 + %0 = stablehlo.constant dense : tensor<64x128xi1> + return %0 : tensor<64x128xi1> } } diff --git a/test/ttmlir/Silicon/StableHLO/Constant/constant_f32.mlir b/test/ttmlir/Silicon/StableHLO/Constant/constant_f32.mlir index 09ce2f18a8..0c6ebdc69f 100644 --- a/test/ttmlir/Silicon/StableHLO/Constant/constant_f32.mlir +++ b/test/ttmlir/Silicon/StableHLO/Constant/constant_f32.mlir @@ -1,22 +1,39 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_constant attributes {} { func.func public @test_float_scalar() -> tensor { - %0 = stablehlo.constant dense<0.3> : tensor + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<1xf32 + %0 = stablehlo.constant dense<3.0> : tensor return %0 : tensor } - func.func public @test_float_empty() -> tensor<64xf32> { - %0 = stablehlo.constant dense<0.0> : tensor<64xf32> - return %0 : tensor<64xf32> + func.func public @test_float_scalar_empty() -> tensor { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<1xf32 + %0 = stablehlo.constant dense<0.0> : tensor + return %0 : tensor + } + + func.func public @test_float_empty() -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<64x128xf32 + %0 = stablehlo.constant dense<0.0> : tensor<64x128xf32> + return %0 : tensor<64x128xf32> } - func.func public @test_float_splat() -> tensor<64xf32> { - %0 = stablehlo.constant dense<0.3> : tensor<64xf32> - return %0 : tensor<64xf32> + func.func public @test_float_splat() -> tensor<64x128xf32> { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<64x128xf32 + %0 = stablehlo.constant dense<3.0> : tensor<64x128xf32> + return %0 : tensor<64x128xf32> } } diff --git a/test/ttmlir/Silicon/StableHLO/Constant/constant_i16.mlir b/test/ttmlir/Silicon/StableHLO/Constant/constant_i16.mlir index 1b757844d5..2968dcc3ed 100644 --- a/test/ttmlir/Silicon/StableHLO/Constant/constant_i16.mlir +++ b/test/ttmlir/Silicon/StableHLO/Constant/constant_i16.mlir @@ -1,21 +1,38 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_constant attributes {} { - func.func public @test_int32_scalar() -> tensor { + func.func public @test_int16_scalar() -> tensor { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<1xi16 %0 = stablehlo.constant dense<3> : tensor return %0 : tensor } - func.func public @test_int32_empty() -> tensor<64x128xi16> { + func.func public @test_int16_scalar_empty() -> tensor { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<1xi16 + %0 = stablehlo.constant dense<0> : tensor + return %0 : tensor + } + + func.func public @test_int16_empty() -> tensor<64x128xi16> { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<64x128xi16 %0 = stablehlo.constant dense<0> : tensor<64x128xi16> return %0 : tensor<64x128xi16> } - func.func public @test_int32_splat() -> tensor<64x128xi16> { + func.func public @test_int16_splat() -> tensor<64x128xi16> { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<64x128xi16 %0 = stablehlo.constant dense<3> : tensor<64x128xi16> return %0 : tensor<64x128xi16> } diff --git a/test/ttmlir/Silicon/StableHLO/Constant/constant_i32.mlir b/test/ttmlir/Silicon/StableHLO/Constant/constant_i32.mlir index db88dc2d9c..554ab12b4b 100644 --- a/test/ttmlir/Silicon/StableHLO/Constant/constant_i32.mlir +++ b/test/ttmlir/Silicon/StableHLO/Constant/constant_i32.mlir @@ -1,21 +1,38 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_constant attributes {} { func.func public @test_int32_scalar() -> tensor { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<1xi32 %0 = stablehlo.constant dense<3> : tensor return %0 : tensor } + func.func public @test_int32_scalar_empty() -> tensor { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<1xi32 + %0 = stablehlo.constant dense<0> : tensor + return %0 : tensor + } + func.func public @test_int32_empty() -> tensor<64x128xi32> { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<64x128xi32 %0 = stablehlo.constant dense<0> : tensor<64x128xi32> return %0 : tensor<64x128xi32> } func.func public @test_int32_splat() -> tensor<64x128xi32> { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<64x128xi32 %0 = stablehlo.constant dense<3> : tensor<64x128xi32> return %0 : tensor<64x128xi32> } diff --git a/test/ttmlir/Silicon/StableHLO/Constant/constant_i64.mlir b/test/ttmlir/Silicon/StableHLO/Constant/constant_i64.mlir index ff1ece580e..62d42a0b1b 100644 --- a/test/ttmlir/Silicon/StableHLO/Constant/constant_i64.mlir +++ b/test/ttmlir/Silicon/StableHLO/Constant/constant_i64.mlir @@ -1,21 +1,38 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_constant attributes {} { func.func public @test_int64_scalar() -> tensor { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<1xi32 %0 = stablehlo.constant dense<3> : tensor return %0 : tensor } + func.func public @test_int64_scalar_empty() -> tensor { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<1xi32 + %0 = stablehlo.constant dense<0> : tensor + return %0 : tensor + } + func.func public @test_int64_empty() -> tensor<64x128xi64> { + // CHECK: ttnn.empty + // CHECK-SAME: -> tensor<64x128xi32 %0 = stablehlo.constant dense<0> : tensor<64x128xi64> return %0 : tensor<64x128xi64> } func.func public @test_int64_splat() -> tensor<64x128xi64> { + // CHECK: ttnn.full + // CHECK-SAME: fillValue = 3.000000e+00 : f32 + // CHECK-SAME: -> tensor<64x128xi32 %0 = stablehlo.constant dense<3> : tensor<64x128xi64> return %0 : tensor<64x128xi64> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/absolute_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/absolute_op.mlir index 81a3c70d84..19dd433d41 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/absolute_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/absolute_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_abs attributes {} { func.func public @test_abs(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.abs + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.abs %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/cbrt_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/cbrt_op.mlir index 3afff418a5..66b0a58ae0 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/cbrt_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/cbrt_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_rsqrt attributes {} { func.func public @test_cbrt(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.cbrt + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.cbrt %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/ceil_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/ceil_op.mlir index fb6a7ef371..4d168e3559 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/ceil_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/ceil_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_ceil attributes {} { func.func public @test_ceil(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.ceil + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.ceil %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/cosine_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/cosine_op.mlir index 8b90d1a7da..cacdcba587 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/cosine_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/cosine_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_cosine attributes {} { func.func public @test_cosine(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.cos + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.cosine %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/exponential_minus_one_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/exponential_minus_one_op.mlir index 5ff8e68658..c22a0ab6d1 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/exponential_minus_one_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/exponential_minus_one_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_expm1 attributes {} { func.func public @test_expm1(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.expm1 + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.exponential_minus_one %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/exponential_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/exponential_op.mlir index cfa263c712..eb00305ad9 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/exponential_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/exponential_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_exp attributes {} { func.func public @test_exp(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.exp + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.exponential %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/floor_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/floor_op.mlir index d0614341fb..8644322ebc 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/floor_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/floor_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_floor attributes {} { func.func public @test_floor(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.floor + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.floor %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/isfinite_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/isfinite_op.mlir index 0de6e7ee9d..5094c6b4c4 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/isfinite_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/isfinite_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_isfinite attributes {} { func.func public @test_isfinite(%arg0: tensor<64x128xf32>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.isfinite + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.is_finite %arg0 : (tensor<64x128xf32>) -> tensor<64x128xi1> return %0 : tensor<64x128xi1> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/log_plus_one_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/log_plus_one_op.mlir index a77ed68e81..f1e21cc0d9 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/log_plus_one_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/log_plus_one_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_log_plus_one attributes {} { func.func public @test_log_plus_one(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.log1p + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.log_plus_one %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/logical_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/logical_op.mlir index 575f9ea34c..0a8d76d369 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/logical_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/logical_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_compare attributes {} { func.func public @logical_not(%arg0: tensor<64x128xi1>) -> tensor<64x128xi1> { + // CHECK: ttnn.empty + // CHECK: ttnn.logical_not + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: tensor<64x128xbf16, + // CHECK-SAME: -> tensor<64x128xbf16, %0 = stablehlo.not %arg0 : tensor<64x128xi1> return %0 : tensor<64x128xi1> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/negate_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/negate_op.mlir index 68ece06ea7..452d867d49 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/negate_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/negate_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_neg attributes {} { func.func public @test_neg(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.neg + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.negate %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/rsqrt_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/rsqrt_op.mlir index f9fdc321ac..3fded1bac8 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/rsqrt_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/rsqrt_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_rsqrt attributes {} { func.func public @test_rsqrt(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.rsqrt + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.rsqrt %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/sign_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/sign_op.mlir index 3f9ea54ec7..80a735e473 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/sign_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/sign_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_sign attributes {} { func.func public @test_sign(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.sign + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.sign %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/sine_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/sine_op.mlir index 19bcec2d46..e1c27a8ff7 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/sine_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/sine_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_sine attributes {} { func.func public @test_sine(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.sin + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.sine %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/sqrt_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/sqrt_op.mlir index 798280d654..1a59a575ef 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/sqrt_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/sqrt_op.mlir @@ -1,11 +1,18 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_sqrt attributes {} { func.func public @test_sqrt(%arg0: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.sqrt + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<64x128xf32, %0 = stablehlo.sqrt %arg0 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/Unary/tranpose_op.mlir b/test/ttmlir/Silicon/StableHLO/Unary/tranpose_op.mlir index 2314dae90d..57679fa7ea 100644 --- a/test/ttmlir/Silicon/StableHLO/Unary/tranpose_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/Unary/tranpose_op.mlir @@ -1,11 +1,17 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_transpose attributes {} { func.func public @test_transpose(%arg0: tensor<64x128xf32>) -> tensor<128x64xf32> { + // CHECK: ttnn.transpose + // CHECK-SAME: {dim0 = 1 : si32, dim1 = 0 : si32} + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: -> tensor<128x64xf32, %0 = stablehlo.transpose %arg0, dims = [1,0] : (tensor<64x128xf32>) -> tensor<128x64xf32> return %0 : tensor<128x64xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/broadcast_op.mlir b/test/ttmlir/Silicon/StableHLO/broadcast_op.mlir index 5236949efe..c6f7f2e972 100644 --- a/test/ttmlir/Silicon/StableHLO/broadcast_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/broadcast_op.mlir @@ -1,11 +1,16 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_broadcast attributes {mhlo.num_partitions = 1 : i32, mhlo.num_replicas = 1 : i32} { func.func public @main(%arg0: tensor<1xf32> {mhlo.layout_mode = "default"}, %arg1: tensor<64x128xf32> {mhlo.layout_mode = "default"}) -> (tensor<64x128xf32> {jax.result_info = "", mhlo.layout_mode = "default"}) { + // CHECK-LABEL: module @jit_broadcast + // CHECK-NOT: broadcast + // CHECK: ttnn.add %0 = stablehlo.broadcast_in_dim %arg0, dims = [1] : (tensor<1xf32>) -> tensor<64x128xf32> %1 = stablehlo.add %0, %arg1 : tensor<64x128xf32> return %1 : tensor<64x128xf32> diff --git a/test/ttmlir/Silicon/StableHLO/composite_op.mlir b/test/ttmlir/Silicon/StableHLO/composite_op.mlir index a44ce1e65e..9bb62c9e8c 100644 --- a/test/ttmlir/Silicon/StableHLO/composite_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/composite_op.mlir @@ -1,16 +1,22 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_add attributes {} { + // CHECK-NOT: func.func.private @add_impl func.func private @add_impl(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { %0 = stablehlo.add %arg0, %arg1 : tensor<64x128xf32> return %0 : tensor<64x128xf32> } + // CHECK-LABEL: func.func public @main func.func public @main(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xf32> { + // CHECK: ttnn.empty + // CEHCK: ttnn.add %results = stablehlo.composite "jit_eltwise_add.my_add" %arg0, %arg1 { decomposition = @add_impl } : (tensor<64x128xf32>, tensor<64x128xf32>) -> tensor<64x128xf32> diff --git a/test/ttmlir/Silicon/StableHLO/conv2d_op.mlir b/test/ttmlir/Silicon/StableHLO/conv2d_op.mlir index a6c0641e09..08b8a07991 100644 --- a/test/ttmlir/Silicon/StableHLO/conv2d_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/conv2d_op.mlir @@ -1,8 +1,10 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s // UNSUPPORTED: true module @jit_convolution attributes {} { diff --git a/test/ttmlir/Silicon/StableHLO/convert_op.mlir b/test/ttmlir/Silicon/StableHLO/convert_op.mlir index 034f783d56..c0825dc0da 100644 --- a/test/ttmlir/Silicon/StableHLO/convert_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/convert_op.mlir @@ -1,11 +1,17 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_convert attributes {} { func.func public @test_convert(%arg0: tensor<64x128xf32>) -> tensor<64x128xbf16> { + // CHECK: ttnn.empty + // CHECK: ttnn.typecast + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, %0 = stablehlo.convert %arg0 : (tensor<64x128xf32>) -> tensor<64x128xbf16> return %0 : tensor<64x128xbf16> } @@ -13,7 +19,15 @@ module @jit_eltwise_convert attributes {} { module @jit_eltwise_add attributes {} { func.func public @test_add(%arg0: tensor<64x128xf32>, %arg1: tensor<64x128xf32>) -> tensor<64x128xbf16> { + // CHECK: ttnn.empty + // CHECK: ttnn.typecast + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, %0 = stablehlo.convert %arg0 : (tensor<64x128xf32>) -> tensor<64x128xbf16> + // CHECK: ttnn.empty + // CHECK: ttnn.typecast + // CHECK-SAME: tensor<64x128xf32, + // CHECK-SAME: tensor<64x128xbf16, %1 = stablehlo.convert %arg1 : (tensor<64x128xf32>) -> tensor<64x128xbf16> %2 = stablehlo.add %0, %1 : tensor<64x128xbf16> return %2 : tensor<64x128xbf16> diff --git a/test/ttmlir/Silicon/StableHLO/dot_general_op.mlir b/test/ttmlir/Silicon/StableHLO/dot_general_op.mlir index 72b6638f32..453604cd69 100644 --- a/test/ttmlir/Silicon/StableHLO/dot_general_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/dot_general_op.mlir @@ -1,11 +1,19 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_dot_general attributes {} { func.func public @test_dot_general(%arg0 : tensor<16x32xf32>, %arg1 : tensor<32x8xf32>) -> tensor<16x8xf32> { + // CHECK: ttnn.empty + // CHECK: ttnn.matmul + // CHECK-SAME: tensor<16x32xf32, + // CHECK-SAME: tensor<32x8xf32, + // CHECK-SAME: tensor<16x8xf32, + // CHECK-SAME: -> tensor<16x8xf32 %0 = stablehlo.dot_general %arg0, %arg1, contracting_dims = [1] x [0] : (tensor<16x32xf32>, tensor<32x8xf32>) -> tensor<16x8xf32> return %0 : tensor<16x8xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/get_dimension_size_op.mlir b/test/ttmlir/Silicon/StableHLO/get_dimension_size_op.mlir index 91f397fee2..68496aa2ee 100644 --- a/test/ttmlir/Silicon/StableHLO/get_dimension_size_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/get_dimension_size_op.mlir @@ -1,11 +1,16 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_get_dimension_size attributes {} { func.func public @test_get_dimension_size(%arg0: tensor<64x128xf32>) -> tensor { + // CHECK: ttnn.full + // CHECK-SAME: {fillValue = 1.280000e+02 : f32} + // CHECK-SAME: -> tensor<1xi32 %0 = stablehlo.get_dimension_size %arg0, dim = 1 : (tensor<64x128xf32>) -> tensor return %0 : tensor } diff --git a/test/ttmlir/Silicon/StableHLO/maxpool2d_op.mlir b/test/ttmlir/Silicon/StableHLO/maxpool2d_op.mlir index c049d12813..e20bb0375a 100644 --- a/test/ttmlir/Silicon/StableHLO/maxpool2d_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/maxpool2d_op.mlir @@ -1,8 +1,10 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s // UNSUPPORTED: true func.func public @test_maxpool2d(%arg0: tensor<1x128x128x32xbf16>) -> tensor<1x64x64x32xbf16> { diff --git a/test/ttmlir/Silicon/StableHLO/mnist_inference.mlir b/test/ttmlir/Silicon/StableHLO/mnist_inference.mlir index 6958d7f928..849f08bb15 100644 --- a/test/ttmlir/Silicon/StableHLO/mnist_inference.mlir +++ b/test/ttmlir/Silicon/StableHLO/mnist_inference.mlir @@ -1,10 +1,13 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s // TODO: Enable when all ops are supported. // UNSUPPORTED: true + module @jit_predict attributes {mhlo.num_partitions = 1 : i32, mhlo.num_replicas = 1 : i32} { func.func public @main(%arg0: tensor<512x784xf32> {mhlo.layout_mode = "default"}, %arg1: tensor<512xf32> {mhlo.layout_mode = "default"}, %arg2: tensor<10x512xf32> {mhlo.layout_mode = "default"}, %arg3: tensor<10xf32> {mhlo.layout_mode = "default"}, %arg4: tensor<128x784xui8> {mhlo.layout_mode = "default"}) -> (tensor<128x10xf32> {jax.result_info = "", mhlo.layout_mode = "default"}) { %0 = stablehlo.convert %arg0 : tensor<512x784xf32> diff --git a/test/ttmlir/Silicon/StableHLO/reduce_add_op.mlir b/test/ttmlir/Silicon/StableHLO/reduce_add_op.mlir index 1fcad52bd8..d9a3a2cddc 100644 --- a/test/ttmlir/Silicon/StableHLO/reduce_add_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/reduce_add_op.mlir @@ -1,11 +1,20 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=/localdev/mmanzoor/compiler/tt-mlir/ttrt-artifacts/system_desc.ttsys" > %t.mlir +// RUN: ttmlir-translate --ttnn-to-flatbuffer %t.mlir > %t.ttnn +// RUN: FileCheck --input-file=%t.mlir %s // UNSUPPORTED: true +// error: keepdim=False is not supported + module @jit_reduce_add attributes {} { func.func public @test_reduce_add(%arg0: tensor<128x10xf32>, %cst_0: tensor) -> tensor<128xf32> { + // CHECK: ttnn.sum + // CHECK-SAME: dim_arg = [1 : i32], + // CHECK-SAME: keep_dim = false + // CHECK-SAME: tensor<128x10xf32 + // CHECK-SAME: -> tensor<128xf32, %0 = stablehlo.reduce(%arg0 init: %cst_0) applies stablehlo.add across dimensions = [1] : (tensor<128x10xf32>, tensor) -> tensor<128xf32> return %0 : tensor<128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/reduce_maximum_op.mlir b/test/ttmlir/Silicon/StableHLO/reduce_maximum_op.mlir index d620c8c509..24b5577985 100644 --- a/test/ttmlir/Silicon/StableHLO/reduce_maximum_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/reduce_maximum_op.mlir @@ -1,11 +1,20 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s // UNSUPPORTED: true +// error: keepdim=False is not supported + module @jit_reduce_maximum attributes {} { func.func public @test_reduce_maximum(%arg0: tensor<128x10xf32>, %cst_0: tensor) -> tensor<128xf32> { + // CHECK: ttnn.max + // CHECK-SAME: dim_arg = [1 : i32], + // CHECK-SAME: keep_dim = false} + // CHECK-SAME: tensor<128x10xf32, + // CHECK-SAME: -> tensor<128xf32 %0 = stablehlo.reduce(%arg0 init: %cst_0) applies stablehlo.maximum across dimensions = [1] : (tensor<128x10xf32>, tensor) -> tensor<128xf32> return %0 : tensor<128xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/rehsape_op.mlir b/test/ttmlir/Silicon/StableHLO/rehsape_op.mlir index 830ba29b85..1214758212 100644 --- a/test/ttmlir/Silicon/StableHLO/rehsape_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/rehsape_op.mlir @@ -1,11 +1,17 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_module_reshape attributes {mhlo.num_partitions = 1 : i32, mhlo.num_replicas = 1 : i32} { func.func public @main(%arg0: tensor<1x64x64x64xf32> {mhlo.layout_mode = "default", mhlo.sharding = "{replicated}"}) -> (tensor<1x1x4096x64xf32> {jax.result_info = "", mhlo.layout_mode = "default"}) { + // CHECK: ttnn.reshape + // CHECK-SAME: {shape = [1 : i32, 1 : i32, 4096 : i32, 64 : i32]} + // CHECK-SAME: tensor<1x64x64x64xf32 + // CHECK-SAME: -> tensor<1x1x4096x64xf32 %0 = stablehlo.reshape %arg0 : (tensor<1x64x64x64xf32>) -> tensor<1x1x4096x64xf32> return %0 : tensor<1x1x4096x64xf32> } diff --git a/test/ttmlir/Silicon/StableHLO/scalar_add_op.mlir b/test/ttmlir/Silicon/StableHLO/scalar_add_op.mlir index 26f697464e..463990d711 100644 --- a/test/ttmlir/Silicon/StableHLO/scalar_add_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/scalar_add_op.mlir @@ -1,11 +1,19 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_scalar_add attributes {} { func.func public @test_scalar_add(%arg0: tensor, %arg1: tensor) -> tensor { + // CHECK: ttnn.empty + // CHECK: ttnn.add + // CHECK-SAME: tensor<1xf32, + // CHECK-SAME: tensor<1xf32, + // CHECK-SAME: tensor<1xf32, + // CHECK-SAME: -> tensor<1xf32, %0 = stablehlo.add %arg0, %arg1 : tensor return %0 : tensor } diff --git a/test/ttmlir/Silicon/StableHLO/slice_op.mlir b/test/ttmlir/Silicon/StableHLO/slice_op.mlir index b62417dd1a..b775b76789 100644 --- a/test/ttmlir/Silicon/StableHLO/slice_op.mlir +++ b/test/ttmlir/Silicon/StableHLO/slice_op.mlir @@ -1,11 +1,21 @@ // REQUIRES: stablehlo // RUN: rm -rf %t.ttnn +// RUN: rm -rf %t.mlir // RUN: ttmlir-opt --stablehlo-to-ttir-pipeline %s | \ -// RUN: ttmlir-opt --ttir-to-ttnn-backend-pipeline="system-desc-path=%system_desc_path%" | \ -// RUN: ttmlir-translate --ttnn-to-flatbuffer > %t.ttnn +// 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 +// RUN: FileCheck --input-file=%t.mlir %s module @jit_eltwise_subtract attributes {} { func.func @slice_op(%arg0: tensor<32x64xbf16>) -> tensor<8x8xbf16> { + // CHECK: ttnn.empty + // CHECK: ttnn.slice + // CHECK-SAME: begins = [0 : i32, 16 : i32], + // CHECK-SAME: ends = [16 : i32, 32 : i32], + // CHECK-SAME: step = [2 : i32, 2 : i32] + // CHECK-SAME: tensor<32x64xbf16, + // CHECK-SAME: tensor<8x8xbf16, + // CHECK-SAME: -> tensor<8x8xbf16 %result = "stablehlo.slice"(%arg0) { start_indices = array, limit_indices = array,