[benlaflen]

Includes patch from Issue #561 as well as a unified address decoding pipeline shared in tracer_nvbit and trace-parser

[JRPan]

Thanks! The #561 logic is right, but this doesn't build yet (the post-processor and trace-parser both fail; inline comments below). I've opened #564 with just the #561 fix so it can merge now. Let's land the shared decoder here afterwards. Please also: fold in or drop the TMA formats (tma_* decoding is still separate, so the title overclaims), squash the two identically named commits, and add a small fixture test (a base_delta ARRIVE that hits an mbarrier only via a delta lane must keep its TRYWAIT). Once it builds, the post-processed .tracez should be byte-identical to #564's.

[JRPan]

Re-review on d71a20e (I pushed a clang-format-only commit; no logic changes). The earlier compile/include comments look addressed, but two blockers remain:

1. decode_base_stride returns 31 entries when the active run ends before lane 31, so trace_parser.cc's memcpy of 32 entries reads out of bounds.
2. The address_format enum was deleted from tracer_tool.cu but is still used there (lines 1458-1468 and 1568-1576), and decode.h does not define it, so the tracer no longer builds.

Also still open: decode() does not handle the TMA modes (3/4), so the title overclaims, and assert(!ss.fail()) vanishes under NDEBUG. See the inline comments. The earlier request for a base_delta fixture test and a byte-identical .tracez check against #564 still stands.

[JRPan]

Bug: the lane that ends the run sets last_bit1_found but pushes nothing, so for any mask whose run ends before lane 31 the vector has 31 entries. The old code assigned addrs[s] by index. trace_parser.cc then memcpys WARP_SIZE * 8 bytes from it, reading 8 bytes past the end. Size the vector up front (std::vector<uint64_t> addrs(WARP_SIZE, 0)) and assign addrs[s], or push 0 here. Please also assert addrs.size() == WARP_SIZE before returning from decode().

[JRPan]

decode() only handles modes 0-2 and throws for 3/4 (TMA), which are still decoded separately in trace_parser.cc and post-traces-processing.cpp. Either fold the TMA formats in here, or drop the unification claim from the title and state they are out of scope. Also, assert(!ss.fail()) is compiled out under NDEBUG, so a truncated line silently decodes to garbage in release builds; throw on failure instead.

[JRPan]

This PR deletes enum address_format from this file, but it is still used at lines 1458, 1462, 1468 (base_stride, base_delta, list_all) and 1568, 1576 (tma_base_delta, tma_list_all). decode.h does not define it, so the tracer does not compile. Move the full enum (0-4) into decode.h and use it from the tracer, the parser and decode().

[JRPan]

The enum is now split: here only the TMA values remain, while the tracer used to carry all five. Keep a single address_format definition in the shared header so the tracer and parser cannot drift apart.