Repository navigation
fabric/tp: HIP device side of the TCP Kind - #2961
Conversation
471f316 to
8255b22
Compare
|
Updated after another review pass; the description is current.
|
|
Thanks @Shoestring-dev. This is careful work, and your update already fixed the two things we'd have raised on the first head: the receipt now carries its six logs, which match We read the device side again on 8255b22. Word reuse across the 256-call ring, the release on failure, the parts (including the gather's rank-major offsets), the part number in the header, the capture refusal and On your questions:
The two-host USB4 run doesn't block this one. Please post it when the cables arrive, and in any case before the pull request that wires TP=2 to it. Next we'll put this head through our landing checks and land it, and we'll write the CHANGELOG line with credit to you.
|
|
The two-host USB4 run, at 8255b22 (this head), as promised.
Both ranks print PASS on both cables, for These are rank 0's numbers, the median and p90 of 500 back-to-back calls. A
The logs, both hosts' environments before and after, the build line and SHA256SUMS are available if useful; nothing here changes the branch. |
stage.zig stages device memory through tcp.zig behind the same Collective VTable (Kind.tcp). A call only queues on its stream: the stream writes the call's ready word, then waits for its done word, 64-bit words in pinned host memory. A proxy thread sees ready, copies the input into a pinned slab on a stream of its own, moves it with tcp.zig, copies the result back and sets done. The caller's stream gets no copies of ours: with copies queued on it around the wait, two ranks of one process hung or read stale bytes on ROCm 10.0. A call larger than Options.part_bytes (default 4 MiB, checked alike when the ranks open) moves in parts, and the frame header's spare word now names the part, so one call in two parts never pairs with two whole calls. A failure sets every done word, so no stream waits on it, and shows at the next call or at check. A capturing stream is refused, and allGatherV moves whole strides. stage_hip.zig is the HIP device: hipStreamWriteValue64 and hipStreamWaitValue64 on mapped coherent host memory, refused by name on a GPU without them, and the proxy's copies on a non-blocking stream of its own. stage_test.zig runs host_test's contract bodies over a device whose copies are memcpy, whole and in 64-byte parts, plus part, size, capture and device-failure refusals; stage_later_test.zig uses streams that run later, so calls return before anything moves and an abort still lets a held stream drain; stage_hip_test.zig checks the same on a GPU, two ranks of one device over loopback, and skips without HIP. `tf-tp-test tcp-hip` runs tcp-host's checks and timings from HIP memory. Nothing selects it yet, and no other backend changes.
8255b22 to
5bf750d
Compare
|
Thank you @Shoestring-dev. This is merged into main with your commits under your own name, so you're listed as a TensorFold contributor. It ships in the next release, credited to you in the changelog and the release notes. |
What this changes
fabric/tp's TCP Kind gets a device side:stage.zigstages device memory throughtcp.zigbehind the sameCollectiveVTable (stillKind.tcp), andstage_hip.zigis HIP's half, so a HIP engine can hand it device pointers and its stream and each call only queues. Nothing selects it yet: only its tests andtf-tp-test tcp-hipopen it. How a call runs is inzig/src/fabric/tp/README.md("Device memory over the TCP Kind").atc overlap --json --pulls 2917(ATC c225834) on this head, base eab305d: complete, no overlap with fabric/mesh: one launch description, --tp up to 64 and the rank aliases (T1a of #2560) #2917, the open pull request touchingzig/src/fabric/tp. fabric/mesh: GLM CUDA shares and sums through one mesh layer (T0 of #2560) #2789 landed after this was first prepared; it rebased with no conflicts (onlytp.zig's neighbouring export line differs).Why the caller's stream gets only two 64-bit words and no copies: with the copies queued on it around the wait, two ranks of one process hung, or read stale bytes, once a call reached 16 KiB (ROCm 10.0, gfx1151). A standalone HIP program with no TensorFold code did the same, and passed with the copies on a stream of their own, so I haven't pinned the cause inside the runtime. The probe sources are available if useful.
Checks on this head
python3 tools/pr_check.py --base eab305db0 --pr 2961: 0 errors; lean 345 vs 345, 0 new; 1 advisory warning (no Unreleased bullet; the maintainer writes it).zig build test-tp --summary allat 954e82d: 210 of 211 pass, 1 skip (roce_test, needs HCAs); the fourstage_hip_testGPU tests are among the passes. The tp test binary with the GPU tests, 20 times in a row: 20 of 20 (at 8566dfde9, which differs from 954e82d by one comment).zig build test-tp --summary all: 141 of 148 pass, 7 skip (the HIP tests without a HIP runtime, and RDMA).zig fmt --check build.zig zigandgit diff --check eab305db0: pass.stage,tcpandstage_hipcode as here apart from the bounded copy wait, the part-size check's order and the GPU tests' skip): the header's part check removed, the stream's wait ondoneremoved, the proxy not waiting forready, the HIP download not waited for, and the part-size check at open removed; each failed exactly the tests that cover it.stage.zigis in the POSIX-only TCP path, liketcp.zig; CI covers the compile); discrete AMD GPUs; ROCm 10.1 (no kernels here, so no HIP records); CUDA (nothing implementsstage.Devicefor it here).New tests:
stage_test.zig(host_test'sCallsandPairwiseon two to four staged ranks, whole and in parts; desync, abort, part numbers and sizes, capture,allGatherV, failed device calls),stage_later_test.zig(calls queued on held streams move nothing until the stream runs, then every result equals the reference),stage_hip_test.zig(the same on a real GPU behind ahipStreamWaitValue64gate, and a proxy copy still queued past its limit fails its call instead of waiting forever), andtcp_test.zig(ranks that differ only in the part number are refused). The GPU tests skip without HIP or a GPU that can wait on a stream value.Known limits
Deviceinterface.doneword holds its hardware queue. Past the runtime's queue count (GPU_MAX_HW_QUEUES, 4 by default) the proxy's stream can share a queue with a parked stream; the proxy waits for each copy at mostHip.limit_ns(30 s), then fails the call (BackendFailed) and frees every stream. That bounded wait costs about 3 µs on the smallest calls againsthipStreamSynchronize(45.5 against 41.3 µs at 16 B); large calls are unchanged.Rank.init(intcp.zig) still checks world <= 8 by debug assert only; nothing selectsKind.tcpyet, and the pull request that does will refuse a larger world by name.Receipt
n/a for inference: no engine calls this yet. Its own receipt is
tf-tp-test tcp-hip, indocs/receipts/tp-hip-device-954e82df/with every log, the environment, a README and SHA256SUMS: two ranks over 127.0.0.1 on one host, both on one GPU, so it shows the device side bit for bit and its cost overtcp-hoston the same link, not the link.Rank 0, median and p90 of 500 back-to-back calls; a
tcp-hiptime is the call and its stream's synchronisation, "queued" the call alone:A peer killed with SIGKILL 3 s into the run: rank 0 ended
error: PeerFailed, exit 1, 0.124 s later. For scale, #2718's receipt over USB4 measured 11 µs at 16 B and 2.4 ms for a 4 MiB gather, againsttcp-host's 2.2 µs and 0.73 ms on loopback here.Pending: the receipt the plan names, two hosts over USB4 networking (two cables, MTU 1500), waits for our second pair's cables. It will be the same command per cable, as for #2718:
Open questions
Hip.init(lib)), instead of addinghipStreamWriteValue64,hipStreamWaitValue64andhipStreamQuerytozig/src/hip/abi.zig'sApi, which every HIP engine requires at load. Would you rather have them inApi?hipStreamWaitValue64is marked beta in HIP's headers, and its docs speak ofhipMallocSignalMemory; on gfx1151 it works on mapped coherent host memory. Is that acceptable upstream, or should the words live in signal memory with the proxy writing them another way?timeout_ns, or a high-priority stream for the proxy so it never shares a queue with a parked stream?Checklist