Repository navigation
perf(rocr): resolve device-memory batch signals without reading the slot - #14115
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
cfreeamd
left a comment
cfreeamd
left a comment
There was a problem hiding this comment.
Looked at all three changed files: the hsa_ext_amd.cpp routing of hsa_amd_signal_wait_all/_any validity checks through Signal::IsValidHandle is a clean consolidation, and the slot-offset bounds check in FindShadowedMember (verifying slot_offset % sizeof(SharedSignal) == offsetof(amd_signal) before dividing) correctly validates handles. Main concern is in signal.cpp (inline).
cfreeamd
left a comment
cfreeamd
left a comment
There was a problem hiding this comment.
🤖 AI-assisted pre-review (for @cfreeamd) — follow-up after author responses
Both prior findings are resolved.
[Important] Use-after-free rebuttal accepted. The lifetime argument is correct: SignalBatch is reference-counted via AcquireSlot/ReleaseSlot in each LocalSignal constructor/destructor, so the batch can only reach refcount 0 when all member LocalSignal objects are destroyed. FindShadowedMember loading a non-null page entry and then racing with batch destruction would require the caller's handle to already be post-hsa_signal_destroy — undefined behavior per HSA spec, and equally a CPU use-after-free on the old path (which read core_signal and called a virtual method on the freed Signal object). The new code also improves the stale-handle case: a destroyed member whose batch is still alive now returns HSA_STATUS_ERROR_INVALID_SIGNAL instead of a dangling pointer. Finding dropped.
[Question] Benchmark data answered. MI300X numbers look good — device batch resolve improves ~3× (3.2–3.4 µs → 1.1–1.2 µs for hsa_signal_load_relaxed, and ~15 µs → ~1.2 µs for wait_any/wait_all). Host-path overhead is +2–3 ns per convert/wait and +10–25 ns for wait_any/wait_all when device batches are registered (proportional to batch count, ~1 radix-tree walk per handle per resolve). Acceptable tradeoff.
One small item before merge: the response says "I've added a comment at the dereference stating this lifetime constraint" but the PR still shows only the original pre-review commit. Please push the lifetime comment (or clarify if it will be a fixup commit).
Every signal API call starts with Signal::Convert, which reads the slot's id and core_signal before the operation reads amd_signal_t::value. For a batch created in a device pool the host reaches over PCIe, each of those reads is an uncached bus round trip, so hsa_signal_load_relaxed cost three reads (3.5 us on MI300X) instead of one (1.2 us), and every wait paid the same two extra reads on entry. Batches whose arena the host reaches across the bus now keep a host copy of each member's Signal*, and their pages are recorded in a page map: a three-level radix tree on the address's page number, as in tcmalloc's PageMap. Arenas are page aligned and own their pages, so each page maps to at most one batch. Convert walks the map first, from host memory only, and falls back to reading the slot for every other handle. Readers take no lock; nodes are created on first use and never freed, so a walk always completes. hsa_amd_signal_wait_any/all validate handles the same way. A destroyed batch member now fails with HSA_STATUS_ERROR_INVALID_SIGNAL: its host entry is cleared, whereas the slot's id is never cleared. MI300X, 300 device batches registered: - device-memory signal load: 3.5 us -> 1.1 us (one bus read) - host-memory signal load: 4 ns -> 7 ns, flat from 1 to 8 threads
6780185 to
a72c529
Compare
TheRock Submodule Bump ActivityNewest first
|
…#14338) ## Motivation Two host-side costs remain for signals in device memory after #14115: 1. **XGMI-attached GPUs get no shadow.** #14115 keeps a host-side copy of each batch member's `Signal*` so `Signal::Convert` resolves a device-memory batch handle without reading the slot's `id` and `core_signal` across the link. The copy was only made when `GpuAgent::PcieStoresNeedFlush()` is true, which is `!xgmi_cpu_gpu_`. On XGMI-attached systems, every signal API call on a device-memory batch member still read two fields across XGMI. 2. **Reading timestamps costs two round trips.** `SharedSignal::GetRawTs` reads `start_ts` and `end_ts` with two 8-byte loads, so `hsa_amd_profiling_get_dispatch_time` and `hsa_amd_profiling_get_async_copy_time` cost two uncached reads per device-memory signal. CLR's built-in profiler (`GPU_CLR_PROFILE_OUTPUT`) attaches a signal to every dispatch, so this cost is paid once per kernel. ## Technical Details - `hsa_amd_signal_batch_create` shadows every batch whose arena is in a discrete GPU's device pool (`SignalBatch::Create` takes the decision as `shadow_members`). The host reads a discrete GPU's memory across PCIe or XGMI either way. `host_access_via_hdp` is unchanged: it still decides HDP flushes after stores and emulated read-modify-writes, which only apply over PCIe. - An APU's device pools (`properties().Integrated`) are not shadowed: the host shares that memory, so reading the slot costs no link round trip. Before, an APU's batches were shadowed when its CPU link was not reported as XGMI. - `SharedSignal::GetRawTs` reads both timestamps with one 16-byte SSE2 load on x86-64. They fill one aligned 16-byte chunk of `amd_signal_t`. `static_assert`s pin the layout and the alignment the load relies on: `alignof(amd_signal_t) >= 16` and `amd_signal` at a 16-byte offset in `SharedSignal`. Both signal allocators align to `SharedSignal`. Other architectures keep the two loads. A host read of device memory on MI300X: 8-byte load 1.1 us, 16-byte load 1.1 us (one transaction), 32- or 64-byte load 2.2 us (two transactions). ## Test Plan - rocrtst `*Signal*` and IPC tests. - Shadow stress (300 device batches, concurrent member loads while batches are recreated). - Per-call signal API cost on host and device batch signals. - `hip_launch_latency -k -g` with `GPU_CLR_PROFILE_OUTPUT` set, CLR device-memory queue signals on and off, against `develop`'s ROCr; kernel durations from the traces compared. - A pageable-copy stress test under the profiler with `HSA_ENABLE_SDMA=1` and `0`, checking copy durations in the trace (`hsa_amd_profiling_get_async_copy_time`). ## Test Result MI300X (PCIe host): - rocrtst: 14/14 pass. Stress: pass (11k batch recreations, 15.7M concurrent member loads, 0 mismatches). - Per call on a device-memory batch signal, median of 15 rounds of 5000 calls, `develop` and this branch interleaved: `hsa_amd_profiling_get_dispatch_time` 2.24-2.47 -> 1.15-1.27 us; `hsa_signal_load_relaxed` unchanged at 1.06-1.25 us. Host-memory signals: 6-7 ns both. - Profiled launches with device-memory signals: 100-node graph enqueue 463 -> 351 us, single launch + sync 11.7 -> 10.5 us. Host-memory signals unchanged. - Traces: no zero or negative kernel or copy durations; host-signal kernel median identical (3.328 us) on both ROCr builds; copy durations consistent with SDMA on and off. Not run on an XGMI-attached system or an APU. On XGMI, device batches now take the shadow path that #14115 measured over PCIe, where dropping the slot's `id` and `core_signal` reads took `hsa_signal_load_relaxed` from 3.5 to 1.1 us; the XGMI saving is two reads across XGMI per call, not measured here. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
Motivation
hsa_amd_signal_batch_create(#12988) can place signals in a GPU pool the host reachesthrough the large BAR. Every signal API call starts with
Signal::Convert, which reads theslot's
id(validation) andcore_signal(thecore::Signal*behind the handle) before theoperation reads
amd_signal_t::value. In device memory each of those is an uncached PCIeround trip, so one
hsa_signal_load_relaxedcosts three bus reads instead of one:hsa_signal_load_relaxedbeforehsa_signal_load_relaxedafterhsa_signal_wait_scacquire(satisfied) beforehsa_signal_wait_scacquire(satisfied) afterhsa_amd_signal_wait_any/wait_all(1 signal) beforehsa_amd_signal_wait_any/wait_all(1 signal) afterEvery wait paid the same two extra reads on entry;
wait_any/wait_allread the slot'sidand
core_signalseveral times per signal. CLR'shipEventSynchronizeon a device-memoryevent signal goes through
Convertand the wait.Technical Details
The handle has to stay the
amd_signal_taddress (packets carry it and the CP reads it), sothe handle-to-object lookup can't become a cast. Instead, batches whose arena the host reaches
across the bus get a host-side copy of what
Convertneeds:SignalBatch::shadow_members_: host array of each member'sSignal*, filled inCreate,cleared in
ReleaseSlot.SignalBatch*: a three-level radix tree on the page number(12 bits per level, 48-bit addresses), as in tcmalloc's PageMap. Arenas are page aligned and
own their pages, so a page belongs to at most one batch. Writers (batch create/destroy) take
a mutex; readers take no lock and write nothing shared. Nodes are created on first use and
never freed, so a walk always completes.
Signal::Convertwalks the map first when any batch is registered and falls back to readingthe slot for every other handle (host signals, IPC, non-batch device memory).
hsa_amd_signal_wait_any/allvalidate handles through the same path(
Signal::IsValidHandle).SharedSignalis unchanged;core_signalandidare still written for every slot and stayat 0x48/0x50 for IPC peers.
Example
A batch of 40 signals whose arena lands at
0x7F12_3456_7000. 40 x 128 bytes = 5120 bytes,so it spans two pages. Each 36-bit page number splits into three 12-bit indices:
0x7F12345670000x7F12345670x7F10x2340x5670x7F12345680000x7F12345680x7F10x2340x568Create. Under the mutex,
root[0x7F1]and itsleaves[0x234]are allocated if missing,then both leaf entries are set to the batch:
Lookup. Member 35's handle is
0x7F1234567000 + 35 * 128 = 0x7F1234568180:0x7F1234568: root0x7F1, mid0x234, leaf0x568gives the batch.0x1180= 4480: a multiple of 128, slot 35, below the batch's 40 slots.shadow_members_[35]is theSignal*. No device memory has been read; the load itselfthen reads
amd_signal_t::value, its one bus read.Misses. A host signal at
0x55D012345040stops atroot[0x55D] == nullptrand resolvesthrough its slot as before. A handle past the last slot (
0x7F1234568400, slot 40) or off aslot boundary also falls back.
Destroy.
hsa_signal_destroyon member 35 clearsshadow_members_[35]; a later call onthat handle fails with
HSA_STATUS_ERROR_INVALID_SIGNAL. Destroying the batch clears bothleaf entries; the nodes stay for the next batch in that range.
Behaviour change
A destroyed batch member now fails with
HSA_STATUS_ERROR_INVALID_SIGNAL. Before, the slot'sidwas never cleared (SharedSignalis trivially destructible for IPC), soConvertreturnedthe stale
core_signal. Using a destroyed handle was already undefined.Issue Tracking
Test Plan
*Signal*and IPC tests.threads resolving members while the main thread destroys and recreates batches for 5 s).
DEBUG_CLR_DEV_SIGNALS=0/1): test(hip): wait on a pending event from a stream with no other work #13818 repro, fork graph.Test Result
MI300X (gfx942):
wait_anyover adevice and a host signal returns the device signal.
develop(925f3d5), satisfied signals, 1 and 300 extra device batches: host-signalload +2-3 ns, wait +1-2 ns,
wait_any/wait_all+10-25 ns; device-signal load and wait3.3 us -> 1.1-1.2 us,
wait_any/wait_all~16 us -> ~1.2 us.Submission Checklist