Record min stack free, not max stack use (#22465)

### Ticket
#19983

### Problem description
Moving to dynamic stack sizing means it is more useful to directly
report the remaining free stack space.

### What's changed
* Record remaining stack free, rather than size used.
* Add stack size test cases
* Implement '*' globing in the debug utils.
* simplify `debug_stack_usage_t`. It was two arrays of uint16_t, now it
is a single array of uint16_t pairs.

There's a little excitement in that we need to record values 0 and up,
but the buffer is also initialized to zero. We do offsets by 1 and
unsigned arithmetic wraparound to avoid some comparisons.

The stack sizes remain fixed, as before.

### Checklist
- [YES] [All post
commit](https://github.com/tenstorrent/tt-metal/actions/workflows/all-post-commit-workflows.yaml)
CI passes
- [YES] [Blackhole Post
commit](https://github.com/tenstorrent/tt-metal/actions/workflows/blackhole-post-commit.yaml)
CI with demo tests 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) [Single-card demo
tests](https://github.com/tenstorrent/tt-metal/actions/workflows/single-card-demo-tests.yaml)
CI passes (if applicable) See [recommended dev
flow](https://github.com/tenstorrent/tt-metal/blob/main/models/MODEL_ADD.md#a-recommended-dev-flow-on-github-for-adding-new-models).
- [ ] New/Existing tests provide coverage for changes

Author: nathan-TT

tests/tt_metal/tt_metal/debug_tools/CMakeLists.txt

@@ -18,6 +18,7 @@ set(UNIT_TESTS_DEBUG_TOOLS_SRC
     ${CMAKE_CURRENT_SOURCE_DIR}/watcher/test_noc_sanitize.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/watcher/test_pause.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/watcher/test_ringbuf.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/watcher/test_stack_usage.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/watcher/test_waypoint.cpp
 )
 

tests/tt_metal/tt_metal/debug_tools/debug_tools_test_utils.hpp

@@ -34,42 +34,90 @@ inline void DumpFile(string file_name) {
     }
 }
 
-// Wildcard is '?', just like a glob.
-inline bool StringCompareWithWildcard(const std::string_view str, const std::string_view pattern) {
-    size_t pattern_size = pattern.size();
-    if (pattern_size != str.size())
-        return false;
-    for (int idx = 0; idx < pattern_size; idx++) {
-        if (str[idx] != pattern[idx] && pattern[idx] != '?') {
-            return false;
+std::string_view::size_type FloatingGlobEndsAt(const std::string_view haystack,
+                                               const std::string_view needle,
+                                               unsigned globs);
+
+// Check of pattern matches at the beginning of str.
+inline std::string_view::size_type AnchoredGlobEndsAt(const std::string_view str,
+                                                      const std::string_view pattern,
+                                                      unsigned globs) {
+    if (str.size() + globs < pattern.size()) {
+        return str.npos;
+    }
+
+    for (std::string_view::size_type idx = 0; idx != pattern.size(); idx++) {
+        if (pattern[idx] == '*') {
+            auto result = FloatingGlobEndsAt(str.substr(idx), pattern.substr(idx + 1), globs - 1);
+            if (result != str.npos) {
+                // An empty suffix matches the whole string.
+                result = result ? result + idx : str.size();
+            }
+            return result;
+        } else if (idx >= str.size()) {
+            return str.npos;
+        } else if (pattern[idx] == '?') {
+            continue;
+        } else if (str[idx] != pattern[idx]) {
+            return str.npos;
         }
     }
-    return true;
+    return pattern.size();;
 }
 
-// Check if haystack contains needle, return true. needle may contain
-// '?' to match any character (just like glob).
-inline bool StringContainsWithWildcard(std::string_view haystack, std::string_view needle) {
-    size_t needle_size = needle.size();
-    if (needle_size == 0 || needle.front() == '?') {
-        // The needle is empty, or begins with '?', fail in order to
-        // force test to be fixed.
-        return false;
-    }
-    if (needle_size > haystack.size()) {
-        return false;
+// Look for needle in haystack. We look backwards through haystack, so
+// that glob use will find the longest match.
+inline std::string_view::size_type FloatingGlobEndsAt(const std::string_view haystack,
+                                                      const std::string_view needle,
+                                                      unsigned globs) {
+    if (needle.empty()) {
+        // Empty needle matches at end.
+        return haystack.size();
     }
     char first = needle.front();
+    if (first == '*') {
+        // '*' at front, handle as an anchored glob.
+        return AnchoredGlobEndsAt(haystack, needle, globs);
+    }
+    if (haystack.size() + globs < needle.size()) {
+        return haystack.npos;
+    }
 
-    for (size_t idx = 0, limit = haystack.size() - needle_size;
-         (idx = haystack.find(first, idx)) <= limit; idx++) {
-        std::string_view substr(&haystack[idx], needle_size);
-        if (StringCompareWithWildcard(substr, needle)) {
-            return true;
+    for (std::string_view::size_type idx = haystack.size() + globs - needle.size();; idx--) {
+        if (first != '?') {
+            // no wildcard at front, scan for first char to begin search.
+            idx = haystack.rfind(first, idx);
+            if (idx == haystack.npos) {
+                break;
+            }
         }
+        // Try an anchored match here.
+        auto result = AnchoredGlobEndsAt(haystack.substr(idx), needle, globs);
+        if (result != haystack.npos) {
+            return result + idx;
+        }
+        if (!idx)
+            break;
     }
 
-    return false;
+    return haystack.npos;
+}
+
+// Count the number of '*' characters.
+inline unsigned GlobCount(const std::string_view glob) {
+    unsigned count = 0;
+    for (std::string_view::size_type idx = 0; (idx = glob.find('*', idx)) != glob.npos; idx++)
+        count++;
+    return count;
+}
+// str matches pattern, allowing '?' and '*' globbing.
+inline bool StringMatchesGlob(const std::string_view str, const std::string_view pattern) {
+    return AnchoredGlobEndsAt(str, pattern, GlobCount(pattern)) == str.size();
+}
+
+// haystack contains needle, allowing '?' and '*' globbing.
+inline bool StringContainsGlob(const std::string_view haystack, const std::string_view needle) {
+    return FloatingGlobEndsAt(haystack, needle, GlobCount(needle)) != haystack.npos;
 }
 
 // Check whether the given file contains a list of strings in any order. Doesn't check for
@@ -99,7 +147,7 @@ inline bool FileContainsAllStrings(string file_name, const std::vector<string> &
         // Check for all target strings in the current line
         std::vector<std::string_view> found_on_current_line;
         for (const auto &s : must_contain_set) {
-            if (StringContainsWithWildcard(line, s)) {
+            if (StringContainsGlob(line, s)) {
                 found_on_current_line.push_back(s);
             }
         }
@@ -149,7 +197,7 @@ inline bool FileContainsAllStringsInOrder(string file_name, const std::vector<st
 
         // Check for all target strings in the current line
         for (; !must_contain_queue.empty(); must_contain_queue.pop_front()) {
-            if (!StringContainsWithWildcard(line, must_contain_queue.front())) {
+            if (!StringContainsGlob(line, must_contain_queue.front())) {
                 break;
             }
         }
@@ -190,7 +238,7 @@ inline bool FilesMatchesString(string file_name, const string& expected) {
     int line_num = 0;
     while (getline(file, line_a) && getline(expect_stream, line_b)) {
         line_num++;
-        if (!StringCompareWithWildcard(line_a, line_b)) {
+        if (!StringMatchesGlob(line_a, line_b)) {
             tt::log_info(
                 tt::LogTest,
                 "Test Error: Line {} of {} did not match expected:\n\t{}\n\t{}",

tests/tt_metal/tt_metal/debug_tools/watcher/test_stack_usage.cpp

@@ -0,0 +1,93 @@
+// SPDX-FileCopyrightText: © 2024 Tenstorrent Inc.
+//
+// SPDX-License-Identifier: Apache-2.0
+
+#include "command_queue_fixture.hpp"
+
+#include <tt-metalium/allocator.hpp>
+#include <tt-metalium/core_coord.hpp>
+#include <tt-metalium/host_api.hpp>
+#include <tt-metalium/kernel.hpp>
+#include <tt-metalium/kernel_types.hpp>
+#include <gtest/gtest.h>
+#include "debug_tools_fixture.hpp"
+#include "debug_tools_test_utils.hpp"
+#include <watcher_server.hpp>
+#include <fmt/base.h>
+#include <string>
+#include <vector>
+
+using namespace tt;
+using namespace tt::tt_metal;
+
+namespace {
+void RunOneTest(WatcherFixture* fixture, IDevice* device, unsigned free) {
+    static const char *const names[] =
+        {"brisc", "ncrisc", "trisc0", "trisc1", "trisc2", "aerisc", "ierisc"};
+    const std::string path = "tests/tt_metal/tt_metal/test_kernels/misc/watcher_stack.cpp";
+    auto msg = [&](std::vector<std::string> &msgs, const char *cpu, unsigned free) {
+        if (msgs.empty()) {
+            msgs.push_back("Stack usage summary:");
+        }
+        msgs.push_back(fmt::format("{} highest stack usage: {} bytes free, on core "
+                                   "* running kernel {} ({})",
+                                   cpu, free, path, !free ? "OVERFLOW" : "Close to overflow"));
+    };
+
+    // Set up program
+    Program program = Program();
+    CoreCoord coord = {0, 0};
+    std::vector<uint32_t> compile_args{free};
+    std::vector<string> expected;
+
+    CreateKernel(program, path, coord,
+                 DataMovementConfig{.processor = DataMovementProcessor::RISCV_0,
+                     .noc = NOC::RISCV_0_default,
+                     .compile_args = compile_args});
+    msg(expected, names[0], free);
+
+    CreateKernel(program, path, coord,
+                 DataMovementConfig{.processor = DataMovementProcessor::RISCV_1,
+                     .noc = NOC::RISCV_1_default,
+                     .compile_args = compile_args});
+    msg(expected, names[1], free);
+
+    CreateKernel(program, path, coord, ComputeConfig{.compile_args = compile_args});
+    for (unsigned ix = 0; ix != 2; ix++) {
+        msg(expected, names[2 + ix], free);
+    }
+
+    // Also run on idle ethernet, if present
+    auto const &inactive_eth_cores = device->get_inactive_ethernet_cores();
+    if (!inactive_eth_cores.empty() && fixture->IsSlowDispatch()) {
+        // Just pick the first core
+        CoreCoord idle_coord = CoreCoord(*inactive_eth_cores.begin());
+        CreateKernel(program, path, idle_coord,
+                     tt_metal::EthernetConfig{.eth_mode = Eth::IDLE, .noc = tt_metal::NOC::NOC_0,
+                         .processor = DataMovementProcessor::RISCV_0, .compile_args = compile_args});
+        msg(expected, names[6], free);
+    }
+
+    fixture->RunProgram(device, program, true);
+
+    EXPECT_TRUE(FileContainsAllStringsInOrder(fixture->log_file_name, expected));
+}
+
+template<uint32_t Free>
+void RunTest(WatcherFixture* fixture, IDevice* device) {
+    RunOneTest(fixture, device, Free);
+}
+
+} // namespace
+
+TEST_F(WatcherFixture, TestWatcherStackUsage0) {
+    for (IDevice* device : this->devices_) {
+        this->RunTestOnDevice(RunTest<0>, device);
+    }
+}
+
+TEST_F(WatcherFixture, TestWatcherStackUsage16) {
+    for (IDevice* device : this->devices_) {
+        this->RunTestOnDevice(RunTest<16>, device);
+    }
+}

tests/tt_metal/tt_metal/test_kernels/misc/watcher_stack.cpp

@@ -0,0 +1,57 @@
+// SPDX-FileCopyrightText: © 2025 Tenstorrent Inc.
+//
+// SPDX-License-Identifier: Apache-2.0
+
+// Scribble on the stack to check stack usage detection.
+
+#include "compile_time_args.h"
+#include <dev_mem_map.h>
+
+static uint32_t get_stack_base() {
+#if defined(COMPILE_FOR_BRISC)
+    return MEM_BRISC_STACK_TOP - MEM_BRISC_STACK_SIZE;
+#elif defined(COMPILE_FOR_NCRISC)
+    return MEM_NCRISC_STACK_TOP - MEM_NCRISC_STACK_SIZE;
+#elif defined(COMPILE_FOR_IDLE_ERISC)
+#if COMPILE_FOR_IDLE_ERISC == 0
+    return MEM_IERISC_STACK_TOP - MEM_IERISC_STACK_SIZE;
+#elif COMPILE_FOR_IDLE_ERISC == 1
+    return MEM_SLAVE_IERISC_STACK_TOP - MEM_SLAVE_IERISC_STACK_SIZE;
+#else
+#error "idle erisc get_stack_base unknown"
+#endif
+#elif defined(COMPILE_FOR_TRISC)
+#if COMPILE_FOR_TRISC == 0
+    return MEM_TRISC0_STACK_TOP - MEM_TRISC0_STACK_SIZE;
+#elif COMPILE_FOR_TRISC == 1
+    return MEM_TRISC1_STACK_TOP - MEM_TRISC1_STACK_SIZE;
+#elif COMPILE_FOR_TRISC == 2
+    return MEM_TRISC2_STACK_TOP - MEM_TRISC2_STACK_SIZE;
+#else
+#error "trisc get_stack_base unknown"
+#endif
+#else
+#error "get_stack_base unknown"
+#endif
+}
+
+#if defined(COMPILE_FOR_TRISC)
+#include "compute_kernel_api/common.h"
+namespace NAMESPACE {
+void MAIN {
+#else
+void kernel_main() {
+#endif
+    uint32_t usage = get_compile_time_arg_val (0);
+    auto base = (uint32_t tt_l1_ptr *)get_stack_base();
+    auto point = &base[usage/sizeof(uint32_t)];
+    uint32_t *sp;
+    asm ("mv %0,sp" : "=r"(sp));
+
+    // Do not scribble above stack pointer.
+    if (sp > point)
+        *point = 0;
+}
+#if defined(COMPILE_FOR_TRISC)
+} // namespace NAMESPACE
+#endif