diff --git a/projects/rocprof-trace-decoder/README.md b/projects/rocprof-trace-decoder/README.md index 62c96080747a..ede13a60f87c 100644 --- a/projects/rocprof-trace-decoder/README.md +++ b/projects/rocprof-trace-decoder/README.md @@ -63,7 +63,7 @@ The chosen `LLVM_DIR` is baked into the generated `rocprof-trace-decoder-config. |---|---| | `rocprof-trace-decoder` | Shared library (`.so`) | | `rocprof-trace-decoder-static` | Static library (`.a`) — built but not installed; consumed via the build-tree CMake export | -| `SQTTInstrumentPass` | LLVM pass plugin for SQTT marker instrumentation (`BUILD_MARKERS=ON`) | +| `SQTTInstrumentPass` | LLVM pass plugin for SQTT marker instrumentation (`libsqttinstrumentpass.so`, `BUILD_MARKERS=ON`) | | `unit_tests` | Unit test executable (requires `BUILD_TESTS=ON`) | | `format` | Run clang-format and cmake-format on all sources | | `docs` | Generate Doxygen API documentation | diff --git a/projects/rocprof-trace-decoder/docs/markers/SPEC.md b/projects/rocprof-trace-decoder/docs/markers/SPEC.md index 3f8adee242a6..5612f82d622f 100644 --- a/projects/rocprof-trace-decoder/docs/markers/SPEC.md +++ b/projects/rocprof-trace-decoder/docs/markers/SPEC.md @@ -10,7 +10,7 @@ function entry/exit, barrier synchronization, and user-defined annotations. The system consists of: -- **`SQTTInstrumentPass.so`** -- LLVM pass plugin loaded via `-fpass-plugin=` +- **`libsqttinstrumentpass.so`** -- LLVM pass plugin loaded via `-fpass-plugin=` - **`markers.hpp`** -- Device-side header for user markers - **`sqtt_flamegraph.py`** -- Post-processing tool: reads SQTT traces and `.sqtt_funcmap` sections to generate flamegraphs @@ -33,13 +33,13 @@ it gets symbols from the host compiler process. ```bash # Minimal (barriers + user markers only) SQTT_INSTRUMENT_BARRIERS=1 \ -hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/SQTTInstrumentPass.so \ +hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel # Full (function entry/exit with threshold, barriers, scope filtering) SQTT_INSTRUMENT_FUNCTIONS=10 SQTT_INSTRUMENT_BARRIERS=1 \ SQTT_SCOPE_CU=0x3 \ -hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/SQTTInstrumentPass.so \ +hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -85,7 +85,8 @@ Where `flags` is the OR of applicable bits: `EXIT_PREV=1`, `ENTER=2`. - If `encoded <= 0xFF`, the target supports it, and shader-clock packing is inactive: emit `s_ttracedata_imm` (8-bit immediate, no m0 setup, ~1 cycle faster) -- Otherwise: emit `s_ttracedata` (32-bit value loaded into m0) +- Otherwise: emit `s_ttracedata` through M0. The pass explicitly emits + `s_mov_b32 m0, ` + `s_nop 0` + `s_ttracedata` on every target. With the 2-bit encoding, IDs 1-63 fit in `s_ttracedata_imm` (previously limited to IDs 1-15 with 4-bit flags). @@ -95,9 +96,10 @@ SQTT hardware only captures the low 8 bits from `s_ttracedata_imm`, so the ### GFX12 Shader Clock Packing -On gfx12+, marker headers include a truncated shader clock from -`s_getreg_b32 hwreg(HW_REG_SHADER_CYCLES_LO, shift, bits)`. The default window -is source bits 4-15 (`SQTT_SHADER_CLOCK_SHIFT=4`, `SQTT_SHADER_CLOCK_BITS=12`), +For the supported gfx12 targets, marker headers include a truncated shader clock. +gfx1250 uses `s_get_shader_cycles_u64`; gfx1200 and gfx1201 use +`s_getreg_b32 hwreg(HW_REG_SHADER_CYCLES_LO, shift, bits)`. When enabled, the +default window is source bits 4-15 (`SQTT_SHADER_CLOCK_SHIFT=4`, `SQTT_SHADER_CLOCK_BITS=12`), stored in marker bits 31-20, leaving 18 marker ID bits plus the two low flags: ```python @@ -108,10 +110,53 @@ id = (val >> 2) & ((1 << (30 - clock_bits)) - 1) clock = val >> (32 - clock_bits) ``` -The clock field is the shader clock sampled with `s_getreg`, not the realtime +Only headers created by the pass use this packed form. Numeric point and enter +markers retain their legacy values so arbitrary user shaderdata is not +reinterpreted as a packed header. Bare exits use the packed form too: their +value is intrinsically indistinguishable from a pass-generated exit. +`att_tool.py` restores packed headers to the legacy form by default; +`--no-decode-markers` preserves the raw value. + +The clock field is the target-specific shader-clock sample, not the realtime clock associated with `s_sendmsg`. The funcmap records the active layout with `M:shader_clock_bits=N;shader_clock_shift=S`. When `SQTT_SHADER_CLOCK_BITS=0`, -gfx12 uses the legacy marker layout and may use `s_ttracedata_imm`. +gfx12 uses the no-clock full-ID layout and may use `s_ttracedata_imm`. +A nonzero `shader_clock_bits` value requires `shader_clock_shift`; a malformed +packed row uses the no-clock full-ID layout. + +Any marker with `R:extra_payload_count > 0`, including address traces and +named `sqtt_marker_data()` records, requires `SQTT_SHADER_CLOCK_BITS=0`. +Clock packing is disabled by default; an explicit nonzero setting is an error. + +The sampled window lets a decoder compensate for the variable delay between +the trace instruction issuing and its shaderdata record appearing in SQTT. +Shaderdata and shader-clock timestamps have the same rate but an unknown fixed +phase per physical SIMD. Do not compare them directly. For each **marker +header** (never an `R:`-declared raw payload), choose any header `r` in the +same physical SIMD/layout only as a coordinate origin and calculate: + +```text +window_mask = (1ULL << (clock_bits + shader_clock_shift)) - 1 +half_window = (window_mask + 1) / 2 +bucket_size = 1ULL << shader_clock_shift +sampled_window = packed_clock_field << shader_clock_shift +relative = (((record_time - r.record_time) - + (sampled_window - r.sampled_window) + half_window) & window_mask) - half_window +``` + +Within each `(shader engine, CU, SIMD, clock_bits, shader_clock_shift)` domain, +the implementation assumes the relative delta stays within the signed sampled +window. Collect the complete domain before selecting `minimum_relative`, then +shift every header earlier by +`max(0, relative - minimum_relative - (bucket_size - 1))`. The coordinate +origin is not a chronological anchor: a later lower value corrects earlier +headers retroactively. + +`att_tool.py` applies this correction while writing JSON by default; +`--no-decode-markers` preserves the packed values and arrival timestamps. It +rewrites each corrected header to the legacy `(id << 2) | flags` form. For an +externally produced packed payload block, JSON output may normalize a +recognized header but does not apply timestamp correction. ### Semantic rules @@ -146,6 +191,44 @@ id = val >> 2 # or masked to 30 - shader_clock_bits when M: is present The marker type (function, user scope, point event) is resolved from the funcmap, not from encoding bits. +### C++ decoder marker encoding + +`` exposes the parsed funcmap and +its `MarkerEncoding`: + +```cpp +const auto& encoding = codeobjs.getFuncmap(codeobj_id).marker_encoding; +const auto marker = codeobjs.decodeMarkerValue(codeobj_id, raw_shaderdata_value); + +uint32_t id = marker.id; +uint32_t marker_id_mask = encoding.marker_id_mask(); +uint32_t packed_clock_mask = encoding.packed_shader_clock_mask(); +uint32_t source_clock_mask = encoding.shader_clock_source_mask(); + +// The sampled field is right-aligned. +uint32_t sampled_lo = 0; +if (encoding.has_shader_clock() && encoding.is_valid()) + sampled_lo = encoding.decode_shader_clock(raw_shaderdata_value) << encoding.shader_clock_shift; +``` + +`marker_id_mask()` and `packed_shader_clock_mask()` apply to the raw 32-bit +shaderdata value. `shader_clock_source_mask()` describes the matching source +low-word window for timestamp comparison (`HW_REG_SHADER_CYCLES_LO` on gfx1200/1201). +`decode_shader_clock()` is meaningful only after the record has been identified +as a packed header. + +`decode_marker_value(raw, funcmap)` (and the `CodeobjMap::decodeMarkerValue()` +wrapper) uses the packed layout for IDs known to the funcmap and for bare +exits; unregistered numeric point/enter values otherwise remain legacy. A +numeric value whose low packed ID collides with a registered ID is inherently +ambiguous from the 32-bit record alone. Call +`decode_marker_value(raw, encoding)` directly when the record is known to be a +packed header. + +With no `M:` row, `MarkerEncoding` uses the no-clock full-ID layout: +`marker_id_mask()` is `0xFFFFFFFC` and the two shader-clock masks and decoded +sampled field are zero. + --- ## ID Allocation @@ -254,8 +337,8 @@ All variables are read at **compile time** by the pass plugin. | `SQTT_INSTRUMENT_FUNCTIONS`| `0` | Function entry/exit threshold (0=disabled) | | `SQTT_INSTRUMENT_MEMORY` | off | Memory op markers. Format: `N:M` (N=ops per marker, M=max gap) | | `SQTT_MEM_BARRIER` | `fence` | Reordering boundary around markers (`none`/`asm`/`fence` or `0`/`1`/`2`) | -| `SQTT_SHADER_CLOCK_BITS` | auto | Gfx12 shader clock bits in marker headers (`12` on gfx12, `0` elsewhere; `0` disables) | -| `SQTT_SHADER_CLOCK_SHIFT` | `4` | Source bit offset in `HW_REG_SHADER_CYCLES_LO` for shader clock packing | +| `SQTT_SHADER_CLOCK_BITS` | `0` | Gfx12 shader clock bits. Set a nonzero value to enable packing; it is invalid with any payload-bearing marker. | +| `SQTT_SHADER_CLOCK_SHIFT` | `4` | Source bit offset in the shader-clock low word for packing | ### Marker reorder boundary (`SQTT_MEM_BARRIER`) @@ -493,6 +576,22 @@ M:shader_clock_bits=12;shader_clock_shift=4 -- gfx12 marker header clock pack The section is added to `llvm.used` to prevent linker stripping. +### `code.json` marker metadata + +`att_tool.py` adds embedded `sqtt_funcmap` data to `code.json` using the +established six-column row shape. Two additive top-level tables carry metadata +that cannot be represented in those rows: + +```json +"sqtt_funcmap_layout": [[codeobj_id, shader_clock_bits, shader_clock_shift]], +"sqtt_funcmap_payloads": [[codeobj_id, marker_id, extra_payload_count]] +``` + +`sqtt_funcmap_layout` contains rows only for code objects with shader-clock +packing; `sqtt_funcmap_payloads` contains nonzero payload counts. Readers that +do not know these optional tables can continue using the unchanged +`sqtt_funcmap` rows. For a packed row, an absent payload count means zero. + ### Extracting the funcmap **Method 1: `sqtt_decode_funcmap.py` (recommended)** @@ -715,7 +814,7 @@ as cache line utilization, stride detection, and memory coalescing analysis. | Variable | Format | Default | Description | |---|---|---|---| -| `SQTT_TRACE_ADDRESSES` | `memory`, `lds`, or `memory,lds` | off | Trace per-lane addresses. `memory` = global/buffer/flat/atomics. `lds` = LDS (AS=3), LDS atomics, ds_permute/ds_bpermute. | +| `SQTT_TRACE_ADDRESSES` | `memory`, `lds`, or `memory,lds` | off | Trace per-lane addresses. `memory` = flat/global (AS=0/1), buffer, and matching atomics. `lds` = LDS (AS=3), LDS atomics, ds_permute/ds_bpermute. | `SQTT_TRACE_ADDRESSES` and `SQTT_INSTRUMENT_MEMORY` are mutually exclusive — setting both produces an error. Address tracing already emits a header marker @@ -733,7 +832,9 @@ for every memory op, so `SQTT_INSTRUMENT_MEMORY` is redundant. | `struct_buffer` | buffer | 32-bit | 2 (voffset+vindex) | 3 (rsrc_lo, rsrc_hi, soffset) | 67 | | `ds_permute`/`ds_bpermute` | LDS hw | 32-bit | 1 (index) | 0 | 32 | -Private/scratch (AS=5) is always excluded. All categories also emit exec_lo + exec_hi (2 tokens) and a header marker (1 token) before the data tokens shown above. +Pointer address spaces other than 0, 1, and selected LDS AS=3 (including +private/scratch AS=5) are excluded. All categories also emit exec_lo + exec_hi +(2 tokens) and a header marker (1 token) before the data tokens shown above. ### Wave size @@ -823,9 +924,9 @@ unified ID counter. The funcmap records each ID with: | Funcmap name prefix | Instruction type | Gated by | |---|---|---| -| `addr_trace_load` | Global/flat load | `memory` | -| `addr_trace_store` | Global/flat store | `memory` | -| `addr_trace_atomic` | Global/flat atomic (AtomicRMW, CmpXchg) | `memory` | +| `addr_trace_load` | Flat/global (AS=0/1) load | `memory` | +| `addr_trace_store` | Flat/global (AS=0/1) store | `memory` | +| `addr_trace_atomic` | Flat/global (AS=0/1) atomic (AtomicRMW, CmpXchg) | `memory` | | `addr_trace_lds_load` | LDS load | `lds` | | `addr_trace_lds_store` | LDS store | `lds` | | `addr_trace_lds_atomic` | LDS atomic | `lds` | @@ -870,9 +971,8 @@ For each qualifying memory operation, the pass emits (before the op): 1. Allocate a unique ID from `NextEventID++` 2. Record the ID, kind name, and source location in `AddrTraceEntries` 3. Header point marker via `emitBareTrace()` with the unique ID -4. EXEC mask via inline asm: `s_mov_b32 $0, exec_lo` / `exec_hi` - (on gfx9, the exec read and `s_ttracedata` are fused into a single - inline asm block to satisfy the s_nop hazard requirement) +4. EXEC mask via inline asm: `s_mov_b32 m0, exec_lo` / `exec_hi`, then + `s_nop 0` and `s_ttracedata`. 5. For memory/LDS/atomics: `ptrtoint` of the pointer operand (i64 for memory, i32 for LDS) For buffers: extract rsrc_lo/hi, soffset as scalar values For permutes: extract the index operand (arg 0 of the intrinsic) @@ -898,11 +998,13 @@ in post-processing. The pass handles operand layout differences across load, store, atomic, and cmpswap variants. The rsrc operand may be `<4 x i32>` (legacy) or `ptr addrspace(8)` (modern) — both are supported. +`*.buffer.load.lds` is excluded because its LDS-destination operand layout is +not this protocol. ### Atomic operations `AtomicRMWInst` and `AtomicCmpXchgInst` are classified by address space the -same way as loads/stores (AS=3 → LDS, AS=5 → excluded, otherwise → memory). +same way as loads/stores (AS=3 → LDS, AS=0/1 → memory, all others excluded). They receive distinct funcmap names (`addr_trace_atomic`, `addr_trace_lds_atomic`) rather than being folded into load/store categories, so the decoder can distinguish atomic access patterns. diff --git a/projects/rocprof-trace-decoder/include/rocprof_trace_decoder/cxx/funcmap.hpp b/projects/rocprof-trace-decoder/include/rocprof_trace_decoder/cxx/funcmap.hpp index 6817c33d86f0..2ca0d0eaf9e2 100644 --- a/projects/rocprof-trace-decoder/include/rocprof_trace_decoder/cxx/funcmap.hpp +++ b/projects/rocprof-trace-decoder/include/rocprof_trace_decoder/cxx/funcmap.hpp @@ -106,16 +106,65 @@ struct FuncmapEntry uint32_t id{0}; // 0 for Kernel rows (no ID) std::string name{}; std::string source_loc{}; // empty if absent + uint64_t vaddr{0}; // resolved by CodeobjDecoderComponent; 0 if unresolved uint32_t extra_payload_count{0}; // following shaderdata records owned by this marker header - uint64_t vaddr{0}; // resolved by CodeobjDecoderComponent; 0 if unresolved }; +namespace detail +{ +constexpr uint32_t bit_mask(unsigned bits) noexcept +{ + return bits >= 32 ? 0xFFFFFFFFu : ((bits == 0) ? 0u : ((uint32_t(1) << bits) - 1u)); +} +} // namespace detail + struct MarkerEncoding { - uint32_t shader_clock_bits{0}; // 0 means legacy encoding + uint32_t shader_clock_bits{0}; // 0 means no packed shader-clock field uint32_t shader_clock_shift{0}; // source bit offset in shader_cycles_lo + // The instrumentation pass only emits layouts that retain at least one + // marker-ID bit and whose shader-clock window fits in a 32-bit source + // clock. + constexpr bool is_valid() const noexcept + { + return shader_clock_bits == 0 || + (shader_clock_bits <= 29 && shader_clock_shift < 32 && + shader_clock_bits <= 32 - shader_clock_shift); + } + constexpr bool has_shader_clock() const noexcept { return shader_clock_bits != 0; } + + /// Extract the right-aligned sampled shader-clock field from a raw marker + /// word. Malformed or legacy layouts return zero. + constexpr uint32_t decode_shader_clock(uint32_t raw) const noexcept + { + return (!has_shader_clock() || !is_valid()) + ? 0u + : (raw >> (32u - shader_clock_bits)) & detail::bit_mask(shader_clock_bits); + } + + /// Raw marker-word bits that contain the ID. The no-clock layout returns + /// 0xFFFFFFFC; malformed packed metadata falls back to that safe layout. + constexpr uint32_t marker_id_mask() const noexcept + { + return (!has_shader_clock() || !is_valid()) ? 0xFFFFFFFCu + : detail::bit_mask(30u - shader_clock_bits) << 2u; + } + + /// Raw marker-word bits occupied by the right-aligned sampled clock. + constexpr uint32_t packed_shader_clock_mask() const noexcept + { + return (!has_shader_clock() || !is_valid()) ? 0u + : detail::bit_mask(shader_clock_bits) << (32u - shader_clock_bits); + } + + /// Corresponding bit window in the producer's low shader-clock word. + constexpr uint32_t shader_clock_source_mask() const noexcept + { + return (!has_shader_clock() || !is_valid()) ? 0u + : detail::bit_mask(shader_clock_bits) << shader_clock_shift; + } }; struct FuncmapDiagnostic @@ -137,8 +186,8 @@ struct Funcmap std::vector entries{}; // owns rows, stable insertion order std::unordered_map by_id{}; // ID -> entry; last-writer-wins on dup uint32_t wave_size{0}; // 0 if no `W:` row - MarkerEncoding marker_encoding{}; std::vector diagnostics{}; // collected during parse / extraction + MarkerEncoding marker_encoding{}; // Returns the entry for `marker_id` (refcount bump, no string copy), or // nullptr if absent. @@ -150,7 +199,6 @@ struct MarkerValue uint32_t id; bool is_enter; bool exit_prev; - uint32_t shader_clock; }; // Parse the `.sqtt_funcmap` ASCII blob. Diagnostics go into @@ -185,31 +233,28 @@ inline Funcmap::EntryPtr Funcmap::find(uint32_t marker_id) const constexpr MarkerValue decode_marker_value(uint32_t v) noexcept { - return MarkerValue{v >> 2, bool((v >> 1) & 1u), bool(v & 1u), 0}; + return MarkerValue{v >> 2, bool((v >> 1) & 1u), bool(v & 1u)}; } -namespace detail -{ -constexpr uint32_t bit_mask(unsigned bits) noexcept -{ - return bits >= 32 ? 0xFFFFFFFFu : ((bits == 0) ? 0u : ((uint32_t(1) << bits) - 1u)); -} -} // namespace detail - constexpr MarkerValue decode_marker_value(uint32_t v, MarkerEncoding encoding) noexcept { - if (!encoding.has_shader_clock()) return decode_marker_value(v); + if (!encoding.has_shader_clock() || !encoding.is_valid()) return decode_marker_value(v); - uint32_t clock_bits = encoding.shader_clock_bits; - uint32_t id_bits = (clock_bits >= 30) ? 0 : (30 - clock_bits); - uint32_t clock = (clock_bits >= 32) ? v : ((v >> (32 - clock_bits)) & detail::bit_mask(clock_bits)); + uint32_t id_bits = 30 - encoding.shader_clock_bits; uint32_t id = (v >> 2) & detail::bit_mask(id_bits); - return MarkerValue{id, bool((v >> 1) & 1u), bool(v & 1u), clock}; + return MarkerValue{id, bool((v >> 1) & 1u), bool(v & 1u)}; } inline MarkerValue decode_marker_value(uint32_t v, const Funcmap& funcmap) noexcept { - return decode_marker_value(v, funcmap.marker_encoding); + MarkerValue marker = decode_marker_value(v, funcmap.marker_encoding); + if (!funcmap.marker_encoding.has_shader_clock() || !funcmap.marker_encoding.is_valid()) return marker; + + // Numeric point/enter markers intentionally retain their legacy values. + // A funcmap identifies the pass-generated headers that actually use the + // packed layout; a bare exit has no ID but is always packed. + if (funcmap.find(marker.id) || (marker.id == 0 && marker.exit_prev && !marker.is_enter)) return marker; + return decode_marker_value(v); } namespace detail @@ -363,6 +408,7 @@ inline Funcmap parse_funcmap_section(std::string_view blob, bool silent) { MarkerEncoding encoding{}; bool saw_clock_bits = false; + bool saw_clock_shift = false; bool ok = true; size_t attr_pos = 0; while (attr_pos <= payload.size()) @@ -391,13 +437,17 @@ inline Funcmap parse_funcmap_section(std::string_view blob, bool silent) if (!detail::parse_u32(value, parsed)) ok = false; else + { encoding.shader_clock_shift = parsed; + saw_clock_shift = true; + } } } if (attr_end >= payload.size()) break; attr_pos = attr_end + 1; } - if (!ok || (saw_clock_bits && encoding.shader_clock_bits > 29) || + if (!ok || (saw_clock_bits && encoding.shader_clock_bits > 0 && !saw_clock_shift) || + (saw_clock_bits && encoding.shader_clock_bits > 29) || (saw_clock_bits && encoding.shader_clock_shift > 32) || (saw_clock_bits && encoding.shader_clock_bits > 32 - encoding.shader_clock_shift)) { diff --git a/projects/rocprof-trace-decoder/markers/CMakeLists.txt b/projects/rocprof-trace-decoder/markers/CMakeLists.txt index d9f3889bca12..e8788eb7112d 100644 --- a/projects/rocprof-trace-decoder/markers/CMakeLists.txt +++ b/projects/rocprof-trace-decoder/markers/CMakeLists.txt @@ -63,7 +63,8 @@ endif() set_target_properties( SQTTInstrumentPass PROPERTIES NO_SONAME ON - PREFIX "" + PREFIX "lib" + OUTPUT_NAME "sqttinstrumentpass" POSITION_INDEPENDENT_CODE ON LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/bin" diff --git a/projects/rocprof-trace-decoder/markers/README.md b/projects/rocprof-trace-decoder/markers/README.md index 4125133516f5..010817bdd8ca 100644 --- a/projects/rocprof-trace-decoder/markers/README.md +++ b/projects/rocprof-trace-decoder/markers/README.md @@ -16,7 +16,7 @@ cmake -B build -DCMAKE_PREFIX_PATH=/opt/rocm cmake --build build ``` -This produces `build/lib/SQTTInstrumentPass.so`. +This produces `build/lib/libsqttinstrumentpass.so`. Use `CMAKE_PREFIX_PATH` to select a specific ROCm installation: @@ -46,7 +46,7 @@ __global__ void my_kernel(float *data, int n) { ``` ```bash -hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/SQTTInstrumentPass.so \ +hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -67,7 +67,7 @@ waves on matching CUs/SIMDs/WGs emit markers: # Only CU 0, SIMD 0, all workgroups SQTT_SCOPE_CU=0x1 SQTT_SCOPE_SIMD=0x1 \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -82,7 +82,7 @@ SQTT already generates wave start/end markers for them. ```bash SQTT_INSTRUMENT_FUNCTIONS=10 \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -91,7 +91,7 @@ Use weighted cost instead of instruction count: ```bash SQTT_INSTRUMENT_FUNCTIONS=cost:50 \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -102,7 +102,7 @@ Insert point markers around `s_barrier_signal` and `s_barrier_wait`: ```bash SQTT_INSTRUMENT_BARRIERS=1 \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -114,7 +114,7 @@ Insert point markers around groups of global/buffer/flat memory operations ```bash SQTT_INSTRUMENT_MEMORY=2:5 \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -134,7 +134,7 @@ cache line utilization analysis, stride detection, and coalescing analysis: ```bash SQTT_TRACE_ADDRESSES=memory \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -143,7 +143,7 @@ Trace LDS addresses only: ```bash SQTT_TRACE_ADDRESSES=lds \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -152,7 +152,7 @@ Trace both global and LDS: ```bash SQTT_TRACE_ADDRESSES=memory,lds \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -177,6 +177,10 @@ to report only active lanes. The parser demultiplexes interleaved records from concurrent waves by grouping on `(cu, simd, wave_id)` before parsing each wave's address block. +Any marker with `extra_payload_count > 0`, including address blocks and +`sqtt_marker_data()`, requires `SQTT_SHADER_CLOCK_BITS=0`. Shader-clock +packing is disabled by default; an explicit nonzero setting is an error. + ### Everything together All options compose. User markers, auto-function, auto-barrier, memory ops, @@ -189,7 +193,7 @@ SQTT_INSTRUMENT_MEMORY=2:5 \ SQTT_SCOPE_CU=0x1 \ SQTT_SCOPE_SIMD=0x1 \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ my_kernel.hip -o my_kernel ``` @@ -232,7 +236,7 @@ __global__ void my_kernel(float *data, int n) { } ``` -Named markers require the pass plugin (`-fpass-plugin=build/lib/SQTTInstrumentPass.so`). +Named markers require the pass plugin (`-fpass-plugin=build/lib/libsqttinstrumentpass.so`). The pass: 1. Finds all `sqtt_marker_enter("...")`, `sqtt_marker_exit("...")`, and @@ -280,10 +284,10 @@ directly via `getenv()`. | `SQTT_INSTRUMENT_FUNCTIONS` | `N` or `cost:N` | off | Instrument device functions exceeding threshold. `20` = instruction count > 20. `cost:100` = weighted cost > 100. | | `SQTT_INSTRUMENT_BARRIERS` | `0`, `1` | `0` | Instrument barriers. Consecutive signal+wait pairs fuse to a single marker. | | `SQTT_INSTRUMENT_MEMORY` | `N:M` | off | Instrument memory ops. N = ops per marker, M = max gap. `2:5` = 1 marker per 2 ops, sequence breaks at gap > 5. Covers global, buffer, flat (not LDS/scratch). | -| `SQTT_TRACE_ADDRESSES` | `memory`, `lds`, or both | off | Trace per-lane virtual addresses. Mutually exclusive with `SQTT_INSTRUMENT_MEMORY`. `memory` = global/buffer/flat, `lds` = LDS (AS=3). Expensive. | +| `SQTT_TRACE_ADDRESSES` | `memory`, `lds`, or both | off | Trace per-lane virtual addresses. Mutually exclusive with `SQTT_INSTRUMENT_MEMORY`. `memory` = flat/global (AS=0/1) and buffer, `lds` = LDS (AS=3). Expensive. | | `SQTT_MEM_BARRIER` | `none` / `asm` / `fence` (or `0` / `1` / `2`) | `fence` | Reordering boundary planted around every marker. `fence` (default) emits `fence syncscope("workgroup") acq_rel` tagged with AMDGPU local/LDS synchronization metadata, anchoring markers against optimizer and scheduler movement without marker-generated global cache invalidation. `asm` plants an empty `~{memory}` inline asm (IR/MIR-only constraint, no machine code). `none` disables both. Default favors marker accuracy; opt down for tight kernels. | -| `SQTT_SHADER_CLOCK_BITS` | unsigned integer | auto (`12` on gfx12, `0` elsewhere) | Number of gfx12 marker header high bits reserved for shader clock. `0` disables clock packing. | -| `SQTT_SHADER_CLOCK_SHIFT` | unsigned integer | `4` | Source bit offset in `HW_REG_SHADER_CYCLES_LO` for the packed shader clock field. | +| `SQTT_SHADER_CLOCK_BITS` | unsigned integer | `0` | Number of gfx12 marker header high bits reserved for shader clock. Set a nonzero value to enable packing; it is invalid with any payload-bearing marker. | +| `SQTT_SHADER_CLOCK_SHIFT` | unsigned integer | `4` | Source bit offset in the shader-clock low word for the packed field. | ## Examples @@ -339,7 +343,7 @@ sqtt_marker_exit("Consumer Thread"); Build with the pass plugin and capture a trace: ```bash -hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/SQTTInstrumentPass.so \ +hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ test/markers/kernels/heavy.cpp -o heavy rocprofv3 --att -d trace -- ./heavy @@ -381,7 +385,7 @@ Build with automatic function instrumentation and capture a trace: ```bash SQTT_INSTRUMENT_FUNCTIONS=10 \ -hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/SQTTInstrumentPass.so \ +hipcc -DSQTT_ENABLED=1 -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ test/markers/kernels/auto.cpp -o auto rocprofv3 --att -d trace -- ./auto @@ -445,6 +449,10 @@ Marker values are packed into 2 flag bits so that small IDs can use the faster packing is inactive. Larger IDs fall back to `s_ttracedata` (32-bit m0, all targets). +For full M0-based traces, the pass explicitly emits `s_mov_b32 m0`, `s_nop 0`, +then `s_ttracedata` on every target. Address EXEC tracing uses the same +sequence. + ``` Both instructions share the same bit layout: @@ -463,12 +471,37 @@ The pass plugin automatically selects `s_ttracedata_imm` when the encoded value fits in 8 bits, the target supports it, and shader-clock packing is inactive. IDs 1-63 use the fast path on RDNA. Exit markers use the same selection rule. -On gfx12+, the default is to pack 12 shader-clock bits from -`s_getreg_b32 hwreg(HW_REG_SHADER_CYCLES_LO, 4, 12)` into marker bits 31-20, -leaving 18 marker ID bits plus the two low flags. Configure this with -`SQTT_SHADER_CLOCK_BITS` and `SQTT_SHADER_CLOCK_SHIFT`; set -`SQTT_SHADER_CLOCK_BITS=0` to use the legacy layout on gfx12. The funcmap emits +On gfx1200, gfx1201, and gfx1250, a nonzero `SQTT_SHADER_CLOCK_BITS` enables +clock packing. With the usual 12-bit setting, it occupies marker bits 31-20 +and leaves 18 marker ID bits plus the two low flags. gfx1250 uses +`s_get_shader_cycles_u64`; gfx1200 and gfx1201 use +`s_getreg_b32 hwreg(HW_REG_SHADER_CYCLES_LO, 4, 12)`. Configure this with +`SQTT_SHADER_CLOCK_BITS` and `SQTT_SHADER_CLOCK_SHIFT`; it defaults to `0`, +which uses the no-clock full-ID layout. The funcmap emits `M:shader_clock_bits=N;shader_clock_shift=S` whenever clock packing is active. +Clock packing cannot be combined with any marker that has an `R:` payload row. +Numeric point and enter markers keep their legacy values; bare exits use the +packed form because their value is indistinguishable from a generated exit. + +The packed clock is a truncated sample at instruction issue. Shaderdata and +shader clocks have an unknown fixed phase per physical SIMD, so `att_tool.py` +cancels it by comparing clock/time deltas to an arbitrary header in the same +SIMD/layout. It takes the minimum from the complete domain and applies the +relative correction retroactively. Apply this only to headers, never `R:` +payloads. `Funcmap::marker_encoding` exposes +`marker_id_mask()`, `packed_shader_clock_mask()`, and +`shader_clock_source_mask()`; without `M:` metadata those describe the legacy +full-width ID layout. + +`att_tool.py` applies that correction by default and masks each recognized +header to its ordinary legacy-form marker value for existing viewers. Use +`--no-decode-markers` to preserve packed values and arrival timestamps. +For externally packed payload blocks it may normalize a recognized header, +but does not apply timestamp correction. + +For packed code objects, `att_tool.py` adds six-column `sqtt_funcmap` rows and +packed-layout/payload tables to `code.json`. Readers that do not know them +continue using the original JSON; legacy `code.json` output is unchanged. The marker type (function, user, barrier, memory) is determined by looking up the ID in the `.sqtt_funcmap` section, not from encoding bits. @@ -508,7 +541,7 @@ Check that markers appear in the IR: ```bash SQTT_INSTRUMENT_FUNCTIONS=5 \ hipcc -DSQTT_ENABLED=1 \ - -fpass-plugin=build/lib/SQTTInstrumentPass.so \ + -fpass-plugin=build/lib/libsqttinstrumentpass.so \ -I include/rocprof_trace_decoder/cxx/ -S -emit-llvm my_kernel.hip -o - \ | grep ttracedata ``` diff --git a/projects/rocprof-trace-decoder/markers/SQTTAddressTrace.cpp b/projects/rocprof-trace-decoder/markers/SQTTAddressTrace.cpp index 277903dc0fd7..0a397e658a3b 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTAddressTrace.cpp +++ b/projects/rocprof-trace-decoder/markers/SQTTAddressTrace.cpp @@ -28,6 +28,28 @@ using namespace llvm; +namespace +{ + +// Keep the explicit EXEC trace sequence together on every target. +void emitExecMaskTrace(IRBuilder<>& B, LLVMContext& Ctx, bool high) +{ + const char* Asm = high ? "s_mov_b32 m0, exec_hi\n" + "s_nop 0\n" + "s_ttracedata" + : "s_mov_b32 m0, exec_lo\n" + "s_nop 0\n" + "s_ttracedata"; + // Model the M0 definition as a fixed output. `m0` is reserved on AMDGPU, + // so a `~{m0}` clobber is diagnosed as undefined behavior by LLVM. + InlineAsm* trace = InlineAsm::get( + FunctionType::get(Type::getInt32Ty(Ctx), false), Asm, "={m0}", /*hasSideEffects=*/true + ); + B.CreateCall(trace); +} + +} // namespace + SQTTInstrumentPass::AddrTraceKind SQTTInstrumentPass::classifyAddrTraceOp( Instruction* I, bool traceMemory, bool traceLDS ) @@ -35,9 +57,12 @@ SQTTInstrumentPass::AddrTraceKind SQTTInstrumentPass::classifyAddrTraceOp( // Helper lambda for pointer-based address space classification auto classifyAS = [&](unsigned AS) -> AddrTraceKind { - if (AS == 5) return AddrTraceKind::None; // private if (AS == 3) return traceLDS ? AddrTraceKind::LDS : AddrTraceKind::None; - return traceMemory ? AddrTraceKind::Memory : AddrTraceKind::None; + // The memory address protocol is defined only for flat (0) and + // global (1) pointers. Other AMDGPU address spaces have different + // representations and must not be reinterpreted as global addresses. + if (AS == 0 || AS == 1) return traceMemory ? AddrTraceKind::Memory : AddrTraceKind::None; + return AddrTraceKind::None; }; if (auto* LI = dyn_cast(I)) return classifyAS(LI->getPointerAddressSpace()); @@ -50,7 +75,12 @@ SQTTInstrumentPass::AddrTraceKind SQTTInstrumentPass::classifyAddrTraceOp( Function* Callee = CI->getCalledFunction(); if (!Callee) return AddrTraceKind::None; StringRef Name = Callee->getName(); - if (traceMemory && isBufferOp(Name)) return AddrTraceKind::Buffer; + // buffer.load.lds has a distinct operand layout (including an LDS + // destination pointer), so the ordinary buffer component protocol is + // not valid for it. Leave it uninstrumented until it has its own + // protocol. + if (traceMemory && isBufferOp(Name) && !Name.ends_with(".buffer.load.lds")) + return AddrTraceKind::Buffer; auto IID = Callee->getIntrinsicID(); if (traceLDS && (IID == Intrinsic::amdgcn_ds_permute || IID == Intrinsic::amdgcn_ds_bpermute || IID == Intrinsic::amdgcn_ds_bpermute_fi_b32)) @@ -132,15 +162,19 @@ void SQTTInstrumentPass::emitAddressTrace( Type* I32 = Type::getInt32Ty(Ctx); Type* I64 = Type::getInt64Ty(Ctx); + emitTraceBlockBoundary(B, /*after=*/false); + // Buffer and permute ops use specialized trace protocols if (kind == AddrTraceKind::Buffer) { emitBufferTrace(B, cast(memOp), headerID, F, gen); + emitTraceBlockBoundary(B, /*after=*/true); return; } if (kind == AddrTraceKind::Permute) { emitPermuteTrace(B, cast(memOp), headerID, F, gen); + emitTraceBlockBoundary(B, /*after=*/true); return; } @@ -168,42 +202,8 @@ void SQTTInstrumentPass::emitAddressTrace( // 3. EXEC mask unsigned waveSize = getWaveSize(gen); - // gfx9: the ISA spec requires s_nop 0 between s_mov_b32 m0 and - // s_ttracedata. We can't insert a nop between the backend-generated - // s_mov_b32 m0 and s_ttracedata pair, so for exec reads we emit the - // entire sequence as a single inline asm block. - if (gen == GfxGen::GFX9) - { - InlineAsm* traceExecLo = InlineAsm::get( - FunctionType::get(Type::getVoidTy(Ctx), false), - "s_mov_b32 m0, exec_lo\n" - "s_nop 0\n" - "s_ttracedata", - "", - /*hasSideEffects=*/true - ); - B.CreateCall(traceExecLo); - - InlineAsm* traceExecHi = InlineAsm::get( - FunctionType::get(Type::getVoidTy(Ctx), false), - "s_mov_b32 m0, exec_hi\n" - "s_nop 0\n" - "s_ttracedata", - "", - /*hasSideEffects=*/true - ); - B.CreateCall(traceExecHi); - } - else - { - InlineAsm* readExecLo = - InlineAsm::get(FunctionType::get(I32, false), "s_mov_b32 $0, exec_lo", "=s", /*hasSideEffects=*/true); - emitBareTraceValue(B, B.CreateCall(readExecLo), M, gen); - - InlineAsm* readExecHi = - InlineAsm::get(FunctionType::get(I32, false), "s_mov_b32 $0, exec_hi", "=s", /*hasSideEffects=*/true); - emitBareTraceValue(B, B.CreateCall(readExecHi), M, gen); - } + emitExecMaskTrace(B, Ctx, /*high=*/false); + emitExecMaskTrace(B, Ctx, /*high=*/true); // 4. Readlane loop Function* ReadLane = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_readlane, {I32}); @@ -256,6 +256,7 @@ void SQTTInstrumentPass::emitAddressTrace( // Restore insertion point to AfterBB B.SetInsertPoint(AfterBB, AfterBB->begin()); + emitTraceBlockBoundary(B, /*after=*/true); } bool SQTTInstrumentPass::instrumentAddressTraces(Function& F, GfxGen gen) @@ -424,38 +425,8 @@ void SQTTInstrumentPass::emitBufferTrace(IRBuilder<>& B, CallInst* bufOp, uint32 // 2. EXEC mask (same as memory/LDS traces) unsigned waveSize = getWaveSize(gen); - if (gen == GfxGen::GFX9) - { - InlineAsm* traceExecLo = InlineAsm::get( - FunctionType::get(Type::getVoidTy(Ctx), false), - "s_mov_b32 m0, exec_lo\n" - "s_nop 0\n" - "s_ttracedata", - "", - /*hasSideEffects=*/true - ); - B.CreateCall(traceExecLo); - - InlineAsm* traceExecHi = InlineAsm::get( - FunctionType::get(Type::getVoidTy(Ctx), false), - "s_mov_b32 m0, exec_hi\n" - "s_nop 0\n" - "s_ttracedata", - "", - /*hasSideEffects=*/true - ); - B.CreateCall(traceExecHi); - } - else - { - InlineAsm* readExecLo = - InlineAsm::get(FunctionType::get(I32, false), "s_mov_b32 $0, exec_lo", "=s", /*hasSideEffects=*/true); - emitBareTraceValue(B, B.CreateCall(readExecLo), M, gen); - - InlineAsm* readExecHi = - InlineAsm::get(FunctionType::get(I32, false), "s_mov_b32 $0, exec_hi", "=s", /*hasSideEffects=*/true); - emitBareTraceValue(B, B.CreateCall(readExecHi), M, gen); - } + emitExecMaskTrace(B, Ctx, /*high=*/false); + emitExecMaskTrace(B, Ctx, /*high=*/true); // 3. rsrc base (lo/hi words) Value* rsrc = ops.Rsrc; @@ -556,38 +527,8 @@ void SQTTInstrumentPass::emitPermuteTrace( // 2. EXEC mask unsigned waveSize = getWaveSize(gen); - if (gen == GfxGen::GFX9) - { - InlineAsm* traceExecLo = InlineAsm::get( - FunctionType::get(Type::getVoidTy(Ctx), false), - "s_mov_b32 m0, exec_lo\n" - "s_nop 0\n" - "s_ttracedata", - "", - /*hasSideEffects=*/true - ); - B.CreateCall(traceExecLo); - - InlineAsm* traceExecHi = InlineAsm::get( - FunctionType::get(Type::getVoidTy(Ctx), false), - "s_mov_b32 m0, exec_hi\n" - "s_nop 0\n" - "s_ttracedata", - "", - /*hasSideEffects=*/true - ); - B.CreateCall(traceExecHi); - } - else - { - InlineAsm* readExecLo = - InlineAsm::get(FunctionType::get(I32, false), "s_mov_b32 $0, exec_lo", "=s", /*hasSideEffects=*/true); - emitBareTraceValue(B, B.CreateCall(readExecLo), M, gen); - - InlineAsm* readExecHi = - InlineAsm::get(FunctionType::get(I32, false), "s_mov_b32 $0, exec_hi", "=s", /*hasSideEffects=*/true); - emitBareTraceValue(B, B.CreateCall(readExecHi), M, gen); - } + emitExecMaskTrace(B, Ctx, /*high=*/false); + emitExecMaskTrace(B, Ctx, /*high=*/true); // 3. Per-lane readlane loop over index operand (arg 0) Value* index = permuteOp->getArgOperand(0); diff --git a/projects/rocprof-trace-decoder/markers/SQTTBarriers.cpp b/projects/rocprof-trace-decoder/markers/SQTTBarriers.cpp index 29b54e7a3866..6c9901d55b1c 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTBarriers.cpp +++ b/projects/rocprof-trace-decoder/markers/SQTTBarriers.cpp @@ -64,7 +64,7 @@ bool SQTTInstrumentPass::instrumentBarriers(Function& F, GfxGen gen) auto [CI, kind] = Barriers[i]; if (kind == BarrierKind::Signal && i + 1 < Barriers.size() && Barriers[i + 1].second == BarrierKind::Wait && - CI->getParent() == Barriers[i + 1].first->getParent()) + CI->getNextNode() == Barriers[i + 1].first) { CallInst* WaitCI = Barriers[i + 1].first; IRBuilder<> B(WaitCI); diff --git a/projects/rocprof-trace-decoder/markers/SQTTConfig.h b/projects/rocprof-trace-decoder/markers/SQTTConfig.h index 4bf70d06aa5a..8c88428983de 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTConfig.h +++ b/projects/rocprof-trace-decoder/markers/SQTTConfig.h @@ -25,7 +25,6 @@ #include #include #include -#include #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringRef.h" @@ -104,8 +103,6 @@ enum class MemBarrierMode struct SQTTConfig { - static constexpr unsigned AutoShaderClockBits = std::numeric_limits::max(); - bool InstrumentBarriers = false; CostMode Mode = CostMode::InstructionCount; unsigned FunctionThreshold = 0; // 0 = disabled @@ -119,7 +116,7 @@ struct SQTTConfig MemBarrierMode MemBarrier = MemBarrierMode::Fence; bool TraceMemoryAddrs = false; // trace global/buffer/flat addresses bool TraceLDSAddrs = false; // trace LDS addresses - unsigned ShaderClockBits = AutoShaderClockBits; // auto: gfx12 defaults to clock packing + unsigned ShaderClockBits = 0; // opt in to clock packing explicitly unsigned ShaderClockShift = 4; bool hasAddressTracing() const { return TraceMemoryAddrs || TraceLDSAddrs; } @@ -197,7 +194,7 @@ struct SQTTConfig c.SimdMask = parseEnvMask("SQTT_SCOPE_SIMD", 0xF); c.CuMask = parseEnvMask("SQTT_SCOPE_CU", 0x3); c.WgMask = parseEnvMask("SQTT_SCOPE_WG", 0xFFFFFFFF); - c.ShaderClockBits = parseEnvUnsigned("SQTT_SHADER_CLOCK_BITS", AutoShaderClockBits); + c.ShaderClockBits = parseEnvUnsigned("SQTT_SHADER_CLOCK_BITS", 0); c.ShaderClockShift = parseEnvUnsigned("SQTT_SHADER_CLOCK_SHIFT", 4); const char* funcEnv = std::getenv("SQTT_INSTRUMENT_FUNCTIONS"); diff --git a/projects/rocprof-trace-decoder/markers/SQTTFunctions.cpp b/projects/rocprof-trace-decoder/markers/SQTTFunctions.cpp index 0889dfc5f30b..3cee231f7ea3 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTFunctions.cpp +++ b/projects/rocprof-trace-decoder/markers/SQTTFunctions.cpp @@ -233,6 +233,7 @@ uint32_t SQTTInstrumentPass::compactFuncIDs(Module& M) auto IID = Callee->getIntrinsicID(); if (IID != Intrinsic::amdgcn_s_ttracedata && IID != Intrinsic::amdgcn_s_ttracedata_imm) continue; if (isRawPayloadTrace(CI)) continue; + if (!isMarkerHeaderTrace(CI)) continue; auto* Arg = dyn_cast(CI->getArgOperand(0)); if (!Arg) continue; uint32_t val = Arg->getZExtValue(); @@ -319,6 +320,7 @@ void SQTTInstrumentPass::rewriteMarkerIDs(Function& F, const std::mapgetIntrinsicID(); if (IID != Intrinsic::amdgcn_s_ttracedata && IID != Intrinsic::amdgcn_s_ttracedata_imm) continue; if (isRawPayloadTrace(CI)) continue; + if (!isMarkerHeaderTrace(CI)) continue; auto* Arg = dyn_cast(CI->getArgOperand(0)); if (!Arg) continue; @@ -353,16 +355,18 @@ void SQTTInstrumentPass::rewriteMarkerIDs(Function& F, const std::map B(CI); + CallInst* replacement; if (canUseImm(newVal) && supportsImmTrace(gen)) { Function* TTDImm = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_ttracedata_imm); - B.CreateCall(TTDImm, {ConstantInt::get(I16, newVal)}); + replacement = B.CreateCall(TTDImm, {ConstantInt::get(I16, newVal)}); } else { Function* TTD = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_ttracedata); - B.CreateCall(TTD, {ConstantInt::get(I32, newVal)}); + replacement = B.CreateCall(TTD, {ConstantInt::get(I32, newVal)}); } + replacement->copyMetadata(*CI); CI->eraseFromParent(); } } @@ -387,6 +391,7 @@ void SQTTInstrumentPass::removeFuncMarkersFromModule(Module& M, uint32_t id) auto IID = Callee->getIntrinsicID(); if (IID != Intrinsic::amdgcn_s_ttracedata && IID != Intrinsic::amdgcn_s_ttracedata_imm) continue; if (isRawPayloadTrace(CI)) continue; + if (!isMarkerHeaderTrace(CI)) continue; auto* Arg = dyn_cast(CI->getArgOperand(0)); if (!Arg) continue; @@ -441,8 +446,18 @@ void SQTTInstrumentPass::removeAdjacentBarriers(CallInst* CI) } } +bool SQTTInstrumentPass::hasMustTailCall(const Function& F) +{ + for (const auto& BB : F) + for (const auto& I : BB) + if (const auto* CB = dyn_cast(&I); CB && CB->isMustTailCall()) return true; + return false; +} + bool SQTTInstrumentPass::instrumentFunctionDirect(Function& F, GfxGen gen) { + if (hasMustTailCall(F)) return false; + unsigned size = computeFunctionSize(F, Config.Mode); if (size <= Config.FunctionThreshold) return false; diff --git a/projects/rocprof-trace-decoder/markers/SQTTMarkers.cpp b/projects/rocprof-trace-decoder/markers/SQTTMarkers.cpp index eaf15911097b..804e3799bc1b 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTMarkers.cpp +++ b/projects/rocprof-trace-decoder/markers/SQTTMarkers.cpp @@ -31,6 +31,7 @@ using namespace llvm; +static constexpr const char* MarkerHeaderMetadataName = "sqtt.marker_header"; static constexpr const char* RawPayloadMetadataName = "sqtt.raw_payload"; static constexpr uint32_t GFX12_SHADER_CYCLES_LO = 29; @@ -88,6 +89,26 @@ static void emitMemBarrier(IRBuilder<>& B, MemBarrierMode mode) F->setMetadata(LLVMContext::MD_mmra, MDNode::get(Ctx, LocalSyncAS)); } +void SQTTInstrumentPass::emitTraceBlockBoundary(IRBuilder<>& B, bool after) +{ + Module* M = B.GetInsertBlock()->getParent()->getParent(); + Type* I32 = Type::getInt32Ty(B.getContext()); + Function* SchedBarrier = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_sched_barrier); + + // Keep the complete address block between two scheduling boundaries, + // rather than adding one around every raw payload. + if (after) + { + emitMemBarrier(B, Config.MemBarrier); + B.CreateCall(SchedBarrier, {ConstantInt::get(I32, 0)}); + } + else + { + B.CreateCall(SchedBarrier, {ConstantInt::get(I32, 0)}); + emitMemBarrier(B, Config.MemBarrier); + } +} + void SQTTInstrumentPass::insertTraceMarker(IRBuilder<>& B, uint32_t markerID, Function& F, GfxGen gen) { insertTraceMarkerWithPayload(B, markerID, nullptr, F, gen); @@ -124,16 +145,18 @@ void SQTTInstrumentPass::insertTraceMarkerWithPayload( } emitMemBarrier(Builder, Config.MemBarrier); + CallInst* trace; if (canUseImm(markerID) && supportsImmTrace(gen)) { Function* TTDImm = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_ttracedata_imm); - Builder.CreateCall(TTDImm, {ConstantInt::get(I16, markerID)}); + trace = Builder.CreateCall(TTDImm, {ConstantInt::get(I16, markerID)}); } else { Function* TTD = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_ttracedata); - Builder.CreateCall(TTD, {ConstantInt::get(I32, markerID)}); + trace = Builder.CreateCall(TTD, {ConstantInt::get(I32, markerID)}); } + markMarkerHeaderTrace(trace); if (payload) { @@ -399,16 +422,19 @@ bool SQTTInstrumentPass::addBarriersToExistingMarkers(Function& F) CallInst* SQTTInstrumentPass::emitBareTrace(IRBuilder<>& B, uint32_t encoded, Module* M, GfxGen gen) { LLVMContext& Ctx = M->getContext(); + CallInst* CI; if (canUseImm(encoded) && supportsImmTrace(gen)) { Function* TTDImm = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_ttracedata_imm); - return B.CreateCall(TTDImm, {ConstantInt::get(Type::getInt16Ty(Ctx), encoded)}); + CI = B.CreateCall(TTDImm, {ConstantInt::get(Type::getInt16Ty(Ctx), encoded)}); } else { Function* TTD = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_ttracedata); - return B.CreateCall(TTD, {ConstantInt::get(Type::getInt32Ty(Ctx), encoded)}); + CI = B.CreateCall(TTD, {ConstantInt::get(Type::getInt32Ty(Ctx), encoded)}); } + markMarkerHeaderTrace(CI); + return CI; } CallInst* SQTTInstrumentPass::emitBareTraceValue(IRBuilder<>& B, Value* val, Module* M, GfxGen gen) @@ -422,6 +448,17 @@ CallInst* SQTTInstrumentPass::emitBareTraceValue(IRBuilder<>& B, Value* val, Mod CallInst* SQTTInstrumentPass::emitRawTracePayload(IRBuilder<>& B, Value* val, Module* M) { + // Full s_ttracedata lowers through M0, so its input must be scalar. The + // intrinsic lowering normally inserts this readfirstlane for a divergent + // named data value; retain that behavior before the explicit asm lowering. + Type* I32 = Type::getInt32Ty(M->getContext()); + if (val->getType() != I32) val = B.CreateZExtOrTrunc(val, I32); + if (!isa(val)) + { + Function* ReadFirstLane = + Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_readfirstlane, {I32}); + val = B.CreateCall(ReadFirstLane, {val}); + } Function* TTD = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_ttracedata); CallInst* CI = B.CreateCall(TTD, {val}); markRawPayloadTrace(CI); @@ -439,6 +476,60 @@ bool SQTTInstrumentPass::isRawPayloadTrace(CallInst* CI) return CI && CI->getMetadata(RawPayloadMetadataName) != nullptr; } +void SQTTInstrumentPass::markMarkerHeaderTrace(CallInst* CI) +{ + if (!CI) return; + CI->setMetadata(MarkerHeaderMetadataName, MDNode::get(CI->getContext(), {})); +} + +bool SQTTInstrumentPass::isMarkerHeaderTrace(CallInst* CI) +{ + return CI && CI->getMetadata(MarkerHeaderMetadataName) != nullptr; +} + +bool SQTTInstrumentPass::lowerFullTracesWithM0Nop(Function& F) +{ + SmallVector traces; + for (auto& BB : F) + { + for (auto& I : BB) + { + auto* CI = dyn_cast(&I); + if (!CI) continue; + Function* Callee = CI->getCalledFunction(); + if (Callee && Callee->getIntrinsicID() == Intrinsic::amdgcn_s_ttracedata) traces.push_back(CI); + } + } + if (traces.empty()) return false; + + LLVMContext& Ctx = F.getContext(); + Type* I32 = Type::getInt32Ty(Ctx); + // Model M0 as a fixed output instead of a clobber. M0 is reserved on + // AMDGPU and LLVM diagnoses `~{m0}` as undefined behavior. + FunctionType* TraceTy = FunctionType::get(I32, {I32}, false); + static constexpr const char TraceAsmText[] = + "s_mov_b32 m0, $1\n" + "s_nop 0\n" + "s_ttracedata"; + // Constants can be encoded directly in the M0 move. Non-constant values + // reaching this helper are required to be scalar. + InlineAsm* ImmediateTrace = + InlineAsm::get(TraceTy, TraceAsmText, "={m0},i", /*hasSideEffects=*/true); + InlineAsm* ScalarTrace = InlineAsm::get(TraceTy, TraceAsmText, "={m0},s", /*hasSideEffects=*/true); + + for (CallInst* CI : traces) + { + Value* value = CI->getArgOperand(0); + InlineAsm* TraceAsm = isa(value) ? ImmediateTrace : ScalarTrace; + IRBuilder<> B(CI); + CallInst* replacement = B.CreateCall(TraceAsm, {value}); + replacement->setDebugLoc(CI->getDebugLoc()); + replacement->copyMetadata(*CI); + CI->eraseFromParent(); + } + return true; +} + bool SQTTInstrumentPass::applyShaderClockPacking(Function& F, GfxGen gen) { unsigned clockBits = getShaderClockBits(Config, gen); @@ -468,9 +559,13 @@ bool SQTTInstrumentPass::applyShaderClockPacking(Function& F, GfxGen gen) if (IID != Intrinsic::amdgcn_s_ttracedata && IID != Intrinsic::amdgcn_s_ttracedata_imm) continue; Value* encodedValue = CI->getArgOperand(0); + bool header = isMarkerHeaderTrace(CI); if (auto* Arg = dyn_cast(encodedValue)) { uint32_t encoded = Arg->getZExtValue(); + // A bare exit has no ID and cannot be distinguished from the + // numeric API in trace data, so keep all of them packed. + if (!header && encoded != FLAG_EXIT_PREV) continue; uint32_t markerID = encoded >> 2; if (markerID > maxID) { @@ -480,6 +575,8 @@ bool SQTTInstrumentPass::applyShaderClockPacking(Function& F, GfxGen gen) ); } } + else if (!header) + continue; Rewrites.push_back({CI, encodedValue}); } } @@ -488,7 +585,19 @@ bool SQTTInstrumentPass::applyShaderClockPacking(Function& F, GfxGen gen) Module* M = F.getParent(); LLVMContext& Ctx = M->getContext(); Type* I32 = Type::getInt32Ty(Ctx); - Function* SGetReg = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_getreg); + const bool useShaderCyclesU64 = hasShaderCyclesU64(F); + Function* SGetReg = + useShaderCyclesU64 ? nullptr : Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_getreg); + InlineAsm* GetShaderCycles = nullptr; + if (useShaderCyclesU64) + { + GetShaderCycles = InlineAsm::get( + FunctionType::get(Type::getInt64Ty(Ctx), false), + "s_get_shader_cycles_u64 $0", + "=s", + /*hasSideEffects=*/true + ); + } Function* TTD = Intrinsic::getOrInsertDeclaration(M, Intrinsic::amdgcn_s_ttracedata); uint32_t hwreg = GETREG_IMMED(clockBits - 1, Config.ShaderClockShift, GFX12_SHADER_CYCLES_LO); uint32_t clockDestShift = 32 - clockBits; @@ -499,10 +608,21 @@ bool SQTTInstrumentPass::applyShaderClockPacking(Function& F, GfxGen gen) Value* markerAndFlags = encoded; if (markerAndFlags->getType() != I32) markerAndFlags = B.CreateZExtOrTrunc(markerAndFlags, I32); markerAndFlags = B.CreateAnd(markerAndFlags, ConstantInt::get(I32, markerAndFlagMask)); - Value* clock = B.CreateCall(SGetReg, {ConstantInt::get(I32, hwreg)}); + Value* clock; + if (GetShaderCycles) + { + Value* cyclesLo = B.CreateTrunc(B.CreateCall(GetShaderCycles), I32); + // The final shift below discards all but clockBits, matching the + // right-aligned field returned by s_getreg on gfx1200 and gfx1201. + clock = B.CreateLShr(cyclesLo, ConstantInt::get(I32, Config.ShaderClockShift)); + } + else + clock = B.CreateCall(SGetReg, {ConstantInt::get(I32, hwreg)}); Value* shiftedClock = B.CreateShl(clock, ConstantInt::get(I32, clockDestShift)); Value* packed = B.CreateOr(shiftedClock, markerAndFlags); - B.CreateCall(TTD, {packed}); + CallInst* packedTrace = B.CreateCall(TTD, {packed}); + packedTrace->setDebugLoc(CI->getDebugLoc()); + packedTrace->copyMetadata(*CI); CI->eraseFromParent(); } diff --git a/projects/rocprof-trace-decoder/markers/SQTTNamedMarkers.cpp b/projects/rocprof-trace-decoder/markers/SQTTNamedMarkers.cpp index f8adddbfb104..457bf242130f 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTNamedMarkers.cpp +++ b/projects/rocprof-trace-decoder/markers/SQTTNamedMarkers.cpp @@ -111,13 +111,12 @@ bool SQTTInstrumentPass::processMarkerCalls(Function& F, GfxGen gen, bool useBar auto& [CI, Type] = Calls[i]; if (!CI) continue; // already consumed by fusion - // Try to fuse exit+enter pairs: if this is an exit and the - // next call is an enter in the same basic block, emit a single - // marker with exit_prev=true. + // Fuse only directly adjacent exit+enter pairs. A marker boundary + // must not absorb work between the calls. if (Type == MarkerType::Exit && i + 1 < Calls.size()) { auto& [NextCI, NextType] = Calls[i + 1]; - if (NextCI && NextType == MarkerType::Enter && CI->getParent() == NextCI->getParent()) + if (NextCI && NextType == MarkerType::Enter && CI->getNextNode() == NextCI) { uint32_t enterEncoded = resolveMarkerString(NextCI, MarkerType::Enter); if (enterEncoded) diff --git a/projects/rocprof-trace-decoder/markers/SQTTPass.cpp b/projects/rocprof-trace-decoder/markers/SQTTPass.cpp index e3c71e86451d..c75c7800a58d 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTPass.cpp +++ b/projects/rocprof-trace-decoder/markers/SQTTPass.cpp @@ -25,8 +25,8 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/Analysis/InlineCost.h" #include "llvm/IR/Constants.h" -#include "llvm/IR/InlineAsm.h" #include "llvm/Passes/PassBuilder.h" +#include "llvm/Support/ErrorHandling.h" #if __has_include("llvm/Plugins/PassPlugin.h") # include "llvm/Plugins/PassPlugin.h" #else @@ -112,11 +112,11 @@ PreservedAnalyses SQTTInstrumentPass::runEarly(Module& M) if (F.getCallingConv() == CallingConv::AMDGPU_KERNEL) continue; if (Config.FunctionThreshold == 0) continue; + if (hasMustTailCall(F)) continue; uint32_t id = NextEventID++; - - insertBareMarkers(F, id, Gen); storeFuncMetadata(F, id, Ctx); + insertBareMarkers(F, id, Gen); Changed = true; } @@ -166,7 +166,7 @@ PreservedAnalyses SQTTInstrumentPass::runLate(Module& M) if (Config.needsScopeCheck()) Changed |= wrapExistingMarkers(F, Gen); - if (hadEarlyPass) Changed |= addBarriersToExistingMarkers(F); + Changed |= addBarriersToExistingMarkers(F); Changed |= processNamedMarkers(F, Gen); @@ -177,8 +177,24 @@ PreservedAnalyses SQTTInstrumentPass::runLate(Module& M) if (Config.hasAddressTracing()) Changed |= instrumentAddressTraces(F, Gen); if (!hadEarlyPass && Config.FunctionThreshold > 0 && !IsKernel) Changed |= instrumentFunctionDirect(F, Gen); + } + + bool HasPayloadMarkers = false; + for (const auto& Entry : UserMarkers) HasPayloadMarkers |= Entry.ExtraPayloadCount != 0; + for (const auto& Entry : AddrTraceEntries) HasPayloadMarkers |= Entry.ExtraPayloadCount != 0; + if (HasPayloadMarkers && Config.ShaderClockBits != 0) + report_fatal_error( + "SQTT payload markers require SQTT_SHADER_CLOCK_BITS=0" + ); + // Packing must happen after the address decision for the entire module. + for (auto& F : M) + { + if (F.isDeclaration()) continue; + GfxGen Gen = getGfxGen(F); + if (Gen == GfxGen::Unknown) continue; Changed |= applyShaderClockPacking(F, Gen); + Changed |= lowerFullTracesWithM0Nop(F); } // Clean up sentinel declarations diff --git a/projects/rocprof-trace-decoder/markers/SQTTPass.h b/projects/rocprof-trace-decoder/markers/SQTTPass.h index c588cf9114fc..1c7a4102e11a 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTPass.h +++ b/projects/rocprof-trace-decoder/markers/SQTTPass.h @@ -177,6 +177,8 @@ class SQTTInstrumentPass : public llvm::PassInfoMixin llvm::CallInst* emitBareTrace(llvm::IRBuilder<>& B, uint32_t encoded, llvm::Module* M, GfxGen gen); llvm::CallInst* emitBareTraceValue(llvm::IRBuilder<>& B, llvm::Value* val, llvm::Module* M, GfxGen gen); llvm::CallInst* emitRawTracePayload(llvm::IRBuilder<>& B, llvm::Value* val, llvm::Module* M); + static void markMarkerHeaderTrace(llvm::CallInst* CI); + static bool isMarkerHeaderTrace(llvm::CallInst* CI); static void markRawPayloadTrace(llvm::CallInst* CI); static bool isRawPayloadTrace(llvm::CallInst* CI); @@ -230,6 +232,9 @@ class SQTTInstrumentPass : public llvm::PassInfoMixin llvm::Function& F, GfxGen gen ); + // One ordering boundary on each side of a complete address-trace block. + // This deliberately does not put a boundary around every raw payload. + void emitTraceBlockBoundary(llvm::IRBuilder<>& B, bool after); // Source location for an instruction. Walks the inline chain via // DILocation::getInlinedAt() and joins entries with " -> " (innermost // first, then each outward call site). Matches the format used by @@ -274,12 +279,16 @@ class SQTTInstrumentPass : public llvm::PassInfoMixin uint32_t compactFuncIDs(llvm::Module& M); void rewriteMarkerIDs(llvm::Function& F, const std::map& IDMap, GfxGen gen); bool applyShaderClockPacking(llvm::Function& F, GfxGen gen); + // Lower full traces after all marker transformations to explicit asm with + // the required M0 hazard spacing. + bool lowerFullTracesWithM0Nop(llvm::Function& F); void removeFuncMarkersFromModule(llvm::Module& M, uint32_t id); void removeAdjacentBarriers(llvm::CallInst* CI); // ----------------------------------------------------------------- // -O0 fallback: direct function instrumentation // ----------------------------------------------------------------- + static bool hasMustTailCall(const llvm::Function& F); bool instrumentFunctionDirect(llvm::Function& F, GfxGen gen); // ----------------------------------------------------------------- diff --git a/projects/rocprof-trace-decoder/markers/SQTTTarget.h b/projects/rocprof-trace-decoder/markers/SQTTTarget.h index a2f6a5e08f5a..5d467cbd0a1f 100644 --- a/projects/rocprof-trace-decoder/markers/SQTTTarget.h +++ b/projects/rocprof-trace-decoder/markers/SQTTTarget.h @@ -55,6 +55,20 @@ inline GfxGen getGfxGen(const llvm::Function& F) return GfxGen::Unknown; } +// gfx1200 and gfx1201 use HW_REG_SHADER_CYCLES_LO through s_getreg; gfx1250 +// and subsequent targets have a dedicated shader-cycle instruction. +inline bool hasShaderCyclesU64(const llvm::Function& F) +{ + llvm::Attribute A = F.getFnAttribute("target-cpu"); + if (!A.isValid()) return false; + llvm::StringRef CPU = A.getValueAsString(); + if (!CPU.consume_front("gfx")) return false; + size_t digits = 0; + while (digits < CPU.size() && CPU[digits] >= '0' && CPU[digits] <= '9') ++digits; + unsigned target = 0; + return digits != 0 && !CPU.take_front(digits).getAsInteger(10, target) && target > 1201; +} + // Does this GfxGen support s_ttracedata_imm? inline bool supportsImmTrace(GfxGen gen) { @@ -82,8 +96,8 @@ inline HwRegEncodings getHwRegEncodings(GfxGen gen) inline unsigned getShaderClockBits(const SQTTConfig& config, GfxGen gen) { - if (config.ShaderClockBits != SQTTConfig::AutoShaderClockBits) return config.ShaderClockBits; - return gen == GfxGen::GFX12 ? 12 : 0; + (void) gen; + return config.ShaderClockBits; } inline bool usesShaderClockPacking(const SQTTConfig& config, GfxGen gen) @@ -137,6 +151,9 @@ inline unsigned computeFunctionSize(const llvm::Function& F, CostMode mode) { for (auto& I : BB) { + // Pass-owned marker calls must not make a function appear large + // enough to retain the instrumentation that introduced them. + if (I.getMetadata("sqtt.marker_header") || I.getMetadata("sqtt.raw_payload")) continue; if (mode == CostMode::WeightedCost) total += instructionCost(I); else diff --git a/projects/rocprof-trace-decoder/python/att_tool.py b/projects/rocprof-trace-decoder/python/att_tool.py index 40f524745d2a..cdd156810cdd 100644 --- a/projects/rocprof-trace-decoder/python/att_tool.py +++ b/projects/rocprof-trace-decoder/python/att_tool.py @@ -40,6 +40,7 @@ def generate_outputs_from_files( lib_path: str | Path | None = None, formats: str = "json,csv", base_name: str | None = None, + decode_markers: bool = True, ) -> list[Path]: inputs = _discover_inputs(files) if not inputs.att_files and not inputs.code_objects and inputs.code_json is None: @@ -91,6 +92,7 @@ def generate_outputs_from_files( formats=formats, base_name=output_base_name, on_warning=_print_warning, + decode_markers=decode_markers, ) if snapshot_source_dir: for output in outputs: @@ -214,6 +216,13 @@ def build_argparser() -> argparse.ArgumentParser: "--base-name", help="Base name for default ui_output_ and stats_<...>.csv output names.", ) + parser.add_argument( + "--no-decode-markers", + action="store_false", + dest="decode_markers", + default=True, + help="Leave packed marker IDs and JSON shaderdata timestamps unchanged.", + ) return parser @@ -225,6 +234,7 @@ def main(argv: list[str] | None = None) -> int: lib_path=args.lib, formats=args.formats, base_name=args.base_name, + decode_markers=args.decode_markers, ) for path in outputs: print(path) diff --git a/projects/rocprof-trace-decoder/python/rocprof_trace_decoder/att.py b/projects/rocprof-trace-decoder/python/rocprof_trace_decoder/att.py index c46d1110fc7a..47920d5fc3dc 100644 --- a/projects/rocprof-trace-decoder/python/rocprof_trace_decoder/att.py +++ b/projects/rocprof-trace-decoder/python/rocprof_trace_decoder/att.py @@ -9,6 +9,7 @@ from .bindings import Decoder from .code_index import CodeIndex from .codegen import CodeObject, generate_code_artifacts +from .records import ShaderDataFlags from .rcv import ( DEFAULT_RUN_BASE_NAME, RcvOutputWriter, @@ -45,6 +46,7 @@ def generate_att_outputs( formats: str | Iterable[str] = "json,csv", base_name: str = DEFAULT_RUN_BASE_NAME, on_warning: Callable[[Path, str], None] | None = None, + decode_markers: bool = True, ) -> list[Path]: output_formats = normalize_output_formats(formats) trace_list = [ @@ -92,6 +94,12 @@ def generate_att_outputs( ) writer = RcvOutputWriter(run_dir, run_code_index, formats=output_formats) + decode_json = ( + decode_markers + and "json" in output_formats + and bool(run_code_index.document.get("sqtt_funcmap_layout")) + ) + decoded = [] if decode_json else None with Decoder(lib_path) as decoder: for trace in run_traces: records = decoder.parse_file( @@ -101,7 +109,15 @@ def generate_att_outputs( if on_warning is not None: for info in records.info: on_warning(trace.path, decoder.info_string(info)) - writer.add_shader_records(trace.shader_engine, records) + if decoded is None: + writer.add_shader_records(trace.shader_engine, records) + else: + decoded.append((trace.shader_engine, records)) + + if decoded is not None: + _correct_marker_timestamps(decoded, run_code_index.document) + for shader_engine, records in decoded: + writer.add_shader_records(shader_engine, records) writer.finish() write_code_json(run_dir, run_code_index) @@ -114,6 +130,135 @@ def generate_att_outputs( return output_dirs +def _correct_marker_timestamps(decoded: list[tuple[int, object]], document: dict) -> None: + """Correct packed marker headers in-place, leaving legacy streams alone.""" + layouts: dict[int, tuple[int, int]] = {} + for row in document.get("sqtt_funcmap_layout", []): + try: + codeobj, bits, shift = map(int, row[:3]) + except (TypeError, ValueError): + continue + if codeobj >= 0 and 0 < bits <= 29 and 0 <= shift and bits + shift <= 32: + layouts[codeobj] = (bits, shift) + if not layouts: + return + + marker_ids: dict[int, set[int]] = defaultdict(set) + for row in document.get("sqtt_funcmap", []): + try: + codeobj, marker_id, kind = int(row[0]), int(row[1]), str(row[2]) + except (IndexError, TypeError, ValueError): + continue + if kind.lower() not in {"k", "kernel"} and marker_id >= 0: + marker_ids[codeobj].add(marker_id) + + payloads: dict[tuple[int, int], int] = {} + for row in document.get("sqtt_funcmap_payloads", []): + try: + codeobj, marker_id, count = map(int, row[:3]) + except (TypeError, ValueError): + continue + if marker_id in marker_ids.get(codeobj, ()) and count > 0: + payloads[codeobj, marker_id] = count + payload_codeobjs = {codeobj for codeobj, _marker_id in payloads} + + shaderdata: dict[tuple[int, int, int, int], list[object]] = defaultdict(list) + occupancy: dict[tuple[int, int, int, int], list[object]] = defaultdict(list) + for se, records in decoded: + for record in records.shaderdata: + shaderdata[se, record.cu, record.simd, record.wave_id].append(record) + for record in records.occupancy: + occupancy[se, record.cu, record.simd, record.wave_id].append(record) + + headers: dict[tuple[int, int, int, int, int], list[tuple[object, int, int]]] = defaultdict(list) + for location, stream in shaderdata.items(): + stream.sort(key=lambda record: record.time) + intervals = _active_codeobj_intervals(occupancy.get(location, ())) + interval = 0 + payload_remaining = 0 + for record in stream: + if payload_remaining: + payload_remaining -= 1 + continue + if record.flags & ((1 << ShaderDataFlags.IMM) | (1 << ShaderDataFlags.PRIV)): + continue + + while interval < len(intervals) and intervals[interval][1] < record.time: + interval += 1 + while interval + 1 < len(intervals) and intervals[interval + 1][0] <= record.time: + interval += 1 + codeobj = ( + intervals[interval][2] + if interval < len(intervals) and intervals[interval][0] <= record.time + else None + ) + layout = layouts.get(codeobj) + known_ids = marker_ids.get(codeobj, ()) + if layout is None: + continue + + bits, shift = layout + raw = int(record.value) & 0xFFFFFFFF + marker_id = (raw >> 2) & ((1 << (30 - bits)) - 1) + if marker_id not in known_ids and not (marker_id == 0 and (raw & 3) == 1): + continue + + record.value = (marker_id << 2) | (raw & 3) + if marker_id in known_ids: + payload_remaining = payloads.get((codeobj, marker_id), 0) + # New producers forbid clock packing with payload protocols. Keep + # old packed traces readable by normalizing their headers only. + if codeobj in payload_codeobjs: + continue + sampled = ((raw >> (32 - bits)) & ((1 << bits) - 1)) << shift + clock_domain = (*location[:3], bits, shift) + headers[clock_domain].append((record, int(record.time), sampled)) + + changed = False + for clock_domain, domain_headers in headers.items(): + bits, shift = clock_domain[-2:] + window_size = 1 << (bits + shift) + window_mask = window_size - 1 + half_window = window_size // 2 + bucket_size = 1 << shift + # A constant field cannot establish a relative issue-time delta. + if len({sampled for _record, _time, sampled in domain_headers}) == 1: + continue + # This is an arbitrary coordinate origin, not a time/minimum anchor. + # Comparing deltas cancels the fixed phase between the two clocks. + _origin, origin_time, origin_clock = domain_headers[0] + delays = [ + (((time - origin_time) - (clock - origin_clock) + half_window) & window_mask) - half_window + for _record, time, clock in domain_headers + ] + minimum = min(delays) + for (record, _time, _clock), delay in zip(domain_headers, delays): + correction = max(0, delay - minimum - (bucket_size - 1)) + if correction: + record.time -= correction + changed = True + # Payload order is semantic. Old traces may mix a no-clock payload code + # object with a packed one, so do not split their blocks by re-sorting. + if changed and not payloads: + for _se, records in decoded: + records.shaderdata.sort(key=lambda record: record.time) + + +def _active_codeobj_intervals(occupancy: Iterable[object]) -> list[tuple[int, int, int]]: + """Associate a wave slot with its last launch until its next launch.""" + events = sorted((record for record in occupancy if record.start), key=lambda record: record.time) + intervals: list[tuple[int, int, int]] = [] + begin = 0 + codeobj: int | None = None + for record in events: + if codeobj is not None: + intervals.append((begin, record.time, codeobj)) + begin, codeobj = record.time, record.pc.code_object_id + if codeobj is not None: + intervals.append((begin, (1 << 63) - 1, codeobj)) + return intervals + + def _group_traces(traces: Iterable[AttTrace]) -> dict[int, list[AttTrace]]: grouped: dict[int, list[AttTrace]] = defaultdict(list) for trace in traces: diff --git a/projects/rocprof-trace-decoder/python/rocprof_trace_decoder/codegen.py b/projects/rocprof-trace-decoder/python/rocprof_trace_decoder/codegen.py index 70cb50e51d0a..c27ec0a4f701 100644 --- a/projects/rocprof-trace-decoder/python/rocprof_trace_decoder/codegen.py +++ b/projects/rocprof-trace-decoder/python/rocprof_trace_decoder/codegen.py @@ -55,6 +55,22 @@ class CodeArtifacts: source_paths: tuple[Path, ...] +@dataclass(frozen=True) +class _SqttFuncmapEntry: + marker_id: int + kind: str + name: str + source_loc: str + + +@dataclass(frozen=True) +class _SqttFuncmap: + entries: tuple[_SqttFuncmapEntry, ...] + shader_clock_bits: int + shader_clock_shift: int + extra_payload_counts: dict[int, int] + + def _find_tool(name: str) -> str: candidates: list[Path] = [] for root in _rocm_roots(): @@ -107,6 +123,134 @@ def is_elf(path: str | Path) -> bool: return False +# --------------------------------------------------------------------------- +# SQTT marker metadata +# --------------------------------------------------------------------------- + + +def _parse_u32(value: str) -> int | None: + if not value or any(char < "0" or char > "9" for char in value): + return None + parsed = int(value) + return parsed if parsed <= 0xFFFFFFFF else None + + +def _split_funcmap_name(value: str) -> tuple[str, str]: + name, sep, source_loc = value.partition("@") + return (name, source_loc if sep else "") + + +def _parse_sqtt_funcmap(data: bytes) -> _SqttFuncmap: + """Parse the subset of ``.sqtt_funcmap`` needed by code.json. + + The output intentionally mirrors the C++ funcmap reader, but malformed + rows are ignored here: generating code metadata must remain best-effort + for ordinary code objects and older marker producers. + """ + + entries: list[_SqttFuncmapEntry] = [] + extra_payload_counts: dict[int, int] = {} + shader_clock_bits = 0 + shader_clock_shift = 0 + + for raw_line in re.split(r"[\n\x00]", data.decode("utf-8", errors="replace")): + line = raw_line.rstrip("\r \t") + if len(line) < 2 or line[1] != ":": + continue + + kind = line[0] + payload = line[2:] + if kind == "M": + bits = 0 + shift = 0 + saw_bits = False + saw_shift = False + valid = True + for attribute in payload.split(";"): + key, separator, value = attribute.partition("=") + if not separator: + continue + parsed = _parse_u32(value) + if key == "shader_clock_bits": + if parsed is None: + valid = False + else: + bits = parsed + saw_bits = True + elif key == "shader_clock_shift": + if parsed is None: + valid = False + else: + shift = parsed + saw_shift = True + + if valid and saw_bits and (bits == 0 or saw_shift) and bits <= 29 and ( + bits == 0 or (shift < 32 and bits <= 32 - shift) + ): + shader_clock_bits = bits + shader_clock_shift = shift if bits else 0 + continue + + if kind == "R": + marker_id_text, separator, attributes = payload.partition(":") + marker_id = _parse_u32(marker_id_text) if separator else None + if marker_id is None: + continue + for attribute in attributes.split(";"): + key, equals, value = attribute.partition("=") + if key != "extra_payload_count" or not equals: + continue + count = _parse_u32(value) + if count is not None: + extra_payload_counts[marker_id] = count + continue + + if kind == "K": + name, source_loc = _split_funcmap_name(payload) + if name: + entries.append(_SqttFuncmapEntry(0, kind, name, source_loc)) + continue + + if kind not in {"F", "U", "P"}: + continue + + marker_id_text, separator, name_and_source = payload.partition(":") + marker_id = _parse_u32(marker_id_text) if separator else None + if marker_id is None: + continue + name, source_loc = _split_funcmap_name(name_and_source) + if name: + entries.append(_SqttFuncmapEntry(marker_id, kind, name, source_loc)) + + return _SqttFuncmap( + entries=tuple(entries), + shader_clock_bits=shader_clock_bits, + shader_clock_shift=shader_clock_shift, + extra_payload_counts=extra_payload_counts, + ) + + +def _read_sqtt_funcmap(elf_path: str) -> _SqttFuncmap | None: + """Return embedded marker metadata, or ``None`` when no section exists.""" + + from elftools.common import exceptions as elftools_exceptions + from elftools.elf.elffile import ELFFile + + try: + with Path(elf_path).open("rb") as file: + section = ELFFile(file).get_section_by_name(".sqtt_funcmap") + if section is None: + return None + data = section.data() + except (OSError, elftools_exceptions.ELFError, ValueError): + return None + except Exception: + LOGGER.debug("Unexpected error reading .sqtt_funcmap from %s", elf_path, exc_info=True) + return None + + return _parse_sqtt_funcmap(data) if data else None + + # --------------------------------------------------------------------------- # DWARF helpers # --------------------------------------------------------------------------- @@ -569,6 +713,9 @@ def generate_code_artifacts(code_objects: Iterable[CodeObject]) -> CodeArtifacts rows: list[list] = [] comments: list[str] = [] kernels_out: list[dict] = [] + sqtt_funcmap_rows: list[list[object]] = [] + sqtt_funcmap_layouts: list[list[int]] = [] + sqtt_funcmap_payloads: list[list[int]] = [] line_no = 1 for code_object in normalized: @@ -577,6 +724,27 @@ def generate_code_artifacts(code_objects: Iterable[CodeObject]) -> CodeArtifacts range_map = build_address_ranges(elf_path) insts = parse_disassembly(elf_path) kernels, debug_labels = read_symbol_labels(elf_path) + funcmap = _read_sqtt_funcmap(elf_path) + + if funcmap is not None: + # Keep the existing sqtt_funcmap row shape used by rocprofv3 + # consumers. The two new side tables below are additive, so older + # readers can ignore them without losing their original schema. + if funcmap.shader_clock_bits: + sqtt_funcmap_layouts.append( + [codeobj_id, funcmap.shader_clock_bits, funcmap.shader_clock_shift] + ) + symbol_vaddrs: dict[str, int] = {} + for address, (name, _demangled) in sorted(kernels.items()): + symbol_vaddrs.setdefault(name, address) + for entry in funcmap.entries: + vaddr = symbol_vaddrs.get(entry.name, 0) if entry.kind in {"F", "K"} else 0 + sqtt_funcmap_rows.append( + [codeobj_id, entry.marker_id, entry.kind, entry.name, entry.source_loc, vaddr] + ) + payload_count = funcmap.extra_payload_counts.get(entry.marker_id, 0) + if entry.kind != "K" and payload_count: + sqtt_funcmap_payloads.append([codeobj_id, entry.marker_id, payload_count]) for vaddr, inst in insts: # Kernel and debug labels are emitted as code rows immediately @@ -639,6 +807,12 @@ def generate_code_artifacts(code_objects: Iterable[CodeObject]) -> CodeArtifacts "kernels": kernels_out, "hsaco": [str(code_object.path) for code_object in normalized], } + if sqtt_funcmap_rows: + code_doc["sqtt_funcmap"] = sqtt_funcmap_rows + if sqtt_funcmap_layouts: + code_doc["sqtt_funcmap_layout"] = sqtt_funcmap_layouts + if sqtt_funcmap_payloads: + code_doc["sqtt_funcmap_payloads"] = sqtt_funcmap_payloads source_paths = tuple(Path(path) for path in extract_paths_from_comments(comments)) return CodeArtifacts(CodeIndex.from_document(code_doc), source_paths) diff --git a/projects/rocprof-trace-decoder/scripts/sqtt_data.py b/projects/rocprof-trace-decoder/scripts/sqtt_data.py index 751cbdaba3be..0265cf8f0173 100644 --- a/projects/rocprof-trace-decoder/scripts/sqtt_data.py +++ b/projects/rocprof-trace-decoder/scripts/sqtt_data.py @@ -139,7 +139,6 @@ def extra_payload_count(self, marker_id: int) -> int: """Returns extra records after this marker header; default is zero.""" return self.extra_payload_counts.get(marker_id, 0) - def merge_funcmaps(funcmaps: Iterable[FuncMap]) -> FuncMap: """Merge funcmaps while preserving all metadata used for marker decoding.""" merged = FuncMap() @@ -203,18 +202,24 @@ def parse_funcmap(raw: str) -> FuncMap: except ValueError: pass elif prefix == "M": + values = {} + valid = True for item in rest.split(";"): key, sep, value = item.partition("=") - if not sep: - continue - try: - parsed = int(value) - except ValueError: + key = key.strip() + if key not in ("shader_clock_bits", "shader_clock_shift"): continue - if key.strip() == "shader_clock_bits": - fm.shader_clock_bits = parsed - elif key.strip() == "shader_clock_shift": - fm.shader_clock_shift = parsed + if not sep or not value.strip().isdigit(): + valid = False + break + values[key] = int(value) + bits = values.get("shader_clock_bits") + shift = values.get("shader_clock_shift") + if valid and bits is not None and ( + bits == 0 or (shift is not None and bits <= 29 and shift < 32 and bits <= 32 - shift) + ): + fm.shader_clock_bits = bits + fm.shader_clock_shift = shift or 0 elif prefix == "K": name, loc = _split_source_loc(rest) fm.kernels.append(name) @@ -315,19 +320,25 @@ class ShaderRecord: simd: int wave_id: int flags: int + shader_engine: int = 0 def decode_marker(value: int, funcmap: FuncMap | None = None) -> tuple[int, bool, bool]: - """Decode marker value -> (id, enter, exit_prev).""" - exit_prev = bool(value & 0x1) - enter = bool(value & 0x2) - clock_bits = funcmap.shader_clock_bits if funcmap else 0 - if clock_bits: - id_bits = max(0, 30 - clock_bits) - marker_id = (value >> 2) & ((1 << id_bits) - 1 if id_bits else 0) - else: - marker_id = value >> 2 - return marker_id, enter, exit_prev + """Decode known packed headers/bare exits; keep numeric markers legacy.""" + raw = value & 0xFFFFFFFF + exit_prev = bool(raw & 0x1) + enter = bool(raw & 0x2) + legacy_id = raw >> 2 + bits = funcmap.shader_clock_bits if funcmap else 0 + shift = funcmap.shader_clock_shift if funcmap else 0 + if not (funcmap and 0 < bits <= 29 and 0 <= shift < 32 and bits <= 32 - shift): + return legacy_id, enter, exit_prev + + id_bits = 30 - bits + packed_id = legacy_id & ((1 << id_bits) - 1) + if packed_id in funcmap.markers or (packed_id == 0 and exit_prev and not enter): + return packed_id, enter, exit_prev + return legacy_id, enter, exit_prev def _load_shaderdata_from_dir(trace_dir: str) -> list[ShaderRecord]: @@ -335,23 +346,28 @@ def _load_shaderdata_from_dir(trace_dir: str) -> list[ShaderRecord]: records = [] filenames_path = os.path.join(trace_dir, "filenames.json") - shaderdata_files = [] + shaderdata_files: list[tuple[int, str]] = [] if os.path.exists(filenames_path): with open(filenames_path) as f: meta = json.load(f) sd_filenames = meta.get("shaderdata_filenames", {}) for se_id, file_list in sd_filenames.items(): + try: + shader_engine = int(se_id) + except (TypeError, ValueError): + continue for entry in file_list: fname = entry[0] if isinstance(entry, list) else entry - shaderdata_files.append(os.path.join(trace_dir, fname)) + shaderdata_files.append((shader_engine, os.path.join(trace_dir, fname))) else: - shaderdata_files = sorted( + paths = sorted( glob(os.path.join(trace_dir, "shaderdata_*.json")), key=lambda p: [int(x) for x in os.path.splitext(os.path.basename(p))[0].split("_")[1:]], ) + shaderdata_files = [(int(os.path.basename(path).split("_")[1]), path) for path in paths] - for path in shaderdata_files: + for shader_engine, path in shaderdata_files: if not os.path.exists(path): continue with open(path) as f: @@ -361,6 +377,7 @@ def _load_shaderdata_from_dir(trace_dir: str) -> list[ShaderRecord]: time=rec[0], value=rec[1], cu=rec[2], simd=rec[3], wave_id=rec[4], flags=rec[5], + shader_engine=shader_engine, )) return records @@ -391,7 +408,7 @@ def load_occupancy(trace_dir: str) -> tuple[dict[str, str], dict[tuple, list[Wav Returns: dispatches: dict mapping dispatch_id (str) -> kernel name/description - wave_spans: dict mapping (cu, simd, wave_slot) -> sorted list of WaveSpan + wave_spans: dict mapping (shader_engine, cu, simd, wave_slot) -> sorted list of WaveSpan """ path = os.path.join(trace_dir, "occupancy.json") if not os.path.exists(path): @@ -404,13 +421,17 @@ def load_occupancy(trace_dir: str) -> tuple[dict[str, str], dict[tuple, list[Wav events: dict[tuple, list[tuple]] = defaultdict(list) for key, val in data.items(): - if key in ("dispatches", "version") or not isinstance(val, list): + try: + shader_engine = int(key) + except (TypeError, ValueError): + continue + if not isinstance(val, list): continue for rec in val: - if len(rec) < 6: + if not isinstance(rec, list) or len(rec) < 6: continue time, cu, simd, wslot, is_launch, dispatch_id = rec[:6] - events[(cu, simd, wslot)].append((time, is_launch, dispatch_id)) + events[(shader_engine, cu, simd, wslot)].append((time, is_launch, dispatch_id)) wave_spans: dict[tuple, list[WaveSpan]] = {} for wave_key, evts in events.items(): @@ -443,7 +464,9 @@ def find_wave_span_at(spans: list[WaveSpan], time: int) -> Optional[WaveSpan]: Trace data records can arrive after a dispatch retires (the wave is still draining), so we match to the most recent span whose launch_time <= time rather than requiring strict [launch, retire] - containment. + containment. Clock-corrected startup markers can also land just before + the first occupancy launch; with no earlier wave in that slot, they + belong to that first span. """ lo, hi = 0, len(spans) - 1 result = None @@ -455,7 +478,27 @@ def find_wave_span_at(spans: list[WaveSpan], time: int) -> Optional[WaveSpan]: lo = mid + 1 else: hi = mid - 1 - return result + return result if result is not None or not spans else spans[0] + + +def make_funcmap_resolver( + per_co: dict[int, FuncMap], + dispatches: dict[str, str], + kernel_to_co: dict[str, int], + wave_spans: dict[tuple, list[WaveSpan]], + fallback: FuncMap, +): + """Resolve a late record to its latest launch until that slot relaunches.""" + def resolve(record: ShaderRecord) -> FuncMap: + span = find_wave_span_at( + wave_spans.get((record.shader_engine, record.cu, record.simd, record.wave_id), []), record.time + ) + if span is None: + return fallback + codeobj = kernel_to_co.get(dispatches.get(str(span.dispatch_id), "")) + return per_co.get(codeobj, fallback) + + return resolve def merge_folded(all_folded: list[dict[str, tuple[int, int]]]) -> dict[str, tuple[int, int]]: @@ -490,6 +533,7 @@ class AddressTrace: addresses: list[int] # active lanes only; 64-bit for memory, 32-bit for LDS marker_id: int = 0 # unique per-op marker ID from funcmap source_loc: str = "" # "file.hip:42" if compiled with -g + shader_engine: int = 0 def active_lane_count(self) -> int: """Return count of active lanes (same as len(addresses)).""" @@ -632,6 +676,7 @@ def parse_address_block( addresses=addresses, marker_id=marker_id, source_loc=source_loc, + shader_engine=header_rec.shader_engine, ) return trace, i @@ -649,8 +694,8 @@ def _parse_buffer_block( """Parse buffer component block after exec mask has been read. Protocol: rsrc_lo, rsrc_hi, soffset, then per-lane voffset. - For struct buffers: additionally per-lane vindex. - Reconstructs full addresses: base + soffset + voffset[lane]. + Struct-buffer records contain vindex; consume them but do not guess an + address without a verified descriptor-stride decoder. """ is_struct = _is_struct_buffer(kind) @@ -672,15 +717,10 @@ def _parse_buffer_block( voffsets.append(records[i].value & 0xFFFFFFFF) i += 1 - # For struct buffers: per-lane vindex - vindices = None if is_struct: if i + wave_size > len(records): return None, len(records) - vindices = [] - for lane in range(wave_size): - vindices.append(records[i].value & 0xFFFFFFFF) - i += 1 + return None, i + wave_size # Reconstruct addresses for active lanes # Base address from rsrc descriptor: bits [47:0] @@ -702,6 +742,7 @@ def _parse_buffer_block( addresses=addresses, marker_id=marker_id, source_loc=source_loc, + shader_engine=header_rec.shader_engine, ) return trace, i @@ -709,13 +750,13 @@ def _parse_buffer_block( def preprocess_records( records: list[ShaderRecord], funcmap: FuncMap, + resolve_funcmap=None, ) -> tuple[list[ShaderRecord], list[AddressTrace]]: """Extract address trace blocks from the record stream. Address trace blocks span multiple consecutive s_ttracedata records (header + exec + addresses), but the raw shaderdata stream interleaves - records from all active waves. We demultiplex by (cu, simd, wave_id) - so each wave's records are contiguous before parsing address blocks. + records from all active waves. We demultiplex by (SE, CU, SIMD, wave). Returns: markers: all non-address-trace records (for build_stacks) @@ -730,26 +771,32 @@ def preprocess_records( # # A record is an address-trace header if decoding its value gives # a marker_id that resolves to an addr_trace_* funcmap entry. - # All subsequent records from the same (cu, simd, wave_id) until + # All subsequent records from the same wave stream until # the next header belong to the same address trace block. # Group all records by wave identity, preserving time order - per_wave: dict[tuple[int, int, int], list[ShaderRecord]] = defaultdict(list) - wave_has_addr_trace: set[tuple[int, int, int]] = set() + per_wave: dict[tuple[int, int, int, int], list[ShaderRecord]] = defaultdict(list) + wave_has_addr_trace: set[tuple[int, int, int, int]] = set() for rec in records: - key = (rec.cu, rec.simd, rec.wave_id) + if rec.flags & 2: # PRIV: generated by the trap handler + continue + key = (rec.shader_engine, rec.cu, rec.simd, rec.wave_id) per_wave[key].append(rec) + def map_for(record: ShaderRecord) -> FuncMap: + return resolve_funcmap(record) if resolve_funcmap is not None else funcmap + for key, wave_records in per_wave.items(): i = 0 while i < len(wave_records): - marker_id, _enter, _exit_prev = decode_marker(wave_records[i].value, funcmap) - name, _mtype = funcmap.resolve(marker_id) + current_map = map_for(wave_records[i]) + marker_id, _enter, _exit_prev = decode_marker(wave_records[i].value, current_map) + name, _mtype = current_map.resolve(marker_id) if _match_addr_trace(name) is not None: wave_has_addr_trace.add(key) break - i += 1 + funcmap.extra_payload_count(marker_id) + i += 1 + current_map.extra_payload_count(marker_id) # Process each wave's records for key, wave_records in per_wave.items(): @@ -759,27 +806,29 @@ def preprocess_records( i = 0 while i < len(wave_records): rec = wave_records[i] - marker_id, _enter, _exit_prev = decode_marker(rec.value, funcmap) + current_map = map_for(rec) + marker_id, _enter, _exit_prev = decode_marker(rec.value, current_map) markers.append(rec) - i += 1 + funcmap.extra_payload_count(marker_id) + i += 1 + current_map.extra_payload_count(marker_id) continue # Parse this wave's contiguous record stream i = 0 while i < len(wave_records): rec = wave_records[i] - marker_id, enter, exit_prev = decode_marker(rec.value, funcmap) - name, mtype = funcmap.resolve(marker_id) + current_map = map_for(rec) + marker_id, enter, exit_prev = decode_marker(rec.value, current_map) + name, mtype = current_map.resolve(marker_id) if _match_addr_trace(name) is not None: trace, i = parse_address_block( - wave_records, i, rec, name, marker_id, funcmap.wave_size, - funcmap.source_locs.get(marker_id, "")) + wave_records, i, rec, name, marker_id, current_map.wave_size, + current_map.source_locs.get(marker_id, "")) if trace: addr_traces.append(trace) else: markers.append(rec) - i += 1 + funcmap.extra_payload_count(marker_id) + i += 1 + current_map.extra_payload_count(marker_id) # Re-sort markers by time since we processed per-wave markers.sort(key=lambda r: r.time) diff --git a/projects/rocprof-trace-decoder/scripts/sqtt_flamegraph.py b/projects/rocprof-trace-decoder/scripts/sqtt_flamegraph.py index ab8c734e1dba..62167514243b 100755 --- a/projects/rocprof-trace-decoder/scripts/sqtt_flamegraph.py +++ b/projects/rocprof-trace-decoder/scripts/sqtt_flamegraph.py @@ -86,6 +86,7 @@ load_occupancy, load_shaderdata, build_kernel_to_co, + make_funcmap_resolver, merge_folded, merge_funcmaps, preprocess_records, @@ -97,7 +98,7 @@ # --------------------------------------------------------------------------- # A unique wave identity -WaveKey = tuple[int, int, int] # (cu, simd, wave_id) +WaveKey = tuple[int, int, int, int, int] # (se, cu, simd, wave_id, instance) def build_stacks( @@ -124,7 +125,7 @@ def build_stacks( # Merged funcmap for fallback resolution. merged = merge_funcmaps(per_co.values()) - # Build sorted launch times per (cu, simd, wave_slot) from occupancy. + # Build sorted launch times per (SE, CU, SIMD, wave_slot) from occupancy. slot_launches: dict[tuple, list[int]] = defaultdict(list) for wk, spans in wave_spans.items(): for span in spans: @@ -132,9 +133,9 @@ def build_stacks( for wk in slot_launches: slot_launches[wk].sort() - def _wave_instance(cu: int, simd: int, wave_id: int, time: int) -> int: + def _wave_instance(shader_engine: int, cu: int, simd: int, wave_id: int, time: int) -> int: """Find which wave instance (launch index) a record belongs to.""" - launches = slot_launches.get((cu, simd, wave_id), []) + launches = slot_launches.get((shader_engine, cu, simd, wave_id), []) if not launches: return 0 lo, hi = 0, len(launches) - 1 @@ -148,17 +149,19 @@ def _wave_instance(cu: int, simd: int, wave_id: int, time: int) -> int: hi = mid - 1 return result - # Group records by (cu, simd, wave_id, instance) + # Group records by (SE, CU, SIMD, wave_id, instance) wave_records: dict[tuple, list[ShaderRecord]] = defaultdict(list) for rec in records: + if rec.flags & 2: + continue if filter_cu is not None and rec.cu != filter_cu: continue if filter_simd is not None and rec.simd != filter_simd: continue if filter_wave is not None and rec.wave_id != filter_wave: continue - inst = _wave_instance(rec.cu, rec.simd, rec.wave_id, rec.time) - key = (rec.cu, rec.simd, rec.wave_id, inst) + inst = _wave_instance(rec.shader_engine, rec.cu, rec.simd, rec.wave_id, rec.time) + key = (rec.shader_engine, rec.cu, rec.simd, rec.wave_id, inst) wave_records[key].append(rec) # (total_cycles, exec_count) per stack @@ -193,15 +196,10 @@ def _flush(prefix, stack, last_time, now): cur_dispatch_id: Optional[int] = None cur_funcmap: FuncMap = merged - base_wk = wave_key[:3] + base_wk = wave_key[:4] spans = wave_spans.get(base_wk, []) for rec in recs: - marker_id, enter, exit_prev = decode_marker(rec.value, cur_funcmap) - - if marker_id == 0 and not exit_prev: - continue - # Check dispatch context span = find_wave_span_at(spans, rec.time) dispatch_id = span.dispatch_id if span else None @@ -212,6 +210,10 @@ def _flush(prefix, stack, last_time, now): stack = [] cur_prefix, cur_funcmap = _dispatch_prefix(dispatch_id) + marker_id, enter, exit_prev = decode_marker(rec.value, cur_funcmap) + if marker_id == 0 and not exit_prev: + continue + name, mtype = cur_funcmap.resolve(marker_id) # Point markers (barriers, memory ops, user points) carry no @@ -568,9 +570,10 @@ def main(): if not v.startswith("0 /")]) # Filter out address trace records before building stacks - records, addr_traces = preprocess_records(records, merged) + resolver = make_funcmap_resolver(per_co, dispatches, kernel_to_co, wave_spans, merged) + records, addr_traces = preprocess_records(records, merged, resolver) - waves = set((r.cu, r.simd, r.wave_id) for r in records) + waves = set((record.shader_engine, record.cu, record.simd, record.wave_id) for record in records) total_records += len(records) all_waves.update(waves) diff --git a/projects/rocprof-trace-decoder/scripts/sqtt_memory_trace.py b/projects/rocprof-trace-decoder/scripts/sqtt_memory_trace.py index c665071c0814..54c1d76346dd 100644 --- a/projects/rocprof-trace-decoder/scripts/sqtt_memory_trace.py +++ b/projects/rocprof-trace-decoder/scripts/sqtt_memory_trace.py @@ -52,10 +52,12 @@ from sqtt_data import ( AddressTrace, + build_kernel_to_co, discover_base_dir, load_funcmaps, load_occupancy, load_shaderdata, + make_funcmap_resolver, merge_funcmaps, preprocess_records, ) @@ -70,6 +72,7 @@ def format_address_trace(trace: AddressTrace) -> dict: "kind": trace.kind, "marker_id": trace.marker_id, "time": trace.time, + "shader_engine": trace.shader_engine, "cu": trace.cu, "simd": trace.simd, "wave": trace.wave_id, @@ -101,17 +104,20 @@ def print_summary(traces: list[AddressTrace], file=sys.stderr): # Global stats total_addrs = sum(len(t.addresses) for t in traces) + shader_engines = set() cus = set() simds = set() by_kind: dict[str, int] = {} for t in traces: by_kind[t.kind] = by_kind.get(t.kind, 0) + 1 + shader_engines.add(t.shader_engine) cus.add(t.cu) simds.add(t.simd) print(f"\nAddress trace summary:", file=file) print(f" Total operations: {len(traces)}", file=file) print(f" Total addresses: {total_addrs}", file=file) + print(f" Shader engines: {sorted(shader_engines)}", file=file) print(f" CUs seen: {sorted(cus)}", file=file) print(f" SIMDs seen: {sorted(simds)}", file=file) print(f" By kind:", file=file) @@ -141,15 +147,15 @@ def print_summary(traces: list[AddressTrace], file=sys.stderr): mid_traces = by_mid[mid] print(f" marker_id={mid} ({len(mid_traces)} samples)", file=file) - # Group by wave (CU, SIMD, wave_id) - by_wave: dict[tuple[int, int, int], list[AddressTrace]] = defaultdict(list) + # Group by wave (SE, CU, SIMD, wave_id) + by_wave: dict[tuple[int, int, int, int], list[AddressTrace]] = defaultdict(list) for t in mid_traces: - by_wave[(t.cu, t.simd, t.wave_id)].append(t) + by_wave[(t.shader_engine, t.cu, t.simd, t.wave_id)].append(t) - for (cu, simd, wid) in sorted(by_wave.keys()): - wave_traces = by_wave[(cu, simd, wid)] + for (se, cu, simd, wid) in sorted(by_wave.keys()): + wave_traces = by_wave[(se, cu, simd, wid)] is_64bit = wave_traces[0].kind in ("load", "store") - print(f" wave CU={cu} SIMD={simd} wave={wid} " + print(f" wave SE={se} CU={cu} SIMD={simd} wave={wid} " f"({len(wave_traces)} samples)", file=file) for t in wave_traces: @@ -217,8 +223,10 @@ def main(): print(f"Found {len(trace_dirs)} trace dir(s), " f"{len(code_objects)} code object(s)", file=sys.stderr) - # Load funcmaps - per_co = load_funcmaps(code_objects, do_demangle=args.demangle) + per_co_raw = load_funcmaps(code_objects, do_demangle=False) + kernel_to_co = build_kernel_to_co(per_co_raw) + per_co = (load_funcmaps(code_objects, do_demangle=True) + if args.demangle else per_co_raw) # Build merged funcmap for preprocessing. merged = merge_funcmaps(per_co.values()) @@ -233,7 +241,10 @@ def main(): continue total_records += len(records) - _, addr_traces = preprocess_records(records, merged) + dispatches, wave_spans = load_occupancy(td) + resolver = make_funcmap_resolver( + per_co, dispatches, kernel_to_co, wave_spans, merged) + _, addr_traces = preprocess_records(records, merged, resolver) # Apply filters for t in addr_traces: diff --git a/projects/rocprof-trace-decoder/scripts/sqtt_perfetto.py b/projects/rocprof-trace-decoder/scripts/sqtt_perfetto.py index 851d34b4c9ac..818a921025a2 100644 --- a/projects/rocprof-trace-decoder/scripts/sqtt_perfetto.py +++ b/projects/rocprof-trace-decoder/scripts/sqtt_perfetto.py @@ -68,6 +68,7 @@ load_funcmaps, load_occupancy, load_shaderdata, + make_funcmap_resolver, merge_funcmaps, preprocess_records, ) @@ -77,13 +78,13 @@ # Wave grouping helpers # --------------------------------------------------------------------------- -WaveKey = tuple[int, int, int, int] # (cu, simd, wave_id, instance) +WaveKey = tuple[int, int, int, int, int] # (se, cu, simd, wave_id, instance) def _build_slot_launches( wave_spans: dict[tuple, list[WaveSpan]], ) -> dict[tuple, list[int]]: - """Sorted launch times per (cu, simd, wave_slot). Used to assign instance + """Sorted launch times per (SE, CU, SIMD, wave_slot). Used to assign instance indices to records that share a physical wave slot across relaunches.""" slot_launches: dict[tuple, list[int]] = defaultdict(list) for wk, spans in wave_spans.items(): @@ -95,9 +96,9 @@ def _build_slot_launches( def _wave_instance(slot_launches: dict[tuple, list[int]], - cu: int, simd: int, wave_id: int, time: int) -> int: + shader_engine: int, cu: int, simd: int, wave_id: int, time: int) -> int: """Index of the most recent launch at or before `time` for this slot.""" - launches = slot_launches.get((cu, simd, wave_id), []) + launches = slot_launches.get((shader_engine, cu, simd, wave_id), []) if not launches: return 0 lo, hi = 0, len(launches) - 1 @@ -119,32 +120,34 @@ def group_by_wave_instance( filter_simd: Optional[int] = None, filter_wave: Optional[int] = None, ) -> dict[WaveKey, list[ShaderRecord]]: - """Group records by (cu, simd, wave_id, instance), preserving time order.""" + """Group records by (SE, CU, SIMD, wave_id, instance), preserving time order.""" slot_launches = _build_slot_launches(wave_spans) grouped: dict[WaveKey, list[ShaderRecord]] = defaultdict(list) for rec in records: + if rec.flags & 2: + continue if filter_cu is not None and rec.cu != filter_cu: continue if filter_simd is not None and rec.simd != filter_simd: continue if filter_wave is not None and rec.wave_id != filter_wave: continue - inst = _wave_instance(slot_launches, rec.cu, rec.simd, + inst = _wave_instance(slot_launches, rec.shader_engine, rec.cu, rec.simd, rec.wave_id, rec.time) - grouped[(rec.cu, rec.simd, rec.wave_id, inst)].append(rec) + grouped[(rec.shader_engine, rec.cu, rec.simd, rec.wave_id, inst)].append(rec) for k in grouped: grouped[k].sort(key=lambda r: r.time) return grouped -def pack_tid(cu: int, simd: int, wave_id: int, instance: int) -> int: - """Pack a wave key into a 32-bit thread id. +def pack_tid(shader_engine: int, cu: int, simd: int, wave_id: int, instance: int) -> int: + """Pack a shader-engine-aware wave key into a 64-bit thread id. - Layout: (cu << 20) | (simd << 16) | (wave_id << 8) | instance. - Collision-free for cu<4096, simd<16, wave_id<256, instance<256 -- all - well above hardware limits. + The lower 32 bits retain the old CU/SIMD/wave/instance layout; shader + engine occupies the high bits so equal wave slots on two engines do not + share a Perfetto track. """ - return (cu << 20) | (simd << 16) | (wave_id << 8) | instance + return (shader_engine << 32) | (cu << 20) | (simd << 16) | (wave_id << 8) | instance # --------------------------------------------------------------------------- @@ -195,9 +198,9 @@ def emit_wave_events( """ events: list[dict] = [] unmatched = 0 - cu, simd, wave_id, instance = wave_key - tid = pack_tid(cu, simd, wave_id, instance) - base_wk = (cu, simd, wave_id) + shader_engine, cu, simd, wave_id, instance = wave_key + tid = pack_tid(shader_engine, cu, simd, wave_id, instance) + base_wk = wave_key[:4] spans = wave_spans.get(base_wk, []) stack: list[tuple[str, str]] = [] # (name, mtype) pairs we've pushed @@ -208,10 +211,6 @@ def _ts(cycles: int) -> float: return cycles * ts_scale for rec in wave_records: - marker_id, enter, exit_prev = decode_marker(rec.value, cur_funcmap) - if marker_id == 0 and not exit_prev: - continue # noop / reserved value - # Update dispatch context. When the dispatch changes mid-stream # we synthesize E events for anything still on the stack -- a # function entered under dispatch A doesn't outlive A. @@ -228,6 +227,10 @@ def _ts(cycles: int) -> float: cur_funcmap = _funcmap_for_dispatch( dispatch_id, dispatches, kernel_to_co, per_co, merged_fm) + marker_id, enter, exit_prev = decode_marker(rec.value, cur_funcmap) + if marker_id == 0 and not exit_prev: + continue # noop / reserved value + name, mtype = cur_funcmap.resolve(marker_id) # Point markers (barriers, memory ops, user points) are instantaneous @@ -274,7 +277,7 @@ def _ts(cycles: int) -> float: def build_metadata_events( pids_seen: set[int], - tids_seen: set[tuple[int, int, tuple[int, int, int, int]]], + tids_seen: set[tuple], dispatches: dict[str, str], do_demangle: bool, ) -> list[dict]: @@ -282,7 +285,7 @@ def build_metadata_events( each thread (wave instance). tids_seen is a set of (pid, tid, wave_key) triples so we can recover - the (cu, simd, wave_id, instance) tuple for the thread name. + the (SE, CU, SIMD, wave_id, instance) tuple for the thread name. """ metadata: list[dict] = [] @@ -299,8 +302,8 @@ def build_metadata_events( }) for pid, tid, wave_key in sorted(tids_seen): - cu, simd, wave_id, instance = wave_key - thread_name = f"CU{cu}/SIMD{simd}/W{wave_id}#{instance}" + shader_engine, cu, simd, wave_id, instance = wave_key + thread_name = f"SE{shader_engine}/CU{cu}/SIMD{simd}/W{wave_id}#{instance}" metadata.append({ "ph": "M", "name": "thread_name", "pid": pid, "tid": tid, "args": {"name": thread_name}, @@ -333,12 +336,12 @@ def export_trace_dir( records = load_shaderdata(trace_dir) if not records: return None - records, _addr = preprocess_records(records, merged_fm) + dispatches, wave_spans = load_occupancy(trace_dir) + resolver = make_funcmap_resolver(per_co, dispatches, kernel_to_co, wave_spans, merged_fm) + records, _addr = preprocess_records(records, merged_fm, resolver) if not records: return None - dispatches, wave_spans = load_occupancy(trace_dir) - grouped = group_by_wave_instance( records, wave_spans, filter_cu=filter_cu, filter_simd=filter_simd, filter_wave=filter_wave) @@ -352,7 +355,7 @@ def export_trace_dir( all_events: list[dict] = [] pids_seen: set[int] = set() - tids_seen: set[tuple[int, int, tuple[int, int, int, int]]] = set() + tids_seen: set[tuple] = set() total_unmatched = 0 for wave_key, recs in grouped.items(): @@ -361,8 +364,8 @@ def export_trace_dir( wave_spans, ts_scale) total_unmatched += unmatched all_events.extend(evs) - cu, simd, wave_id, instance = wave_key - tid = pack_tid(cu, simd, wave_id, instance) + shader_engine, cu, simd, wave_id, instance = wave_key + tid = pack_tid(shader_engine, cu, simd, wave_id, instance) for ev in evs: pids_seen.add(ev["pid"]) tids_seen.add((ev["pid"], tid, wave_key)) @@ -424,19 +427,19 @@ def _self_test() -> int: exit_prev = 0x1 point_barrier = (2 << 2) - # One wave (cu=0, simd=1, wave_id=3) running dispatch 7. + # One wave (SE=0, cu=0, simd=1, wave_id=3) running dispatch 7. recs = [ ShaderRecord(time=100, value=enter_compute, cu=0, simd=1, wave_id=3, flags=0), ShaderRecord(time=200, value=point_barrier, cu=0, simd=1, wave_id=3, flags=0), ShaderRecord(time=300, value=exit_prev, cu=0, simd=1, wave_id=3, flags=0), ] - wave_spans = {(0, 1, 3): [WaveSpan(launch_time=0, retire_time=1000, dispatch_id=7)]} + wave_spans = {(0, 0, 1, 3): [WaveSpan(launch_time=0, retire_time=1000, dispatch_id=7)]} dispatches = {"7": "compute_kernel"} grouped = group_by_wave_instance(recs, wave_spans) assert len(grouped) == 1, f"expected 1 wave group, got {len(grouped)}" wave_key = next(iter(grouped)) - assert wave_key == (0, 1, 3, 0), wave_key + assert wave_key == (0, 0, 1, 3, 0), wave_key events, unmatched = emit_wave_events( grouped[wave_key], wave_key, dispatches, kernel_to_co={}, @@ -450,7 +453,7 @@ def _self_test() -> int: assert events[1]["name"] == "barrier" and events[1]["cat"] == "sqtt.point" assert events[1]["s"] == "t" assert all(e["pid"] == 7 for e in events) - expected_tid = pack_tid(0, 1, 3, 0) + expected_tid = pack_tid(0, 0, 1, 3, 0) assert all(e["tid"] == expected_tid for e in events) # B/E pairing (per (pid, tid)) diff --git a/projects/rocprof-trace-decoder/test/CMakeLists.txt b/projects/rocprof-trace-decoder/test/CMakeLists.txt index b714fbf869bb..f4d6c0c08038 100644 --- a/projects/rocprof-trace-decoder/test/CMakeLists.txt +++ b/projects/rocprof-trace-decoder/test/CMakeLists.txt @@ -111,6 +111,8 @@ if(BUILD_INTEGRATION_TESTS) set(CSV_PY ${CMAKE_CURRENT_SOURCE_DIR}/csv_test.py) set(ATT_TOOL_PY ${CMAKE_CURRENT_SOURCE_DIR}/../python/att_tool.py) set(PYTHON_API_PY ${CMAKE_CURRENT_SOURCE_DIR}/python_api_test.py) + set(CODEGEN_TEST_PY ${CMAKE_CURRENT_SOURCE_DIR}/codegen_test.py) + set(ATT_TOOL_TEST_PY ${CMAKE_CURRENT_SOURCE_DIR}/att_tool_test.py) set(LIB_PATH ${CMAKE_BINARY_DIR}/lib/librocprof-trace-decoder${CMAKE_SHARED_LIBRARY_SUFFIX}) set(SANITIZED_LIB_PATH ${CMAKE_BINARY_DIR}/lib/librocprof-trace-decoder-sanitized${CMAKE_SHARED_LIBRARY_SUFFIX}) @@ -146,6 +148,27 @@ if(BUILD_INTEGRATION_TESTS) ENVIRONMENT "${PYTHON_API_TEST_ENV}") + add_test(NAME codegen-funcmap-metadata COMMAND ${Python3_EXECUTABLE} ${CODEGEN_TEST_PY}) + set_tests_properties( + codegen-funcmap-metadata + PROPERTIES TIMEOUT + 30 + LABELS + "integration-tests" + ENVIRONMENT + "PYTHONPATH=${CMAKE_CURRENT_SOURCE_DIR}/../python") + + add_test(NAME att-marker-clock-correction COMMAND ${Python3_EXECUTABLE} ${ATT_TOOL_TEST_PY}) + set_tests_properties( + att-marker-clock-correction + PROPERTIES + TIMEOUT + 30 + LABELS + "integration-tests" + ENVIRONMENT + "PYTHONPATH=${CMAKE_CURRENT_SOURCE_DIR}/../python") + if(DEFINED ROCPROF_TRACE_DECODER_PYTHON_INSTALL_DIR) configure_file(python_api_install_test.cmake.in python_api_install_test.cmake @ONLY) add_test( diff --git a/projects/rocprof-trace-decoder/test/att_tool_test.py b/projects/rocprof-trace-decoder/test/att_tool_test.py new file mode 100644 index 000000000000..76eb6267c18f --- /dev/null +++ b/projects/rocprof-trace-decoder/test/att_tool_test.py @@ -0,0 +1,404 @@ +#!/usr/bin/env python3 +"""Focused checks for packed-marker JSON normalization and timestamp correction.""" + +from __future__ import annotations + +import sys +import unittest +import xml.etree.ElementTree as ET +from pathlib import Path +from tempfile import TemporaryDirectory +from unittest.mock import patch + +ROOT = Path(__file__).resolve().parents[1] +sys.path[:0] = [str(ROOT / "python"), str(ROOT / "scripts")] + +from att_tool import build_argparser +from rocprof_trace_decoder.att import AttTrace, _correct_marker_timestamps, generate_att_outputs +from rocprof_trace_decoder.code_index import CodeIndex +from rocprof_trace_decoder.records import Occupancy, Pc, ShaderData, ShaderDataFlags, TraceRecords +from sqtt_data import FuncMap, ShaderRecord, WaveSpan, find_wave_span_at, preprocess_records +from sqtt_flamegraph import build_stacks, render_flamegraph_svg +from sqtt_perfetto import emit_wave_events + + +def _shaderdata(time: int, value: int, flags: int = 0, simd: int = 2) -> ShaderData: + return ShaderData(time, value, cu=1, simd=simd, wave_id=3, flags=flags, reserved=0) + + +def _occupancy(start: int, time: int, codeobj: int = 7, simd: int = 2) -> Occupancy: + return Occupancy(Pc(0, codeobj), time, 0, 1, simd, 3, start, 0, 0, 0, 0, 0) + + +class _Decoder: + def __init__(self, records: TraceRecords): + self.records = records + + def __enter__(self): + return self + + def __exit__(self, *_args): + return False + + def parse_file(self, *_args, **_kwargs): + return self.records + + +class _Writer: + records: TraceRecords | None = None + + def __init__(self, *_args, **_kwargs): + pass + + def add_shader_records(self, _se: int, records: TraceRecords): + type(self).records = records + + def finish(self): + pass + + +class MarkerTimestampCorrectionTest(unittest.TestCase): + def test_payload_headers_are_normalized_without_clock_correction(self): + first_value = (0x100 << 20) | (3 << 2) + second_value = (0x101 << 20) | (4 << 2) + first = _shaderdata(0x1000, first_value) + payload = _shaderdata(0x1020, 0xDEADBEEF) + second = _shaderdata(0x1040, second_value) + records = TraceRecords( + occupancy=[_occupancy(1, 0)], + shaderdata=[first, payload, second], + ) + document = { + "sqtt_funcmap": [[7, 3, "P", "first", "", 0], [7, 4, "P", "second", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + "sqtt_funcmap_payloads": [[7, 3, 1]], + } + + _correct_marker_timestamps([(0, records)], document) + + self.assertEqual([record.time for record in records.shaderdata], [0x1000, 0x1020, 0x1040]) + self.assertIs(records.shaderdata[0], first) + self.assertIs(records.shaderdata[1], payload) + self.assertIs(records.shaderdata[2], second) + self.assertEqual([record.value for record in records.shaderdata], [12, 0xDEADBEEF, 16]) + + def test_payload_records_stay_adjacent_when_another_code_object_is_corrected(self): + header = _shaderdata(0x1000, 3 << 2) + payload = _shaderdata(0x1020, 0xDEADBEEF) + first = _shaderdata(0x1040, (0x100 << 20) | (4 << 2)) + second = _shaderdata(0x1080, (0x107 << 20) | (4 << 2)) + records = TraceRecords( + occupancy=[_occupancy(1, 0, codeobj=7), _occupancy(1, 0x1030, codeobj=8)], + shaderdata=[header, payload, first, second], + ) + + _correct_marker_timestamps( + [(0, records)], + { + "sqtt_funcmap": [[7, 3, "P", "payload", "", 0], [8, 4, "P", "marker", "", 0]], + "sqtt_funcmap_layout": [[8, 12, 4]], + "sqtt_funcmap_payloads": [[7, 3, 1]], + }, + ) + + self.assertEqual([record.time for record in records.shaderdata], [0x1000, 0x1020, 0x101F, 0x1080]) + self.assertEqual([record.value for record in records.shaderdata], [12, 0xDEADBEEF, 16, 16]) + + def test_imm_record_is_not_a_packed_header(self): + record = _shaderdata(0x1010, 3 << 2, flags=1 << ShaderDataFlags.IMM) + records = TraceRecords(occupancy=[_occupancy(1, 0)], shaderdata=[record]) + + _correct_marker_timestamps( + [(0, records)], + {"sqtt_funcmap": [[7, 3, "P", "marker"]], "sqtt_funcmap_layout": [[7, 12, 4]]}, + ) + + self.assertEqual(record.time, 0x1010) + self.assertEqual(record.value, 3 << 2) + + def test_legacy_metadata_leaves_records_unchanged(self): + records = TraceRecords(shaderdata=[_shaderdata(20, 12), _shaderdata(10, 16)]) + original = list(records.shaderdata) + + _correct_marker_timestamps([(0, records)], {"sqtt_funcmap": [[7, 3, "P", "marker"]]}) + + self.assertEqual(records.shaderdata, original) + + def test_no_active_code_object_is_not_code_object_zero(self): + value = (0x100 << 20) | (3 << 2) + record = _shaderdata(0x1080, value) + records = TraceRecords(shaderdata=[record]) + + _correct_marker_timestamps( + [(0, records)], + {"sqtt_funcmap": [[0, 3, "P", "marker"]], "sqtt_funcmap_layout": [[0, 12, 4]]}, + ) + + self.assertEqual(record.time, 0x1080) + self.assertEqual(record.value, value) + + def test_exit_without_a_marker_row_is_decoded(self): + record = _shaderdata(0x1010, (0x100 << 20) | 1) + records = TraceRecords(occupancy=[_occupancy(1, 0)], shaderdata=[record]) + + _correct_marker_timestamps( + [(0, records)], + { + "sqtt_funcmap": [[7, 0, "K", "kernel", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + }, + ) + + self.assertEqual(record.value, 1) + + def test_unknown_packed_values_are_unchanged(self): + records = TraceRecords( + occupancy=[_occupancy(1, 0)], + shaderdata=[_shaderdata(20, 4 << 20), _shaderdata(10, 4 << 20)], + ) + original = list(records.shaderdata) + + _correct_marker_timestamps( + [(0, records)], + {"sqtt_funcmap": [[7, 3, "P", "marker"]], "sqtt_funcmap_layout": [[7, 12, 4]]}, + ) + + self.assertEqual(records.shaderdata, original) + + def test_correction_is_default(self): + early_value = (0x100 << 20) | (3 << 2) + late_value = (0x101 << 20) | (3 << 2) + records = TraceRecords( + occupancy=[_occupancy(1, 0)], + shaderdata=[_shaderdata(0x1080, late_value), _shaderdata(0x1010, early_value)], + ) + code_index = CodeIndex.from_document( + { + "sqtt_funcmap": [[7, 3, "P", "marker", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + } + ) + + with TemporaryDirectory() as directory: + with patch("rocprof_trace_decoder.att.Decoder", return_value=_Decoder(records)), \ + patch("rocprof_trace_decoder.att.RcvOutputWriter", _Writer), \ + patch("rocprof_trace_decoder.att.write_code_json"): + generate_att_outputs( + [AttTrace(Path(directory) / "trace.att", shader_engine=0)], + code_index=code_index, + output_dir=directory, + formats="json", + ) + + self.assertEqual([record.time for record in records.shaderdata], [0x1010, 0x102F]) + self.assertEqual([record.value for record in records.shaderdata], [12, 12]) + + def test_constant_clock_samples_are_normalized_without_correction(self): + late = _shaderdata(0x1080, 3 << 2) + early = _shaderdata(0x1010, 4 << 2) + records = TraceRecords( + occupancy=[_occupancy(1, 0)], shaderdata=[late, early] + ) + + _correct_marker_timestamps( + [(0, records)], + { + "sqtt_funcmap": [[7, 3, "P", "first", "", 0], [7, 4, "P", "second", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + }, + ) + + self.assertEqual([late.time, early.time], [0x1080, 0x1010]) + self.assertEqual([late.value, early.value], [12, 16]) + + def test_correction_does_not_share_the_reference_between_simds(self): + slow_first = _shaderdata(0x1000, (0x100 << 20) | (3 << 2), simd=2) + slow_second = _shaderdata(0x1080, (0x101 << 20) | (3 << 2), simd=2) + fast_first = _shaderdata(0x3000, (0x100 << 20) | (3 << 2), simd=3) + fast_second = _shaderdata(0x3080, (0x101 << 20) | (3 << 2), simd=3) + records = TraceRecords( + occupancy=[_occupancy(1, 0, simd=2), _occupancy(1, 0, simd=3)], + shaderdata=[slow_first, slow_second, fast_first, fast_second], + ) + + _correct_marker_timestamps( + [(0, records)], + { + "sqtt_funcmap": [[7, 3, "P", "marker", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + }, + ) + + self.assertEqual([record.time for record in records.shaderdata], [0x1000, 0x101F, 0x3000, 0x301F]) + self.assertEqual([record.value for record in records.shaderdata], [12, 12, 12, 12]) + + def test_correction_cancels_unknown_clock_phase(self): + # The fixed inter-clock phase crosses the sampled window. The + # correction must use clock/time deltas, not absolute values. + marker_id = 3 << 2 + delayed = _shaderdata(0x101D0, marker_id) + reference = _shaderdata(0x10DFC, (0x100 << 20) | marker_id) + records = TraceRecords( + occupancy=[_occupancy(1, 0)], shaderdata=[delayed, reference]) + + _correct_marker_timestamps( + [(0, records)], + { + "sqtt_funcmap": [[7, 3, "P", "marker", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + }, + ) + + self.assertEqual([record.time for record in records.shaderdata], [0xFE0B, 0x10DFC]) + self.assertEqual([record.value for record in records.shaderdata], [marker_id, marker_id]) + + def test_packed_bare_exit_after_retirement_is_normalized_and_corrected(self): + before_retirement = _shaderdata(0x1010, (0x100 << 20) | (3 << 2)) + after_retirement = _shaderdata(0x1080, (0x101 << 20) | 1) + records = TraceRecords( + occupancy=[_occupancy(1, 0), _occupancy(0, 0x1020)], + shaderdata=[before_retirement, after_retirement], + ) + + _correct_marker_timestamps( + [(0, records)], + { + "sqtt_funcmap": [[7, 3, "P", "marker", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + }, + ) + + self.assertEqual(before_retirement.time, 0x1010) + self.assertEqual(after_retirement.time, 0x102F) + self.assertEqual([record.value for record in records.shaderdata], [12, 1]) + + def test_later_launch_supersedes_a_retired_wave_slot(self): + value = (0x80 << 21) | (4 << 2) + first = _shaderdata(0x1020, value) + second = _shaderdata(0x1080, value) + records = TraceRecords( + occupancy=[ + _occupancy(1, 0, codeobj=7), + _occupancy(0, 0x20, codeobj=7), + _occupancy(1, 0x40, codeobj=8), + ], + shaderdata=[first, second], + ) + + _correct_marker_timestamps( + [(0, records)], + { + "sqtt_funcmap": [[7, 3, "P", "first", "", 0], [8, 4, "P", "second", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4], [8, 11, 5]], + }, + ) + + self.assertEqual([record.time for record in records.shaderdata], [0x1020, 0x1080]) + self.assertEqual([record.value for record in records.shaderdata], [16, 16]) + + def test_multi_payload_headers_are_normalized_without_clock_correction(self): + value = (0x100 << 20) | (3 << 2) + header = _shaderdata(0x1080, value) + payload0 = _shaderdata(0x1084, 0xDEADBEEF) + payload1 = _shaderdata(0x1088, 0xCAFEBABE) + records = TraceRecords( + occupancy=[_occupancy(1, 0)], + shaderdata=[header, payload0, payload1], + ) + + _correct_marker_timestamps( + [(0, records)], + { + "sqtt_funcmap": [[7, 3, "P", "multi", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + "sqtt_funcmap_payloads": [[7, 3, 2]], + }, + ) + + self.assertEqual(header.value, 12) + self.assertEqual(header.time, 0x1080) + self.assertEqual([payload0.time, payload1.time], [0x1084, 0x1088]) + self.assertEqual([payload0.value, payload1.value], [0xDEADBEEF, 0xCAFEBABE]) + + def test_no_decode_markers_preserves_packed_values(self): + value = (0x100 << 20) | (3 << 2) + records = TraceRecords( + occupancy=[_occupancy(1, 0)], + shaderdata=[_shaderdata(0x1080, value), _shaderdata(0x1010, value)], + ) + code_index = CodeIndex.from_document( + { + "sqtt_funcmap": [[7, 3, "P", "marker", "", 0]], + "sqtt_funcmap_layout": [[7, 12, 4]], + } + ) + + _Writer.records = None + with TemporaryDirectory() as directory: + with patch("rocprof_trace_decoder.att.Decoder", return_value=_Decoder(records)), \ + patch("rocprof_trace_decoder.att.RcvOutputWriter", _Writer), \ + patch("rocprof_trace_decoder.att.write_code_json"): + generate_att_outputs( + [AttTrace(Path(directory) / "trace.att", shader_engine=0)], + code_index=code_index, + output_dir=directory, + formats="json", + decode_markers=False, + ) + + self.assertIs(_Writer.records, records) + self.assertEqual([record.time for record in records.shaderdata], [0x1080, 0x1010]) + self.assertEqual([record.value for record in records.shaderdata], [value, value]) + + def test_no_decode_markers_cli_flag(self): + parser = build_argparser() + self.assertTrue(parser.parse_args(["trace.att"]).decode_markers) + self.assertFalse(parser.parse_args(["--no-decode-markers", "trace.att"]).decode_markers) + + +class MarkerPipelineTest(unittest.TestCase): + def test_late_records_recover_after_an_orphan_and_render(self): + funcmap = FuncMap(markers={1: ("outer", "function"), 2: ("inner", "function")}) + spans = [WaveSpan(100, 130, 1), WaveSpan(300, 360, 2)] + wave_spans = {(0, 1, 2, 3): spans} + dispatches = {"1": "kernel"} + records = [ + ShaderRecord(160, 1, 1, 2, 3, 0), # orphan exit after retirement + ShaderRecord(191, 6, 1, 2, 3, 0), # enter outer + ShaderRecord(207, 10, 1, 2, 3, 0), # enter inner + ShaderRecord(223, 1, 1, 2, 3, 0), # exit inner + ShaderRecord(239, 1, 1, 2, 3, 0), # exit outer + ] + + self.assertIs(find_wave_span_at(spans, 99), spans[0]) + self.assertIs(find_wave_span_at(spans, 299), spans[0]) + self.assertIs(find_wave_span_at(spans, 300), spans[1]) + + markers, addresses = preprocess_records(records, funcmap) + self.assertFalse(addresses) + folded = build_stacks(markers, {7: funcmap}, dispatches, wave_spans, {"kernel": 7}) + self.assertEqual( + folded, + { + "code_object_7;kernel;outer": (32, 2), + "code_object_7;kernel;outer;inner": (16, 1), + }, + ) + ET.fromstring(render_flamegraph_svg(folded)) + + events, unmatched = emit_wave_events( + markers, + (0, 1, 2, 3, 0), + dispatches, + {"kernel": 7}, + {7: funcmap}, + funcmap, + wave_spans, + 1.0, + ) + self.assertEqual(unmatched, 1) + self.assertEqual([event["ph"] for event in events], ["B", "B", "E", "E"]) + + +if __name__ == "__main__": + unittest.main() diff --git a/projects/rocprof-trace-decoder/test/codegen_test.py b/projects/rocprof-trace-decoder/test/codegen_test.py new file mode 100644 index 000000000000..625547ac656e --- /dev/null +++ b/projects/rocprof-trace-decoder/test/codegen_test.py @@ -0,0 +1,74 @@ +#!/usr/bin/env python3 +"""Focused schema checks for code.json SQTT funcmap metadata.""" + +from __future__ import annotations + +import sys +import unittest +from pathlib import Path +from unittest.mock import patch + +sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "python")) + +import rocprof_trace_decoder.codegen as codegen + + +class CodegenFuncmapTest(unittest.TestCase): + def generate(self, funcmaps: list[codegen._SqttFuncmap]): + code_objects = [ + codegen.CodeObject("packed.co", 10), + codegen.CodeObject("legacy.co", 11), + ][: len(funcmaps)] + with ( + patch.object(codegen, "_normalize_code_objects", return_value=code_objects), + patch.object(codegen, "build_address_ranges", return_value=[]), + patch.object(codegen, "parse_disassembly", return_value=[]), + patch.object(codegen, "read_symbol_labels", return_value=({}, {})), + patch.object(codegen, "_read_sqtt_funcmap", side_effect=funcmaps), + ): + return codegen.generate_code_artifacts(code_objects).code_index.document + + def test_adds_packed_metadata_without_changing_funcmap_rows(self): + doc = self.generate( + [ + codegen._parse_sqtt_funcmap( + b"M:shader_clock_bits=12;shader_clock_shift=4\n" + b"F:7:scope@source.cpp:1\n" + b"U:8:user\n" + b"R:8:extra_payload_count=1\n" + ) + ] + ) + + self.assertEqual( + doc["sqtt_funcmap"], + [[10, 7, "F", "scope", "source.cpp:1", 0], [10, 8, "U", "user", "", 0]], + ) + self.assertTrue(all(len(row) == 6 for row in doc["sqtt_funcmap"])) + self.assertEqual(doc["sqtt_funcmap_layout"], [[10, 12, 4]]) + self.assertEqual(doc["sqtt_funcmap_payloads"], [[10, 8, 1]]) + + def test_legacy_funcmap_exports_rows_without_a_packed_layout(self): + doc = self.generate( + [codegen._parse_sqtt_funcmap(b"F:7:scope\nR:7:extra_payload_count=1\n")] + ) + + self.assertEqual(doc["sqtt_funcmap"], [[10, 7, "F", "scope", "", 0]]) + self.assertNotIn("sqtt_funcmap_layout", doc) + self.assertEqual(doc["sqtt_funcmap_payloads"], [[10, 7, 1]]) + + def test_incomplete_packed_layout_is_ignored(self): + doc = self.generate([codegen._parse_sqtt_funcmap(b"M:shader_clock_bits=12\nF:7:scope\n")]) + + self.assertEqual(doc["sqtt_funcmap"], [[10, 7, "F", "scope", "", 0]]) + self.assertNotIn("sqtt_funcmap_layout", doc) + + def test_packed_layout_does_not_require_named_entries(self): + doc = self.generate([codegen._parse_sqtt_funcmap(b"M:shader_clock_bits=12;shader_clock_shift=4\n")]) + + self.assertNotIn("sqtt_funcmap", doc) + self.assertEqual(doc["sqtt_funcmap_layout"], [[10, 12, 4]]) + + +if __name__ == "__main__": + unittest.main() diff --git a/projects/rocprof-trace-decoder/test/data/navi4_heavy.tar.gz b/projects/rocprof-trace-decoder/test/data/navi4_heavy.tar.gz new file mode 100644 index 000000000000..dc1ea9189d7f Binary files /dev/null and b/projects/rocprof-trace-decoder/test/data/navi4_heavy.tar.gz differ diff --git a/projects/rocprof-trace-decoder/test/data/navi4_marker.tar.gz b/projects/rocprof-trace-decoder/test/data/navi4_marker.tar.gz new file mode 100644 index 000000000000..5b543a69bac4 Binary files /dev/null and b/projects/rocprof-trace-decoder/test/data/navi4_marker.tar.gz differ diff --git a/projects/rocprof-trace-decoder/test/markers/kernels/heavy.cpp b/projects/rocprof-trace-decoder/test/markers/kernels/heavy.cpp index 57d69131b6b3..81759dffe04a 100644 --- a/projects/rocprof-trace-decoder/test/markers/kernels/heavy.cpp +++ b/projects/rocprof-trace-decoder/test/markers/kernels/heavy.cpp @@ -98,25 +98,60 @@ using IVec4 = __attribute__((__vector_size__(4 * sizeof(uint32_t)))) uint32_t using Vec16 = __attribute__((__vector_size__(16 * sizeof(float)))) float; using Vec8 = __attribute__((__vector_size__(8 * sizeof(float)))) float; using float8x8 = __attribute__((__vector_size__(8 * sizeof(uint8_t)))) uint8_t; +using IVec2 = __attribute__((__vector_size__(2 * sizeof(int)))) int; +using IVec8 = __attribute__((__vector_size__(8 * sizeof(int)))) int; static_assert(sizeof(float8) == sizeof(uint8_t)); +#if defined(__gfx942__) || defined(__gfx950__) + inline __device__ uint64_t castfrom8x8(const float8x8& f) { return *reinterpret_cast(&f); } -inline __device__ Vec4 mfma(const float8x8& a, const float8x8& b) +inline __device__ Vec4 mfma(const float8x8& a, const float8x8& b, const Vec4& initial = Vec4{}) { - return __builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(castfrom8x8(a), castfrom8x8(b), Vec4{}, 0, 0, 0); + return __builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(castfrom8x8(a), castfrom8x8(b), initial, 0, 0, 0); } inline __device__ Vec4 mfma(const float8x8& a0, const float8x8& a1, const float8x8& b0, const float8x8& b1) { - Vec4 ret = __builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(castfrom8x8(a0), castfrom8x8(b0), Vec4{}, 0, 0, 0); - return __builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(castfrom8x8(a1), castfrom8x8(b1), ret, 0, 0, 0); + Vec4 ret = mfma(a0, b0); + return mfma(a1, b1, ret); } +#else + +#if defined(__gfx1200__) || defined(__gfx1201__) +inline __device__ IVec2 castfrom8x8(const float8x8& f) +{ + return *reinterpret_cast(&f); +} +inline __device__ Vec8 mfma(const float8x8& a, const float8x8& b, const Vec8& initial = Vec8{}) +{ + return __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12(castfrom8x8(a), castfrom8x8(b), initial); +} +#else +inline __device__ IVec8 castfrom8x8(const float8x8& f) +{ + return *reinterpret_cast(&f); +} +inline __device__ Vec8 mfma(const IVec8& a, const IVec8& b, const Vec8& initial = Vec8{}) +{ + return __builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8(a, b, 0, initial, 0, 0); +} +#endif + +inline __device__ Vec4 mfma(const float8x8& a0, const float8x8& a1, const float8x8& b0, const float8x8& b1) +{ + Vec8 ret = mfma(castfrom8x8(a0), castfrom8x8(b0)); + ret = mfma(castfrom8x8(a1), castfrom8x8(b1), ret); + return {ret[0], ret[1], ret[2], ret[3]}; +} + +#endif + // HEIGHT <= 4 // WIDTH > 4 causes register spill template @@ -309,94 +344,59 @@ void launchHip( int M, int N, int K ) { dim3 block(TBLOCK, 32, 1); + dim3 grid((M + 4*SHMBLOCK - 1) / SHMBLOCK / 4, N / SHMBLOCK); - int mi300_threads = SHMBLOCK * SHMBLOCK * 304; - int bpcu = M * N / mi300_threads; - - if (bpcu >= 8 && M >= 4*SHMBLOCK) - { - dim3 grid((M + 4*SHMBLOCK - 1) / SHMBLOCK / 4, N / SHMBLOCK); - - fp8_gemm_kernel<4><<>>( - reinterpret_cast(a), - reinterpret_cast(b), - reinterpret_cast(c), - scale_a, - scale_b, - M, N, K); - } - else if (bpcu >= 3 && M >= 2*SHMBLOCK) - { - dim3 grid((M + 2*SHMBLOCK - 1) / SHMBLOCK / 2, N / SHMBLOCK); - - fp8_gemm_kernel<2><<>>( - reinterpret_cast(a), - reinterpret_cast(b), - reinterpret_cast(c), - scale_a, - scale_b, - M, N, K); - } - else if (bpcu >= 1 || N > M || 2 * M * N > mi300_threads) - { - dim3 grid((M + SHMBLOCK - 1) / SHMBLOCK, N / SHMBLOCK); - - fp8_gemm_kernel<1><<>>( - reinterpret_cast(a), - reinterpret_cast(b), - reinterpret_cast(c), - scale_a, - scale_b, - M, N, K); - } - else - { - dim3 grid((M + SHMBLOCK - 1) / SHMBLOCK, 2 * N / SHMBLOCK); - - fp8_gemm_kernel<1, 2><<>>( - reinterpret_cast(a), - reinterpret_cast(b), - reinterpret_cast(c), - scale_a, - scale_b, - M, N, K); - } + fp8_gemm_kernel<4><<>>( + reinterpret_cast(a), + reinterpret_cast(b), + reinterpret_cast(c), + scale_a, + scale_b, + M, N, K); HIP_API_CALL(hipGetLastError()); HIP_API_CALL(hipDeviceSynchronize()); } + int main() { const int device = 0; - const int M = 6144; - const int K = 7168; - const int minN = 576; - const int maxN = 4608; + int M = 6144; + int K = 7168; + int N = 4608; hipDeviceProp_t devProp{}; HIP_API_CALL(hipGetDeviceProperties(&devProp, device)); HIP_API_CALL(hipSetDevice(device)); + if (devProp.major > 9 && devProp.minor == 0) + { + M /= 3; + K /= 3; + N /= 3; + } + else if (devProp.major > 9 && devProp.minor == 5) + { + M /= 2; + K /= 2; + } + Matrix a(K, M); - Matrix b(K, maxN); - Matrix c(M, maxN); + Matrix b(K, N); + Matrix c(M, N); Matrix scale_a(K/128, M); - Matrix scale_b(K/128, maxN/128); + Matrix scale_b(K/128, N/128); // warmup - launchHip(a.dev, b.dev, c.dev, scale_a.dev, scale_b.dev, M, minN, K); - launchHip(a.dev, b.dev, c.dev, scale_a.dev, scale_b.dev, M, maxN, K); + launchHip(a.dev, b.dev, c.dev, scale_a.dev, scale_b.dev, M, N, K); HIP_API_CALL(hipDeviceSynchronize()); roctxProfilerResume(0); int runs = 10; for (int i=0; i diff --git a/projects/rocprof-trace-decoder/test/markers/run_tests.py b/projects/rocprof-trace-decoder/test/markers/run_tests.py index b1954cf8be55..f9f4bf5c4889 100644 --- a/projects/rocprof-trace-decoder/test/markers/run_tests.py +++ b/projects/rocprof-trace-decoder/test/markers/run_tests.py @@ -51,7 +51,7 @@ PROJECT_ROOT = os.path.dirname(os.path.dirname(SCRIPT_DIR)) BUILD_DIR = os.environ.get("SQTT_BUILD_DIR", os.path.join(PROJECT_ROOT, "build")) PASS_PLUGIN = os.environ.get( - "SQTT_PASS_PLUGIN", os.path.join(BUILD_DIR, "lib", "SQTTInstrumentPass.so")) + "SQTT_PASS_PLUGIN", os.path.join(BUILD_DIR, "lib", "libsqttinstrumentpass.so")) INCLUDE_DIR = os.path.join(PROJECT_ROOT, "include", "rocprof_trace_decoder", "cxx") TEST_SOURCE = os.path.join(SCRIPT_DIR, "kernels", "auto.cpp") ADDR_TRACE_SOURCE = os.path.join(SCRIPT_DIR, "kernels", "addr_trace.cpp") @@ -119,15 +119,15 @@ def extract_funcmap(binary: str, tmpdir: str) -> str: # Find extracted code object import glob - cos = glob.glob(os.path.join(tmpdir, "*.hipv4-amdgcn-amd-amdhsa--*")) + cos = glob.glob(os.path.join(tmpdir, "*.hip*-amdgcn-amd-amdhsa--*")) if not cos: # llvm-objdump extracts relative to cwd or next to the binary base = os.path.basename(binary) - cos = glob.glob(os.path.join(tmpdir, f"{base}*.hipv4*")) + cos = glob.glob(os.path.join(tmpdir, f"{base}*.hip*-amdgcn-amd-amdhsa--*")) if not cos: # Try in the binary's directory bdir = os.path.dirname(binary) - cos = glob.glob(os.path.join(bdir, f"{os.path.basename(binary)}*.hipv4*")) + cos = glob.glob(os.path.join(bdir, f"{os.path.basename(binary)}*.hip*-amdgcn-amd-amdhsa--*")) if not cos: return "" @@ -368,8 +368,8 @@ def extract_funcmap(binary: str, tmpdir: str) -> str: "reject_funcmap": [], }, { - "name": "gfx9_addr_trace_nop_hazard", - "desc": "gfx9: s_nop 0 between s_ttracedata and s_mov_b32 m0 in addr trace", + "name": "gfx9_exec_m0_trace_spacing", + "desc": "gfx9: address EXEC traces use the explicit mov/nop/trace sequence", "env": {"SQTT_TRACE_ADDRESSES": "memory"}, "mode": "asm", "flags": ["--offload-arch=gfx90a"], @@ -379,7 +379,50 @@ def extract_funcmap(binary: str, tmpdir: str) -> str: "reject_ir": [], "expect_funcmap": [], "reject_funcmap": [], - "custom_asm_check": "nop_between_ttracedata_and_mov_m0", + "custom_asm_check": "gfx9_m0_trace_spacing", + }, + { + "name": "gfx9_full_m0_trace_nop", + "desc": "gfx9: every full M0 trace has s_nop 0 before s_ttracedata", + "env": {}, + "mode": "asm", + "source": "marker", + "flags": ["--offload-arch=gfx90a"], + "expect_ir": [ + r"s_ttracedata\b", + ], + "reject_ir": [], + "expect_funcmap": [], + "reject_funcmap": [], + "custom_asm_check": "m0_nop_before_ttracedata", + }, + { + "name": "gfx10_full_m0_trace_nop", + "desc": "gfx10: every full M0 trace has s_nop 0 before s_ttracedata", + "env": {"SQTT_TRACE_ADDRESSES": "memory"}, + "mode": "asm", + "flags": ["--offload-arch=gfx1030"], + "expect_ir": [ + r"s_ttracedata\b", + ], + "reject_ir": [], + "expect_funcmap": [], + "reject_funcmap": [], + "custom_asm_check": "m0_nop_before_ttracedata", + }, + { + "name": "gfx12_full_m0_trace_nop", + "desc": "gfx12: every full M0 trace has s_nop 0 before s_ttracedata", + "env": {"SQTT_TRACE_ADDRESSES": "memory"}, + "mode": "asm", + "flags": ["--offload-arch=gfx1200"], + "expect_ir": [ + r"s_ttracedata\b", + ], + "reject_ir": [], + "expect_funcmap": [], + "reject_funcmap": [], + "custom_asm_check": "m0_nop_before_ttracedata", }, { "name": "addr_trace_unique_ids", @@ -435,7 +478,7 @@ def extract_funcmap(binary: str, tmpdir: str) -> str: { "name": "gfx12_func_instrumentation", "desc": "Function instrumentation on gfx12", - "env": {"SQTT_INSTRUMENT_FUNCTIONS": "10"}, + "env": {"SQTT_INSTRUMENT_FUNCTIONS": "10", "SQTT_SHADER_CLOCK_BITS": "12"}, "mode": "obj", "flags": ["--offload-arch=gfx1200"], "expect_ir": [], @@ -493,7 +536,7 @@ def extract_funcmap(binary: str, tmpdir: str) -> str: { "name": "gfx12_asm_shader_clock", "desc": "gfx12 packs shader clock bits into full s_ttracedata markers", - "env": {"SQTT_INSTRUMENT_FUNCTIONS": "10"}, + "env": {"SQTT_INSTRUMENT_FUNCTIONS": "10", "SQTT_SHADER_CLOCK_BITS": "12"}, "mode": "asm", "flags": ["--offload-arch=gfx1200"], "expect_ir": [ @@ -504,6 +547,53 @@ def extract_funcmap(binary: str, tmpdir: str) -> str: "expect_funcmap": [], "reject_funcmap": [], }, + { + "name": "gfx1250_asm_shader_clock", + "desc": "gfx1250 uses s_get_shader_cycles_u64 for packed shader clocks", + "env": {"SQTT_INSTRUMENT_FUNCTIONS": "10", "SQTT_SHADER_CLOCK_BITS": "12"}, + "mode": "asm", + "flags": ["--offload-arch=gfx1250"], + "expect_ir": [ + r"s_get_shader_cycles_u64\b", + r"s_ttracedata\b", + ], + "reject_ir": [ + r"HW_REG_SHADER_CYCLES_LO", + r"s_ttracedata_imm\b", + ], + "expect_funcmap": [], + "reject_funcmap": [], + "custom_asm_check": "m0_nop_before_ttracedata", + }, + { + "name": "gfx12_default_no_clock_asm", + "desc": "gfx12: default layout keeps immediate headers and dynamic payload traces", + "env": {}, + "mode": "asm", + "source": "marker", + "flags": ["--offload-arch=gfx1200"], + "expect_ir": [ + r"s_ttracedata_imm\b", + r"v_readfirstlane_b32\b", + ], + "reject_ir": [r"HW_REG_SHADER_CYCLES_LO"], + "expect_funcmap": [], + "reject_funcmap": [], + "custom_asm_check": "m0_nop_before_ttracedata", + }, + { + "name": "gfx12_default_no_clock_funcmap", + "desc": "gfx12: default layout emits no funcmap layout row", + "env": {}, + "mode": "obj", + "source": "marker", + "flags": ["--offload-arch=gfx1200"], + "expect_ir": [], + "reject_ir": [], + "expect_funcmap": ["payload_value"], + "reject_funcmap": ["M:shader_clock_bits"], + "custom_funcmap_check": "data_marker_payload", + }, { "name": "gfx12_markers_no_global_inv", "desc": "gfx12 marker fences do not generate global cache invalidates", @@ -592,7 +682,7 @@ def run_test(test: dict, hipcc: str, verbose: bool) -> TestResult: try: if mode == "ir" or mode == "ir+obj": ir_path = os.path.join(tmpdir, f"{test['name']}.ll") - r = compile_test(env_vars, ["-S", "-emit-llvm"] + extra_flags, + r = compile_test(env_vars, ["-S", "-emit-llvm", "--offload-device-only"] + extra_flags, ir_path, hipcc, source=source) if r.returncode != 0: result.fail(f"Compilation failed (exit {r.returncode})") @@ -674,7 +764,7 @@ def run_test(test: dict, hipcc: str, verbose: bool) -> TestResult: ir_path = os.path.join(tmpdir, f"{test['name']}.ll") # Compile without pass plugin or SQTT_ENABLED env = os.environ.copy() - cmd = [hipcc, f"-I{INCLUDE_DIR}", "-S", "-emit-llvm", + cmd = [hipcc, f"-I{INCLUDE_DIR}", "-S", "-emit-llvm", "--offload-device-only", TEST_SOURCE, "-o", ir_path] r = subprocess.run(cmd, capture_output=True, text=True, env=env, timeout=120, cwd=PROJECT_ROOT) @@ -722,45 +812,52 @@ def check_funcmap(result: TestResult, funcmap: str, def check_custom_asm(result: TestResult, asm: str, check: str): """Run custom assembly validation checks.""" - if check == "nop_between_ttracedata_and_mov_m0": - # gfx9 hazard: s_mov_b32 m0 immediately followed by s_ttracedata - # requires an s_nop 0 in between. Extract instruction lines only. - insn_lines = [] - for line in asm.splitlines(): - stripped = line.strip() - if not stripped or stripped.startswith(";") or stripped.startswith("."): - continue - if stripped.endswith(":"): + if check not in {"m0_nop_before_ttracedata", "gfx9_m0_trace_spacing"}: + return + + # The pass emits direct M0 traces as mov/nop/trace. On gfx9, the backend + # may instead schedule an unrelated scalar instruction in that hazard + # slot, which is safe but means the final sequence is not contiguous. + strict_nop = check == "m0_nop_before_ttracedata" + insn_lines = [] + for line in asm.splitlines(): + stripped = line.strip() + if not stripped or stripped.startswith(";") or stripped.startswith("."): + continue + if stripped.endswith(":"): + continue + insn_lines.append(stripped) + + if strict_nop: + trace_count = 0 + for i, line in enumerate(insn_lines): + if not re.match(r"s_ttracedata\b", line): continue - insn_lines.append(stripped) - - violations = 0 - for i in range(len(insn_lines) - 1): - if re.match(r"s_mov_b32\s+m0\b", insn_lines[i]): - nxt = insn_lines[i + 1] - if re.match(r"s_ttracedata\b", nxt): - violations += 1 - result.fail( - f"s_mov_b32 m0 immediately followed by s_ttracedata " - f"(no s_nop): {insn_lines[i]!r} -> {nxt!r}") - if violations == 0: - # Verify we actually saw the pattern with the nop - saw_mov_m0 = False - saw_nop = False - for line in insn_lines: - if re.match(r"s_mov_b32\s+m0\b", line): - saw_mov_m0 = True - saw_nop = False - elif saw_mov_m0 and re.match(r"s_nop\b", line): - saw_nop = True - elif saw_mov_m0 and saw_nop and re.match(r"s_ttracedata\b", line): - break - else: - saw_mov_m0 = False - saw_nop = False - else: - result.fail("No s_mov_b32 m0 -> s_nop -> s_ttracedata sequence " - "found (expected at least one for exec mask hazard)") + trace_count += 1 + if (i < 2 or not re.match(r"s_nop\s+0\b", insn_lines[i - 1]) or + not re.match(r"s_mov_b32\s+m0\b", insn_lines[i - 2])): + before = insn_lines[max(0, i - 2):i + 1] + result.fail(f"full M0 trace does not use mov/nop 0/trace: {before!r}") + if trace_count == 0: + result.fail("No full M0 trace found") + return + + valid_sequences = 0 + for i in range(len(insn_lines) - 1): + if not re.match(r"s_mov_b32\s+m0\b", insn_lines[i]): + continue + nxt = insn_lines[i + 1] + if re.match(r"s_ttracedata\b", nxt): + result.fail( + f"s_mov_b32 m0 immediately followed by s_ttracedata " + f"(no spacing): {insn_lines[i]!r} -> {nxt!r}") + continue + if i + 2 < len(insn_lines) and re.match(r"s_nop\s+0\b", nxt) and re.match( + r"s_ttracedata\b", insn_lines[i + 2]): + valid_sequences += 1 + + if valid_sequences == 0: + result.fail("No s_mov_b32 m0 -> s_nop 0 -> s_ttracedata sequence found") def check_custom_funcmap(result: TestResult, funcmap: str, check: str): diff --git a/projects/rocprof-trace-decoder/test/unit/funcmap_test.cpp b/projects/rocprof-trace-decoder/test/unit/funcmap_test.cpp index c883b55ab3d5..c3746956f0b1 100644 --- a/projects/rocprof-trace-decoder/test/unit/funcmap_test.cpp +++ b/projects/rocprof-trace-decoder/test/unit/funcmap_test.cpp @@ -29,6 +29,7 @@ #include #include #include +#include #include using namespace rocprof_trace_decoder::codeobj; @@ -108,13 +109,111 @@ TEST(DecodeMarkerValue, ShaderClockEncoding) uint32_t raw = (0xABCu << 20) | (42u << 2) | 0b10u; auto m = decode_marker_value(raw, enc); EXPECT_EQ(m.id, 42u); - EXPECT_EQ(m.shader_clock, 0xABCu); + EXPECT_EQ(enc.decode_shader_clock(raw), 0xABCu); EXPECT_TRUE(m.is_enter); EXPECT_FALSE(m.exit_prev); Funcmap fm; fm.marker_encoding = enc; - EXPECT_EQ(decode_marker_value(raw, fm).id, 42u); + EXPECT_EQ(decode_marker_value(raw, fm).id, raw >> 2); +} + +TEST(MarkerEncoding, ExposesRawAndSourceClockMasks) +{ + MarkerEncoding enc{12, 4}; + + // The raw shaderdata word reserves bits [31:20] for the packed clock, + // bits [19:2] for the ID, and preserves the two low flags. + EXPECT_TRUE(enc.is_valid()); + EXPECT_TRUE(enc.has_shader_clock()); + EXPECT_EQ(enc.marker_id_mask(), 0x000FFFFCu); + EXPECT_EQ(enc.packed_shader_clock_mask(), 0xFFF00000u); + + // The sampled source-clock low word uses bits [15:4]. + EXPECT_EQ(enc.shader_clock_source_mask(), 0x0000FFF0u); + + uint32_t raw = (0xABCu << 20) | (42u << 2) | 0b10u; + EXPECT_EQ((raw & enc.marker_id_mask()) >> 2, 42u); + EXPECT_EQ(raw & enc.packed_shader_clock_mask(), 0xABC00000u); +} + +TEST(MarkerEncoding, LegacyEncodingRemainsFullWidthIdOnly) +{ + MarkerEncoding enc{}; + + EXPECT_TRUE(enc.is_valid()); + EXPECT_FALSE(enc.has_shader_clock()); + EXPECT_EQ(enc.marker_id_mask(), 0xFFFFFFFCu); + EXPECT_EQ(enc.packed_shader_clock_mask(), 0u); + EXPECT_EQ(enc.shader_clock_source_mask(), 0u); + + uint32_t raw = (0x3FFFFFFFu << 2) | 0b01u; + EXPECT_EQ(decode_marker_value(raw, enc).id, 0x3FFFFFFFu); + EXPECT_EQ(enc.decode_shader_clock(raw), 0u); + + MarkerEncoding invalid{30, 0}; + EXPECT_FALSE(invalid.is_valid()); + EXPECT_EQ(decode_marker_value(raw, invalid).id, 0x3FFFFFFFu); +} + +TEST(MarkerEncoding, SupportsBoundaryClockWindows) +{ + MarkerEncoding high_bit{1, 31}; + EXPECT_TRUE(high_bit.is_valid()); + EXPECT_EQ(high_bit.marker_id_mask(), 0x7FFFFFFCu); + EXPECT_EQ(high_bit.packed_shader_clock_mask(), 0x80000000u); + EXPECT_EQ(high_bit.shader_clock_source_mask(), 0x80000000u); + + MarkerEncoding wide{29, 3}; + EXPECT_TRUE(wide.is_valid()); + EXPECT_EQ(wide.marker_id_mask(), 0x00000004u); + EXPECT_EQ(wide.packed_shader_clock_mask(), 0xFFFFFFF8u); + EXPECT_EQ(wide.shader_clock_source_mask(), 0xFFFFFFF8u); +} + +TEST(DecodeMarkerValue, FuncmapLeavesUnregisteredValuesLegacy) +{ + Funcmap fm; + fm.marker_encoding = {12, 4}; + auto entry = std::make_shared(); + entry->kind = FuncmapEntryKind::Point; + entry->id = 42; + fm.entries.push_back(entry); + fm.by_id.emplace(entry->id, entry); + + uint32_t known = (0xABCu << 20) | (42u << 2); + EXPECT_EQ(decode_marker_value(known, fm).id, 42u); + EXPECT_EQ(fm.marker_encoding.decode_shader_clock(known), 0xABCu); + + uint32_t numeric = (0xABCu << 20) | (123u << 2); + EXPECT_EQ(decode_marker_value(numeric, fm).id, numeric >> 2); + + uint32_t exit = (0xABCu << 20) | 1u; + EXPECT_EQ(decode_marker_value(exit, fm).id, 0u); + EXPECT_EQ(fm.marker_encoding.decode_shader_clock(exit), 0xABCu); +} + +TEST(FuncmapCompatibility, ExistingAggregateInitializersStillWork) +{ + // Keep source compatibility for consumers that aggregate-initialize the + // public structs from the pre-packed-marker layout. + std::vector diagnostics; + Funcmap map{ + std::vector{}, + std::unordered_map{}, + 0, + diagnostics, + }; + FuncmapEntry entry{FuncmapEntryKind::Point, 7, "point", "source.cpp:1", 0x100000000ull}; + MarkerValue marker{7, true, false}; + + EXPECT_TRUE(map.diagnostics.empty()); + EXPECT_FALSE(map.marker_encoding.has_shader_clock()); + EXPECT_EQ(entry.vaddr, 0x100000000ull); + EXPECT_EQ(entry.extra_payload_count, 0u); + EXPECT_EQ(marker.id, 7u); + EXPECT_TRUE(marker.is_enter); + EXPECT_FALSE(marker.exit_prev); } // ─── parse_funcmap_section ────────────────────────────────────────────────── @@ -169,6 +268,17 @@ TEST(ParseFuncmap, ParsesExtraPayloadAndMarkerEncoding) EXPECT_EQ(p->extra_payload_count, 1u); EXPECT_EQ(m.marker_encoding.shader_clock_bits, 12u); EXPECT_EQ(m.marker_encoding.shader_clock_shift, 4u); + EXPECT_EQ(m.marker_encoding.marker_id_mask(), 0x000FFFFCu); + EXPECT_EQ(m.marker_encoding.packed_shader_clock_mask(), 0xFFF00000u); + EXPECT_EQ(m.marker_encoding.shader_clock_source_mask(), 0x0000FFF0u); +} + +TEST(ParseFuncmap, PackedMarkerEncodingRequiresClockShift) +{ + Funcmap m = parse_funcmap_section("M:shader_clock_bits=12\n", /*silent=*/true); + + EXPECT_FALSE(m.marker_encoding.has_shader_clock()); + EXPECT_TRUE(has_warning(m.diagnostics, "malformed M:")); } TEST(ParseFuncmap, ToleratesBlankLinesCRLFAndTrailingNUL) diff --git a/projects/rocprof-trace-decoder/test/unit/markers_test.cpp b/projects/rocprof-trace-decoder/test/unit/markers_test.cpp index 7d5d8ca4ecd5..7b06e815f729 100644 --- a/projects/rocprof-trace-decoder/test/unit/markers_test.cpp +++ b/projects/rocprof-trace-decoder/test/unit/markers_test.cpp @@ -29,6 +29,7 @@ #include "llvm/IR/Constants.h" #include "llvm/IR/Function.h" +#include "llvm/IR/InlineAsm.h" #include "llvm/IR/IRBuilder.h" #include "llvm/IR/IntrinsicsAMDGPU.h" #include "llvm/IR/Metadata.h" @@ -138,13 +139,36 @@ void useFullScopeMasks(SQTTConfig& config) config.MemBarrier = MemBarrierMode::None; } -CallInst* insertTraceCallBefore(Instruction* insertPt, uint32_t encoded) +CallInst* insertTraceCallBefore(Instruction* insertPt, uint32_t encoded, bool passHeader = false) { Module* module = insertPt->getModule(); LLVMContext& ctx = module->getContext(); IRBuilder<> builder(insertPt); Function* trace = Intrinsic::getOrInsertDeclaration(module, Intrinsic::amdgcn_s_ttracedata); - return builder.CreateCall(trace, {ConstantInt::get(Type::getInt32Ty(ctx), encoded)}); + CallInst* call = builder.CreateCall(trace, {ConstantInt::get(Type::getInt32Ty(ctx), encoded)}); + if (passHeader) call->setMetadata("sqtt.marker_header", MDNode::get(ctx, {})); + return call; +} + +Function* makeNamedMarkerSentinel(Module& module, StringRef name) +{ + LLVMContext& ctx = module.getContext(); + auto* type = FunctionType::get(Type::getVoidTy(ctx), {PointerType::get(ctx, 0)}, false); + return Function::Create(type, GlobalValue::ExternalLinkage, name, module); +} + +GlobalVariable* makeMarkerString(Module& module, StringRef value) +{ + LLVMContext& ctx = module.getContext(); + auto* initializer = ConstantDataArray::getString(ctx, value, true); + return new GlobalVariable( + module, + initializer->getType(), + true, + GlobalValue::PrivateLinkage, + initializer, + ".sqtt.marker.string" + ); } void addEarlyFunctionMetadata(Function& function, uint32_t id, unsigned preOptSize, StringRef sourceLoc) @@ -235,6 +259,29 @@ size_t countIntrinsicCalls(const Function& function, Intrinsic::ID id) return count; } +size_t countFences(const Function& function) +{ + size_t count = 0; + for (const BasicBlock& block : function) + for (const Instruction& inst : block) + if (isa(inst)) ++count; + return count; +} + +std::optional earlyFunctionPreOptSize(const Module& module, StringRef name) +{ + const NamedMDNode* entries = module.getNamedMetadata("sqtt.funcmap.early"); + if (!entries) return std::nullopt; + for (const MDNode* entry : entries->operands()) + { + if (entry->getNumOperands() < 3) continue; + auto* entryName = dyn_cast(entry->getOperand(1)); + auto* size = mdconst::dyn_extract(entry->getOperand(2)); + if (entryName && size && entryName->getString() == name) return size->getZExtValue(); + } + return std::nullopt; +} + std::vector traceMarkerValues(const Function& function) { std::vector values; @@ -255,6 +302,25 @@ std::vector traceMarkerValues(const Function& function) return values; } +const CallInst* findM0NopTrace(const Function& function) +{ + constexpr const char TraceAsm[] = + "s_mov_b32 m0, $1\n" + "s_nop 0\n" + "s_ttracedata"; + for (const BasicBlock& block : function) + { + for (const Instruction& inst : block) + { + auto* call = dyn_cast(&inst); + if (!call) continue; + auto* asmCall = dyn_cast(call->getCalledOperand()); + if (asmCall && asmCall->hasSideEffects() && asmCall->getAsmString() == TraceAsm) return call; + } + } + return nullptr; +} + bool hasPtrToIntFromAddressSpace(const Module& module, unsigned addressSpace, unsigned resultBits) { for (const Function& function : module) @@ -325,6 +391,25 @@ std::optional pointEntryId(const std::string& funcMap, StringRef name) return std::nullopt; } +size_t countPointEntries(const std::string& funcMap, StringRef name) +{ + size_t count = 0; + SmallVector lines; + StringRef(funcMap).split(lines, '\n', /*MaxSplit=*/-1, /*KeepEmpty=*/false); + for (StringRef line : lines) + { + line = line.rtrim("\r"); + if (!line.consume_front("P:")) continue; + auto [idText, rest] = line.split(':'); + unsigned id = 0; + if (idText.getAsInteger(10, id)) continue; + auto [entryName, sourceLoc] = rest.split('@'); + (void) sourceLoc; + if (entryName == name) ++count; + } + return count; +} + std::optional extraPayloadCountForId(const std::string& funcMap, unsigned markerId) { SmallVector lines; @@ -406,7 +491,7 @@ TEST(MarkerConfig, ParsesEnvironmentAndRejectsConflictingModes) EXPECT_EQ(config.MemoryMaxGap, 7u); EXPECT_FALSE(config.TraceMemoryAddrs); EXPECT_FALSE(config.TraceLDSAddrs); - EXPECT_EQ(config.ShaderClockBits, SQTTConfig::AutoShaderClockBits); + EXPECT_EQ(config.ShaderClockBits, 0u); EXPECT_EQ(config.ShaderClockShift, 8u); EXPECT_EQ(config.WaveMask, 0xFFFFFFFFu); EXPECT_EQ(config.SimdMask, 0x5u); @@ -433,6 +518,10 @@ TEST(MarkerTarget, ClassifiesArchitecturesAndInstructionCosts) EXPECT_EQ(getGfxGen(*makeVoidFunction(*module, "gfx11_func", "gfx1100")), GfxGen::RDNA); EXPECT_EQ(getGfxGen(*makeVoidFunction(*module, "gfx12_func", "gfx1200")), GfxGen::GFX12); EXPECT_EQ(getGfxGen(*makeVoidFunction(*module, "unknown_func", "notgfx")), GfxGen::Unknown); + EXPECT_FALSE(hasShaderCyclesU64(*makeVoidFunction(*module, "gfx1200_cycles", "gfx1200"))); + EXPECT_FALSE(hasShaderCyclesU64(*makeVoidFunction(*module, "gfx1201_cycles", "gfx1201"))); + EXPECT_TRUE(hasShaderCyclesU64(*makeVoidFunction(*module, "gfx1250_cycles", "gfx1250"))); + EXPECT_TRUE(hasShaderCyclesU64(*makeVoidFunction(*module, "gfx1250_features", "gfx1250:xnack+"))); EXPECT_EQ(getWaveSize(GfxGen::GFX9), 64u); EXPECT_EQ(getWaveSize(GfxGen::RDNA), 32u); @@ -440,10 +529,12 @@ TEST(MarkerTarget, ClassifiesArchitecturesAndInstructionCosts) EXPECT_TRUE(supportsImmTrace(GfxGen::GFX12)); SQTTConfig config; - EXPECT_EQ(getShaderClockBits(config, GfxGen::GFX12), 12u); + EXPECT_EQ(getShaderClockBits(config, GfxGen::GFX12), 0u); EXPECT_EQ(getShaderClockBits(config, GfxGen::RDNA), 0u); + EXPECT_FALSE(usesShaderClockPacking(config, GfxGen::GFX12)); config.ShaderClockBits = 5; EXPECT_EQ(getShaderClockBits(config, GfxGen::RDNA), 5u); + EXPECT_EQ(getShaderClockBits(config, GfxGen::GFX12), 5u); EXPECT_TRUE(usesShaderClockPacking(config, GfxGen::GFX12)); EXPECT_FALSE(usesShaderClockPacking(config, GfxGen::RDNA)); @@ -547,6 +638,177 @@ TEST(MarkerPass, AddressTracingHandlesBuffersAndPermutes) EXPECT_EQ(countIntrinsicCalls(*module, Intrinsic::amdgcn_readlane), 5u); } +TEST(MarkerPass, AddressTracingOnlyUsesFlatAndGlobalAddressSpaces) +{ + LLVMContext ctx; + auto module = makeModule(ctx); + Type* i32 = Type::getInt32Ty(ctx); + Type* voidTy = Type::getVoidTy(ctx); + FunctionType* type = FunctionType::get( + voidTy, + {PointerType::get(ctx, 0), + PointerType::get(ctx, 1), + PointerType::get(ctx, 2), + PointerType::get(ctx, 4), + PointerType::get(ctx, 5)}, + false + ); + Function* function = makeFunction(*module, "address_spaces", "gfx1100", type); + BasicBlock* entry = BasicBlock::Create(ctx, "entry", function); + IRBuilder<> builder(entry); + builder.CreateLoad(i32, function->getArg(0)); + builder.CreateStore(ConstantInt::get(i32, 1), function->getArg(1)); + builder.CreateLoad(i32, function->getArg(2)); + builder.CreateLoad(i32, function->getArg(3)); + builder.CreateLoad(i32, function->getArg(4)); + builder.CreateRetVoid(); + + SQTTConfig config; + useFullScopeMasks(config); + config.TraceMemoryAddrs = true; + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + + const std::string funcMap = getFuncMap(*module); + EXPECT_EQ(countPointEntries(funcMap, "addr_trace_load"), 1u); + EXPECT_EQ(countPointEntries(funcMap, "addr_trace_store"), 1u); + EXPECT_EQ(countPointEntries(funcMap, "addr_trace_atomic"), 0u); +} + +TEST(MarkerPass, AddressTracingSkipsBufferLoadLds) +{ + LLVMContext ctx; + auto module = makeModule(ctx); + Type* voidTy = Type::getVoidTy(ctx); + Function* function = makeVoidFunction(*module, "buffer_load_lds", "gfx1200"); + IRBuilder<> builder(function->getEntryBlock().getTerminator()); + Function* bufferLoadLds = Function::Create( + FunctionType::get(voidTy, false), + GlobalValue::ExternalLinkage, + "llvm.amdgcn.raw.buffer.load.lds", + module.get() + ); + builder.CreateCall(bufferLoadLds); + Function* point = makeNamedMarkerSentinel(*module, "__sqtt_named_marker_point"); + builder.CreateCall(point, {makeMarkerString(*module, "ordinary_point")}); + + SQTTConfig config; + useFullScopeMasks(config); + config.TraceMemoryAddrs = true; + // This would fail if .buffer.load.lds were misclassified as an address + // block during the preflight scan. + config.ShaderClockBits = 12; + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + + const std::string funcMap = getFuncMap(*module); + expectContains(funcMap, "M:shader_clock_bits=12;shader_clock_shift=4"); + expectNotContains(funcMap, "addr_trace_buffer_load"); + expectNotContains(funcMap, "W:"); +} + +TEST(MarkerPass, AddressTraceBlocksKeepDefaultGfx12ClockPackingDisabledAndUseBlockBoundaries) +{ + LLVMContext ctx; + auto module = makeModule(ctx); + Type* i32 = Type::getInt32Ty(ctx); + Type* voidTy = Type::getVoidTy(ctx); + Function* function = makeFunction( + *module, + "gfx12_address_trace", + "gfx1200", + FunctionType::get(voidTy, {PointerType::get(ctx, 1)}, false) + ); + BasicBlock* entry = BasicBlock::Create(ctx, "entry", function); + IRBuilder<> builder(entry); + builder.CreateLoad(i32, function->getArg(0)); + builder.CreateRetVoid(); + + SQTTConfig config; + useFullScopeMasks(config); + config.MemBarrier = MemBarrierMode::Fence; + config.TraceMemoryAddrs = true; + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + + const std::string funcMap = getFuncMap(*module); + expectContains(funcMap, "W:32"); + expectPointEntryWithPayload(funcMap, "addr_trace_load", 66); + expectNotContains(funcMap, "M:shader_clock_bits="); + EXPECT_EQ(countIntrinsicCalls(*function, Intrinsic::amdgcn_s_getreg), 0u); + EXPECT_EQ(countFences(*function), 2u); + EXPECT_EQ(countIntrinsicCalls(*function, Intrinsic::amdgcn_sched_barrier), 2u); +} + +TEST(MarkerPass, ForcedGfx12ClockPackingFailsForPayloadMarkers) +{ + EXPECT_DEATH( + { + LLVMContext ctx; + auto module = makeModule(ctx); + Type* i32 = Type::getInt32Ty(ctx); + Type* voidTy = Type::getVoidTy(ctx); + Function* function = makeFunction( + *module, + "forced_clock_address_trace", + "gfx1200", + FunctionType::get(voidTy, {PointerType::get(ctx, 1)}, false) + ); + BasicBlock* entry = BasicBlock::Create(ctx, "entry", function); + IRBuilder<> builder(entry); + builder.CreateLoad(i32, function->getArg(0)); + builder.CreateRetVoid(); + + SQTTConfig config; + useFullScopeMasks(config); + config.TraceMemoryAddrs = true; + config.ShaderClockBits = 12; + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + }, + "SQTT payload markers require SQTT_SHADER_CLOCK_BITS=0" + ); +} + +TEST(MarkerPass, ForcedGfx12ClockPackingFailsForNamedData) +{ + EXPECT_DEATH( + { + LLVMContext ctx; + auto module = makeModule(ctx); + Type* i32 = Type::getInt32Ty(ctx); + Type* voidTy = Type::getVoidTy(ctx); + Function* data = Function::Create( + FunctionType::get(voidTy, {PointerType::get(ctx, 0), i32}, false), + GlobalValue::ExternalLinkage, + "__sqtt_named_marker_data", + module.get() + ); + GlobalVariable* name = makeMarkerString(*module, "one_payload"); + Function* function = makeVoidFunction(*module, "gfx12_named_data", "gfx1200"); + IRBuilder<> builder(function->getEntryBlock().getTerminator()); + builder.CreateCall(data, {name, ConstantInt::get(i32, 17)}); + + SQTTConfig config; + useFullScopeMasks(config); + config.ShaderClockBits = 12; + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + }, + "SQTT payload markers require SQTT_SHADER_CLOCK_BITS=0" + ); +} + TEST(MarkerPass, Gfx9BufferAndPermuteAddressTracesUseInlineAsmExecProtocol) { LLVMContext ctx; @@ -590,6 +852,77 @@ TEST(MarkerPass, Gfx9BufferAndPermuteAddressTracesUseInlineAsmExecProtocol) expectContains(ir, "s_mov_b32 m0, exec_lo"); expectContains(ir, "s_nop 0"); expectContains(ir, "s_ttracedata"); + expectContains(ir, "={m0}"); +} + +TEST(MarkerPass, FullM0TracesUseExplicitNopSequenceOnEveryArchitecture) +{ + LLVMContext ctx; + auto module = makeModule(ctx); + const uint32_t markerValue = encodeMarker(64, false, false); + MDNode* traceMetadata = MDNode::get(ctx, {}); + + std::vector functions; + for (const auto& [name, cpu] : {std::pair{"gfx9_trace", "gfx90a"}, + std::pair{"gfx10_trace", "gfx1030"}, + std::pair{"gfx12_trace", "gfx1200"}}) + { + Function* function = makeVoidFunction(*module, name, cpu); + CallInst* trace = insertTraceCallBefore(function->getEntryBlock().getTerminator(), markerValue); + trace->setMetadata("sqtt.test.trace", traceMetadata); + functions.push_back(function); + } + + SQTTConfig config; + useFullScopeMasks(config); + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + + for (Function* function : functions) + { + const CallInst* trace = findM0NopTrace(*function); + ASSERT_NE(trace, nullptr) << function->getName().str(); + auto* inlineAsm = dyn_cast(trace->getCalledOperand()); + ASSERT_NE(inlineAsm, nullptr); + EXPECT_EQ(inlineAsm->getConstraintString(), "={m0},i"); + EXPECT_EQ(trace->getMetadata("sqtt.test.trace"), traceMetadata); + } +} + +TEST(MarkerPass, Gfx12PacksOnlyPassGeneratedHeaders) +{ + LLVMContext ctx; + auto module = makeModule(ctx); + Type* i32 = Type::getInt32Ty(ctx); + Function* function = makeFunction(*module, "numeric_marker", "gfx1200", FunctionType::get(i32, {i32}, false)); + BasicBlock* entry = BasicBlock::Create(ctx, "entry", function); + IRBuilder<> builder(entry); + Value* value = builder.CreateAdd(function->getArg(0), ConstantInt::get(i32, 1)); + builder.CreateAdd(value, ConstantInt::get(i32, 2)); + Instruction* ret = builder.CreateRet(value); + + // This mirrors the numeric marker API: it has no funcmap entry and must + // keep its full legacy value even when the same module has packed headers. + insertTraceCallBefore(ret, encodeMarker(1u << 20, false, false)); + Function* trace = Intrinsic::getOrInsertDeclaration(module.get(), Intrinsic::amdgcn_s_ttracedata); + builder.SetInsertPoint(ret); + builder.CreateCall(trace, {value}); + + SQTTConfig config; + useFullScopeMasks(config); + config.FunctionThreshold = 1; + config.ShaderClockBits = 12; + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + + // Entry and exit function markers read the shader clock; neither + // unregistered numeric marker reads it or trips the packed-ID range check. + EXPECT_EQ(countIntrinsicCalls(*function, Intrinsic::amdgcn_s_getreg), 2u); + expectContains(getFuncMap(*module), "M:shader_clock_bits=12;shader_clock_shift=4"); } TEST(MarkerPass, BarrierInstrumentationHandlesSplitAndStandaloneBarriers) @@ -606,12 +939,19 @@ TEST(MarkerPass, BarrierInstrumentationHandlesSplitAndStandaloneBarriers) FunctionCallee signal = Intrinsic::getOrInsertDeclaration(module.get(), Intrinsic::amdgcn_s_barrier_signal); FunctionCallee wait = Intrinsic::getOrInsertDeclaration(module.get(), Intrinsic::amdgcn_s_barrier_wait); FunctionCallee full = Intrinsic::getOrInsertDeclaration(module.get(), Intrinsic::amdgcn_s_barrier); + Function* work = Function::Create( + FunctionType::get(Type::getVoidTy(ctx), false), + GlobalValue::ExternalLinkage, + "barrier_work", + module.get() + ); builder.CreateCall(signal, {ConstantInt::get(i32, 0)}); builder.CreateCall(wait, {ConstantInt::get(i16, 0)}); builder.CreateCall(signal, {ConstantInt::get(i32, 0)}); - builder.CreateCall(full); + builder.CreateCall(work); builder.CreateCall(wait, {ConstantInt::get(i16, 0)}); + builder.CreateCall(full); SQTTConfig config; useFullScopeMasks(config); @@ -637,6 +977,67 @@ TEST(MarkerPass, BarrierInstrumentationHandlesSplitAndStandaloneBarriers) EXPECT_EQ(traceCount, 4u); } +TEST(MarkerPass, NamedExitEnterFusionRequiresDirectAdjacency) +{ + LLVMContext ctx; + auto module = makeModule(ctx); + Type* i32 = Type::getInt32Ty(ctx); + Function* exit = makeNamedMarkerSentinel(*module, "__sqtt_named_marker_exit"); + Function* enter = makeNamedMarkerSentinel(*module, "__sqtt_named_marker_enter"); + GlobalVariable* oldName = makeMarkerString(*module, "old"); + GlobalVariable* newName = makeMarkerString(*module, "new"); + + Function* separated = makeFunction( + *module, "separated_named_markers", "gfx1100", FunctionType::get(Type::getVoidTy(ctx), {i32}, false) + ); + BasicBlock* separatedEntry = BasicBlock::Create(ctx, "entry", separated); + IRBuilder<> separatedBuilder(separatedEntry); + separatedBuilder.CreateCall(exit, {oldName}); + separatedBuilder.CreateAdd(separated->getArg(0), ConstantInt::get(i32, 1)); + separatedBuilder.CreateCall(enter, {newName}); + separatedBuilder.CreateRetVoid(); + + Function* adjacent = makeVoidFunction(*module, "adjacent_named_markers", "gfx1100"); + Instruction* adjacentRet = adjacent->getEntryBlock().getTerminator(); + IRBuilder<> adjacentBuilder(adjacentRet); + adjacentBuilder.CreateCall(exit, {oldName}); + adjacentBuilder.CreateCall(enter, {newName}); + + SQTTConfig config; + useFullScopeMasks(config); + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + + std::vector separatedMarkers = traceMarkerValues(*separated); + ASSERT_EQ(separatedMarkers.size(), 2u); + EXPECT_EQ(separatedMarkers[0], FLAG_EXIT_PREV); + EXPECT_EQ(separatedMarkers[1] & FLAG_MASK, FLAG_ENTER); + + std::vector adjacentMarkers = traceMarkerValues(*adjacent); + ASSERT_EQ(adjacentMarkers.size(), 1u); + EXPECT_EQ(adjacentMarkers[0] & FLAG_MASK, FLAG_ENTER | FLAG_EXIT_PREV); +} + +TEST(MarkerPass, NumericMarkersAlwaysReceiveConfiguredMemoryBarriers) +{ + LLVMContext ctx; + auto module = makeModule(ctx); + Function* function = makeVoidFunction(*module, "numeric_marker", "gfx1100"); + insertTraceCallBefore(function->getEntryBlock().getTerminator(), encodeMarker(17, false, false)); + + SQTTConfig config; + useFullScopeMasks(config); + config.MemBarrier = MemBarrierMode::Fence; + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass pass(config, SQTTInstrumentPass::Mode::Late); + pass.run(*module, analysisManager); + + EXPECT_EQ(countFences(*function), 2u); +} + TEST(MarkerPass, DirectFunctionInstrumentationHandlesO0Fallback) { LLVMContext ctx; @@ -679,6 +1080,69 @@ TEST(MarkerPass, DirectFunctionInstrumentationHandlesO0Fallback) EXPECT_TRUE(traceMarkerValues(*kernel).empty()); } +TEST(MarkerPass, FunctionThresholdIgnoresPassMarkersBeforeAndAfterOptimization) +{ + LLVMContext ctx; + auto module = makeModule(ctx); + Function* function = makeVoidFunction(*module, "small_function", "gfx1100"); + + SQTTConfig config; + useFullScopeMasks(config); + config.FunctionThreshold = 1; + + ModuleAnalysisManager analysisManager; + SQTTInstrumentPass early(config, SQTTInstrumentPass::Mode::Early); + early.run(*module, analysisManager); + + EXPECT_EQ(earlyFunctionPreOptSize(*module, function->getName()), 1u); + + SQTTInstrumentPass late(config, SQTTInstrumentPass::Mode::Late); + late.run(*module, analysisManager); + + EXPECT_TRUE(traceMarkerValues(*function).empty()); + expectNotContains(getFuncMap(*module), "small_function"); +} + +TEST(MarkerPass, AutomaticFunctionInstrumentationSkipsMustTailFunctions) +{ + auto makeMustTailFunction = [](Module& module, StringRef name) + { + LLVMContext& ctx = module.getContext(); + Type* i32 = Type::getInt32Ty(ctx); + FunctionType* type = FunctionType::get(i32, {i32}, false); + Function* callee = + Function::Create(type, GlobalValue::ExternalLinkage, name.str() + ".callee", &module); + Function* function = makeFunction(module, name, "gfx1100", type); + BasicBlock* entry = BasicBlock::Create(ctx, "entry", function); + IRBuilder<> builder(entry); + CallInst* call = builder.CreateCall(callee, {function->getArg(0)}); + call->setTailCallKind(CallInst::TCK_MustTail); + builder.CreateRet(call); + return function; + }; + + SQTTConfig config; + useFullScopeMasks(config); + config.FunctionThreshold = 1; + ModuleAnalysisManager analysisManager; + + LLVMContext earlyCtx; + auto earlyModule = makeModule(earlyCtx); + Function* earlyFunction = makeMustTailFunction(*earlyModule, "early_musttail"); + SQTTInstrumentPass early(config, SQTTInstrumentPass::Mode::Early); + early.run(*earlyModule, analysisManager); + EXPECT_TRUE(traceMarkerValues(*earlyFunction).empty()); + EXPECT_FALSE(earlyModule->getNamedMetadata("sqtt.funcmap.early")); + + LLVMContext lateCtx; + auto lateModule = makeModule(lateCtx); + Function* lateFunction = makeMustTailFunction(*lateModule, "late_musttail"); + SQTTInstrumentPass late(config, SQTTInstrumentPass::Mode::Late); + late.run(*lateModule, analysisManager); + EXPECT_TRUE(traceMarkerValues(*lateFunction).empty()); + expectNotContains(getFuncMap(*lateModule), "late_musttail"); +} + TEST(MarkerPass, FunctionThresholdPrunesMarkersAndPreservesExistingLlvmUsed) { LLVMContext ctx; @@ -689,8 +1153,8 @@ TEST(MarkerPass, FunctionThresholdPrunesMarkersAndPreservesExistingLlvmUsed) Function* small = makeVoidFunction(*module, "small_function", "gfx1100"); Instruction* smallRet = small->getEntryBlock().getTerminator(); - insertTraceCallBefore(smallRet, encodeMarker(smallId, true, false)); - insertTraceCallBefore(smallRet, encodeMarker(smallId, false, true)); + insertTraceCallBefore(smallRet, encodeMarker(smallId, true, false), true); + insertTraceCallBefore(smallRet, encodeMarker(smallId, false, true), true); addEarlyFunctionMetadata(*small, smallId, 1, "small.hip:3"); Function* large = @@ -702,10 +1166,15 @@ TEST(MarkerPass, FunctionThresholdPrunesMarkersAndPreservesExistingLlvmUsed) builder.CreateRetVoid(); Instruction* firstLargeInst = &*large->getEntryBlock().getFirstInsertionPt(); Instruction* largeRet = large->getEntryBlock().getTerminator(); - insertTraceCallBefore(firstLargeInst, encodeMarker(largeId, true, false)); - insertTraceCallBefore(largeRet, encodeMarker(largeId, false, true)); + insertTraceCallBefore(firstLargeInst, encodeMarker(largeId, true, false), true); + insertTraceCallBefore(largeRet, encodeMarker(largeId, false, true), true); addEarlyFunctionMetadata(*large, largeId, 40, "large.hip:17"); + // This unregistered numeric marker happens to use the pruned function's + // old ID. It must not be removed or rewritten with pass-owned headers. + Function* numeric = makeVoidFunction(*module, "numeric_marker", "gfx90a"); + insertTraceCallBefore(numeric->getEntryBlock().getTerminator(), encodeMarker(smallId, true, false)); + addEarlyFunctionMapEntry(*module, 99, "inlined_large_function", 40, "inlined.hip:21"); addEarlyFunctionMapEntry(*module, 100, "inlined_small_function", 1, "inlined.hip:4"); addExistingLlvmUsed(*module); @@ -727,4 +1196,9 @@ TEST(MarkerPass, FunctionThresholdPrunesMarkersAndPreservesExistingLlvmUsed) EXPECT_EQ(countIntrinsicCalls(*small, Intrinsic::amdgcn_s_ttracedata), 0u); EXPECT_EQ(countIntrinsicCalls(*small, Intrinsic::amdgcn_s_ttracedata_imm), 0u); + const CallInst* numericTrace = findM0NopTrace(*numeric); + ASSERT_NE(numericTrace, nullptr); + auto* numericValue = dyn_cast(numericTrace->getArgOperand(0)); + ASSERT_NE(numericValue, nullptr); + EXPECT_EQ(numericValue->getZExtValue(), encodeMarker(smallId, true, false)); }