Skip to content

perftest: Fix ping-pong flag reset ordering#98

Open
Lyapunova-cd wants to merge 1 commit into
NVIDIA:develfrom
Lyapunova-cd:fix-ping-pong-flag-reset-ordering
Open

perftest: Fix ping-pong flag reset ordering#98
Lyapunova-cd wants to merge 1 commit into
NVIDIA:develfrom
Lyapunova-cd:fix-ping-pong-flag-reset-ordering

Conversation

@Lyapunova-cd

Copy link
Copy Markdown

Complete each local flag reset before entering a global barrier. This prevents a slower PE from overwriting the peer's first signal after the timed kernel launches begin.## Summary

  • synchronize each PE after resetting its local ping-pong flag
  • add a global barrier before starting each timed collective launch
  • keep both operations outside the measured interval

Why

The timed loop resets each PE's symmetric flag immediately before a
collective launch. Unlike the warmup path and the other ping-pong tests in
this directory, it did not synchronize the reset across PEs.

This permits the following ordering:

  1. PE 0 finishes its reset and launches the kernel.
  2. PE 0 writes the first signal to PE 1's flag.
  3. PE 1 performs its delayed reset and overwrites that signal with zero.
  4. PE 1's kernel waits indefinitely for the lost signal.

A device-side atomic signal store, volatile store, or fence cannot repair
this race because the missing happens-before relationship is between the
host-initiated reset on one PE and the kernel launch on another PE.

The issue was reproduced on a downstream CUDA-compatible backend. This
change mirrors the reset ordering already used by the warmup path and the
other symmetric-flag ping-pong tests.

Testing

  • downstream full library and perftest build
  • 20 consecutive shmem_p_ping_pong_latency runs:
    • 10 with the legacy/static symmetric heap
    • 10 with the VMM symmetric heap
    • sizes 4 B through 16 KiB, 50 warmup iterations, 500 timed iterations
  • three runs each of the four sibling ping-pong latency tests

Notes

The local device synchronization is needed to complete the local reset.
The following nvshmem_barrier_all() is needed to establish cross-PE
ordering. Neither operation is included in the latency measurement.

Complete each local flag reset before entering a global barrier. This
prevents a slower PE from overwriting the peer's first signal after the
timed kernel launches begin.

Signed-off-by: hongwei.liu <hongwei.liu@mthreads.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants