fix sliding window hash calculus (#18053)

### Problem description
Sliding Window Infra hash calculus didn't account for is_bilinear,
is_transpose and snap_to_tile. This was causing a customer model to
fail.

### What's changed
Updated SlidingWindowConfig::to_string(), that is used by
SlidingWindowConfig::to_hash()

### Checklist
- [ ] [All post
commit](https://github.com/tenstorrent/tt-metal/actions/workflows/all-post-commit-workflows.yaml)
CI passes
- [ ] [Blackhole Post
commit](https://github.com/tenstorrent/tt-metal/actions/workflows/blackhole-post-commit.yaml)
CI passes (if applicable)
- [ ] [Model
regression](https://github.com/tenstorrent/tt-metal/actions/workflows/perf-models.yaml)
CI passes (if applicable)
- [ ] [Device performance
regression](https://github.com/tenstorrent/tt-metal/actions/workflows/perf-device-models.yaml)
CI passes (if applicable)
- [ ] **(For models and ops writers)** Full [new models
tests](https://github.com/tenstorrent/tt-metal/actions/workflows/full-new-models-suite.yaml)
CI passes (if applicable)
- [ ] New/Existing tests provide coverage for changes

tests/ttnn/unit_tests/gtests/CMakeLists.txt

@@ -9,6 +9,7 @@ set(TTNN_UNIT_TESTS_SRC
     ${CMAKE_CURRENT_SOURCE_DIR}/test_graph_query_op_runtime.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_reflect.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_to_and_from_json.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/test_sliding_window_infra.cpp
 )
 
 set(TTNN_CCL_UNIT_TESTS_SRC

tests/ttnn/unit_tests/gtests/test_sliding_window_infra.cpp

@@ -0,0 +1,73 @@
+// SPDX-FileCopyrightText: © 2025 Tenstorrent Inc.
+//
+// SPDX-License-Identifier: Apache-2.0
+
+#include "gtest/gtest.h"
+
+#include "ttnn/operations/sliding_window/sliding_window.hpp"
+#include "tt_metal/api/tt-metalium/core_coord.hpp"
+
+namespace ttnn::operations::sliding_window::test {
+
+using namespace tt::tt_metal;
+
+class SlidingWindowTestFixture : public testing::TestWithParam<SlidingWindowConfig> {};
+
+TEST_P(SlidingWindowTestFixture, SlidingWindowHash) {
+    auto sliding_window_a = GetParam();
+
+    // start of same input
+    auto sliding_window_b = sliding_window_a;
+    log_info(tt::LogTest, "sliding_window_a:[{}] {}", sliding_window_a.get_hash(), sliding_window_a.to_string());
+    log_info(tt::LogTest, "sliding_window_b:[{}] {}", sliding_window_b.get_hash(), sliding_window_b.to_string());
+    EXPECT_EQ(sliding_window_a.get_hash(), sliding_window_b.get_hash());
+
+    // flip snap_to_tile
+    sliding_window_b.snap_to_tile = !sliding_window_a.snap_to_tile;
+    log_info(tt::LogTest, "sliding_window_a:[{}] {}", sliding_window_a.get_hash(), sliding_window_a.to_string());
+    log_info(tt::LogTest, "sliding_window_b:[{}] {}", sliding_window_b.get_hash(), sliding_window_b.to_string());
+    EXPECT_NE(sliding_window_a.get_hash(), sliding_window_b.get_hash());
+    sliding_window_b.snap_to_tile = !sliding_window_a.snap_to_tile;
+
+    // flip is_bilinear
+    sliding_window_b.is_bilinear = !sliding_window_a.is_bilinear;
+    log_info(tt::LogTest, "sliding_window_a:[{}] {}", sliding_window_a.get_hash(), sliding_window_a.to_string());
+    log_info(tt::LogTest, "sliding_window_b:[{}] {}", sliding_window_b.get_hash(), sliding_window_b.to_string());
+    EXPECT_NE(sliding_window_a.get_hash(), sliding_window_b.get_hash());
+    sliding_window_b.is_bilinear = !sliding_window_a.is_bilinear;
+
+    // flip is_transpose
+    sliding_window_b.is_transpose = !sliding_window_a.is_transpose;
+    log_info(tt::LogTest, "sliding_window_a:[{}] {}", sliding_window_a.get_hash(), sliding_window_a.to_string());
+    log_info(tt::LogTest, "sliding_window_b:[{}] {}", sliding_window_b.get_hash(), sliding_window_b.to_string());
+    EXPECT_NE(sliding_window_a.get_hash(), sliding_window_b.get_hash());
+    sliding_window_b.is_transpose = !sliding_window_a.is_transpose;
+
+    // flip ceil_mode
+    sliding_window_b.ceil_mode = !sliding_window_a.ceil_mode;
+    log_info(tt::LogTest, "sliding_window_a:[{}] {}", sliding_window_a.get_hash(), sliding_window_a.to_string());
+    log_info(tt::LogTest, "sliding_window_b:[{}] {}", sliding_window_b.get_hash(), sliding_window_b.to_string());
+    EXPECT_NE(sliding_window_a.get_hash(), sliding_window_b.get_hash());
+    sliding_window_b.ceil_mode = !sliding_window_a.ceil_mode;
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    SlidingWindowHashTests,
+    SlidingWindowTestFixture,
+    ::testing::Values(SlidingWindowConfig{
+        .batch_size = 1,
+        .input_hw = {32, 32},
+        .window_hw = {3, 3},
+        .stride_hw = {1, 1},
+        .pad_hw = {1, 1},
+        .output_pad_hw = {0, 0},
+        .dilation_hw = {1, 1},
+        .num_cores_nhw = 1,
+        .num_cores_c = 1,
+        .core_range_set = tt::tt_metal::CoreRangeSet(tt::tt_metal::CoreRange({0, 0}, {7, 7})),
+        .snap_to_tile = false,
+        .is_bilinear = false,
+        .is_transpose = false,
+        .ceil_mode = false}));
+
+}  // namespace ttnn::operations::sliding_window::test

ttnn/cpp/ttnn/operations/sliding_window/sliding_window.cpp

@@ -699,7 +699,8 @@ std::string SlidingWindowConfig::to_string() const {
            std::to_string(std::get<1>(stride_hw)) + "_" + std::to_string(std::get<0>(pad_hw)) + "_" +
            std::to_string(std::get<1>(pad_hw)) + "_" + std::to_string(std::get<0>(dilation_hw)) + "_" +
            std::to_string(std::get<1>(dilation_hw)) + "_" + std::to_string(num_cores_nhw) + "_" +
-           std::to_string(num_cores_c) + "_" + std::to_string(ceil_mode) + "_" + core_range_set.str();
+           std::to_string(num_cores_c) + "_" + core_range_set.str() + (snap_to_tile ? "_snap_to_tile" : "") +
+           (is_bilinear ? "_bilinear" : "") + (is_transpose ? "_transpose" : "") + (ceil_mode ? "_ceil_mode" : "");
 }
 
 }  // namespace ttnn::operations::sliding_window