Using the CCO SDMA device API

ccoSdma drives the GPU's SDMA copy engines from inside a kernel: intra-node peer-to-peer copies that do not occupy CUs. It is a device-side API in include/mori/cco/cco.hpp — one header, no separate library call per op.

Only for peers on the same node (LSA). Cross-node goes through ccoGda.

Every number quoted below was measured on MI355X (gfx950), two ranks over intra-node xGMI. Treat them as shape, not as a spec.

Turning it on

Build: SDMA is compiled out unless both are set.

cmake -B build -DBUILD_CCO=ON -DBUILD_CCO_SDMA=ON -DGPU_TARGETS=gfx950

Kernel TUs that include cco.hpp need -DBUILD_CCO_SDMA=1 too, or ccoSdma will not be declared.

Runtime:

Queue count is a property of the comm, not the environment: set reqs.sdmaQueueCount before ccoDevCommCreate. MORI_SDMA_NUM_CHANNELS is only the fallback when it is left at 0. Default 2, hardware cap 8 (clamped with a warning above that).

The count is fixed by the first ccoDevCommCreate on a comm; later ones with a different value warn and keep the original.

Host setup

ccoCommCreate(uid, nRanks, rank, vmmSize, &comm);
ccoMemAlloc(comm, bytes, &buf);              // both src and dst
ccoWindowRegister(comm, buf, bytes, &win);   // both src and dst

ccoDevCommRequirements reqs = CCO_DEV_COMM_REQUIREMENTS_INITIALIZER;
reqs.gdaConnectionType = CCO_GDA_CONNECTION_NONE;   // SDMA needs no GDA
reqs.gdaContextCount = reqs.gdaSignalCount = reqs.gdaCounterCount = 0;
reqs.sdmaQueueCount = 8;                            // 0 = env / default 2
ccoDevComm devComm{};
ccoDevCommCreate(comm, &reqs, &devComm);

if (devComm.sdma.sdmaNumQueue == 0) { /* SDMA unavailable — MORI_ENABLE_SDMA? */ }

Pass devComm to the kernel by value; construct ccoSdma sdma{devComm} inside.

The API

Everything hangs off ccoSdma sdma{devComm}, constructed per kernel.

put/get are non-blocking: they place a packet and (unless Aggregate) ring the doorbell. Nothing has happened when they return — you must drain.

All indices are LSA ranks (devComm.lsaRank, devComm.lsaSize), never world ranks. queueId must be < devComm.sdma.sdmaNumQueue; out of range is a silent no-op.

put / get

template <typename Coop            = ccoCoopThread,
          bool localSignal         = false,
          bool remoteSignal        = false,
          uint32_t optFlags        = ccoSdmaOptFlagsDefault,
          ccoSdmaThreadMode mode   = ccoSdmaThreadIndependent>
void put(int peer, ccoWindow_t dstWin, size_t dstOffset,
         ccoWindow_t srcWin, size_t srcOffset, size_t bytes, int queueId = 0);

peer is an LSA rank (devComm.lsaRank is yours, devComm.lsaSize the count) — not the world rank. They only coincide on a single node.

get is the same with the direction reversed.

Coop scope

ccoCoopThread (default), ccoCoopWarp, ccoCoopBlock. Warp and block are leader-only: lane 0 does the whole copy, the rest return. They generate identical code to thread scope and measure the same — pick whichever matches how your kernel is structured, not for speed.

Use warp/block scope when the surrounding code is warp- or block-collective and you want one op per group without writing the if (lane == 0) yourself.

ccoSdmaThreadMode

Only meaningful for ccoCoopThread (static_assert enforces it), and only when several lanes of one wave call put at once.

Reach for SameQueue first. Lanes sharing a queue post as one group — one reservation, one doorbell — so the issue cost grows very slowly with the number of lanes. Trigger cycles for one warp, one packet per lane:

64 packets cost 43% more to issue than one, so per packet it is 45x cheaper.

Spreading lanes over queues is worse, not better: distinct queues avoid the commit chain but each lane still issues its own uncached reservation and doorbell, and those queue up in the memory pipeline. Queues buy parallelism between warps, not between lanes of one warp.

Use Independent when the shape forces it — one lane per peer in an alltoall, say. Lanes on distinct queues then post in parallel, and any that do share a queue are grouped for you.

Breaking the SameQueue promise makes lanes write slots reserved on another lane's queue. Silent corruption in release; traps under MORI_CCO_SDMA_DEBUG.

optFlags

A bitmask, and a template argument. Two bits:

ccoSdmaOptFlagsAggregate posts without ringing the doorbell — issue N puts, then one commit(peer, queueId) rings the batch. Amortises the doorbell; sweet spot around N = 8–16. A batch must fit in the ring with one unit to spare: at most 256 KB − 32 B on a queue before commit (8191 single-copy puts without signals). The put that would fill the ring waits for the engine, which an uncommitted batch never starts, so it hangs.

ccoSdmaOptFlagsSignalPerCopy gives every copy its own signal packet. You rarely want this. By default a group of lanes signals once, after all of its copies, so expected advances by one per put() call whatever the scope and however many lanes joined. Per-copy signals make it advance by the lane count instead, and cost a packet each — an ATOMIC costs the engine about what a COPY does, so 64 signalled copies take 248 µs instead of 132 µs. Set it only if a receiver has to observe copies landing one at a time. No effect at warp or block scope, which post a single copy either way.

Note the two unrelated "aggregate" names: ccoSdmaOptFlagsAggregate batches doorbells, ccoSdmaThreadSameQueue groups lanes.

Completion

Three ways to know an op is done. All are sender-side: they tell you the engine consumed the command, not that a receiver observed the data.

sdma.put(peer, dst, 0, src, 0, bytes, q);
sdma.quietQueue(peer, q);          // our send is out; src is reusable

quiet/quietQueue take a Coop like put does, and it should match the scope the surrounding code runs at — at warp/block scope only the leader polls, and the call syncs the group before returning, so the completion is visible to everyone.

waitSignal is slightly faster than quiet at every size (0.24 µs at 8B, 0.62 at 64KB) because the signal slot is ordinary device memory while the queue read pointer is not. quiet needs no bookkeeping — prefer it unless the difference matters.

Do not reuse a source buffer before the op completes.

commit only matters with ccoSdmaOptFlagsAggregate: it rings what was posted but not yet rung. quiet drains to the last rung packet, so an aggregate batch that was never committed makes quiet return with nothing sent.

for (int i = 0; i < 8; i++)
  sdma.put<ccoCoopThread, false, false, ccoSdmaOptFlagsAggregate>(
      peer, dst, i * chunk, src, i * chunk, chunk, q);
sdma.commit(peer, q);       // one doorbell for all eight
sdma.quietQueue(peer, q);

Signals

A put can fire trailing atomics in addition to the copy:

• localSignal — increments your own signalBuf[myLsaRank * n + q]
• remoteSignal — increments the peer's signalBuf[myLsaRank * n + q]

Slots are indexed by the sender, so your own local signals and a peer's remote signals land in different slots and never collide.

// sender: notify the peer that the data landed
sdma.put<ccoCoopThread, false, true>(peer, dstWin, off, srcWin, off, bytes, q);

// receiver: no barrier needed — the signal alone proves the data arrived
sdma.waitSignal(senderLsaRank, q, expected);

The ATOMIC follows the COPY on the same queue over the same link, so seeing the signal means the data is there.

waitSignal rules:

• expected is yours to maintain, monotonic, per (srcRank, queueId) pair. Slots are never reset. One running total per pair — a single counter across queues will wait forever.
• One put() call advances the counter by one, no matter how many lanes joined the group, unless the sender set ccoSdmaOptFlagsSignalPerCopy.
• srcRank is an LSA rank. To wait on your own localSignal, pass comm.lsaRank, not comm.rank.
• The comparison is >=, so you cannot miss an increment.

Queues and issuers

The queue is identified by (peer, queueId) — different peers are different queues even at the same queueId.

Keep concurrent issuing warps per queue at or below one. Several warps posting to the same queue is correct but serialises linearly on the in-order commit chain (8 warps on one queue cost ~3.2x the same 8 spread over 8 queues). Raise reqs.sdmaQueueCount so every issuing warp gets its own queue.

This is about warps, not lanes: lanes of one warp should share a queue and use SameQueue. The two combine — a warp aggregates its lanes into one reservation, and each warp owns a queue.

Queues do not make a single transfer faster: one peer is one xGMI link, so splitting a copy across queues gains nothing. They exist to let independent issuers proceed in parallel.

Transfer size

Bandwidth is a function of bytes per op and nothing else. Scope, direction and queue count do not enter into it:

put and get land within 1% of each other, and thread/warp/block are indistinguishable at equal bytes per op. Half of the achievable bandwidth is already gone at 512 KB, and below ~256 KB the fixed ~6 µs dispatch cost dominates whatever the copy itself does.

So the only lever for bandwidth is not splitting the transfer. Two ops of 4 MB reach 57 GB/s where one op of 8 MB reaches 60, even on separate queues — splitting always costs, in proportion to how small the pieces get. Issue the largest contiguous copy you have and let one engine drain it.

Below ~256 KB you are latency-bound, not bandwidth-bound; that is the regime where warp aggregation and SameQueue matter and transfer size does not.

Many small copies

The engine costs about 2 µs per packet, near enough the same for 64 B as for 4 KB. Two consequences.

Issuing is not the bottleneck: 64 packets take 1.3 µs to issue and 128 µs for the engine to run, so SameQueue aggregation buys latency and CU time, not throughput.

Scattering is expensive: 64 x 4 KB costs 138 µs, where the same 256 KB in one copy costs about 11 µs.

So if the data is non-contiguous, packing it first is usually faster: a local gather runs at ~7 TB/s, adding only a couple of percent to the transfer. Scatter directly when the chunks are large (≳1 MB, where the 2 µs is noise), when the CU must stay free for compute, or when you want the data to land in place with no unpack on the receiver.

Engine-side wait and broadcast

Two more packets, with the gfx9 layouts. Validated on gfx950 and gfx1250.

Unlike put, these post from a single issuer and do no lane grouping: thread scope posts one active lane at a time (several lanes are serialized, so they may share a queue), warp/block scope lets the leader post. Only ccoSdmaOptFlagsAggregate is accepted.

postWaitMem / postWaitSignal

// park queue (next, q) until the local flag reads 1; the put behind it then runs
sdma.postWaitMem(next, q, flagWin, 0, 1u);                  // cmp defaults to ccoSdmaCmpEq
sdma.put<ccoCoopThread, false, true>(next, dstWin, 0, srcWin, 0, bytes, q);
// kernel can return here: no CU is polling or has to ring anything later

The engine runs the queue in order, so everything posted after the wait runs only once it holds. win/offset address memory on this GPU (a flag a peer writes with put, a CU store, a host write). cmp is ccoSdmaCmp{Always,Lt,Le,Eq,Ne,Ge,Gt} on (*addr & mask); only 32 bits are compared.

postWaitSignal(peer, q, srcRank, sigQueueId, expected) waits on a signal slot, so one engine can forward what another delivered with no CU in the path:

sdma.postWaitSignal(next, relayQ, prev, hopQ, 1);   // prev's hop landed
sdma.put<ccoCoopThread, false, true>(next, fwdWin, 0, recvWin, 0, bytes, relayQ);

The engine compares 32 bits, so expected must stay below 2^32. A larger value is truncated and releases early: 2^32 + 5 waits for low32 >= 5, which a counter of 100 already meets. Traps under MORI_CCO_SDMA_DEBUG.

Pre-posting everything and ringing once is postWaitMem<.., Aggregate> + put<.., Aggregate> + commit.

Measured on gfx950, flag set to copy complete, median of 100: a CU polling the flag and then issuing a put costs 7.1 µs at 4 KB, 8.5 µs at 64 KB; pre-placing the packets and having the CU only ring costs 5.9 / 7.3 µs; postWaitMem costs 4.4 / 5.8 µs. The gap is a fixed ~1.5–3 µs and is gone by 16 MB. What it really buys is that no CU waits.

Sharp edges: the queue is blocked until the condition holds, so anything else posted to it waits too, and a flag that never arrives parks it for good (retry is unbounded). A full ring spins the poster until the wait clears. Give a waiting queue its own queueId.

postWaitMemPutPtr

// lane p: copy my block to peer p once p has stored this call's epoch in my ready[p];
// the engine waits, the kernel exits right after posting
sdma.postWaitMemPutPtr<false, true>(p, q, ctrlWin, READY_OFF + 4 * p, epoch, ccoSdmaCmpEq,
                                    0xFFFFFFFFu, dstOnPeer, src, bytes);

postWaitMem<Aggregate> + putPtr<Aggregate> + commit in one reservation and one doorbell: cheaper to post, same flag-to-copy latency. Thread scope; lanes on distinct queues post in parallel. dst is the peer's memory as mapped here (lsa_ptr), win/offset the local flag; bytes == 0 posts wait + signals, and nothing without a signal. Over CCO_SDMA_MAX_COPY_BYTES (~1 GiB) it posts wait + chunked put + commit, one lane at a time, so on a shared queue each lane's copy waits only for its own flag and those of lanes posted before it. Use it when the flag is usually not there yet; a wait that already holds still costs ~1.6 µs over a plain put, so check the flag on the CU first if it usually is: read it once, plain put if it already holds, this post if not (example 12). Use Eq on a per-call epoch: the engine compares unsigned, so Ge on a running total misbehaves across the u32 wrap. The postWaitMem caveats apply.

putBroadcast

sdma.putBroadcast(next, recvWin, 0,  /*peer2=*/me, localCopyWin, 0,  sendWin, 0, bytes, q);

One COPY packet, two destinations; peer/dstWin is the first and selects the queue, peer2 may be your own rank. remoteSignal bumps each distinct destination peer other than yourself once; localSignal is one atomic for the call.

The two destination addresses must be congruent mod 32 ((dst1 ^ dst2) & 31 == 0), as in ROCr's broadcast packet. If not, the second destination receives the data shifted by the offset difference: the whole range is wrong and up to 31 bytes next to it are overwritten. Nothing checks this unless you build with MORI_CCO_SDMA_DEBUG, which traps.

A signal slot counts per (source rank, queueId), not per physical queue. The broadcast is posted on queue (peer, q) but its remote signal bumps peer2's [you, q] slot, which put only advances from queue (peer2, q). A waiter that reads the count as "the Nth packet on (peer2, q) has landed" can be released by a broadcast's signal while an earlier put on (peer2, q) is still in flight; only the cumulative total is reliable. If you wait on per-queue completion, give broadcasts their own queueId.

Do not reach for it to save time. One engine does not overlap the two writes: broadcast to two remote GPUs ran at 38 GB/s against 61 GB/s for two puts on one queue (122 GB/s for two puts on two queues), and local + remote was the same 43 vs 61. It was faster only when both destinations are local (68 vs 61 GB/s). For fan-out to peers, issue puts on separate queues.

The packet is 36 B, so it takes two ring units plus one of slack that is padded with a NOP, so the packet is never split across the ring end.

Failure modes

Build with -DMORI_CCO_SDMA_DEBUG when any of these appear: it turns the misuse checks into traps at the offending call instead of a hang or corruption.

Worked example

Each lane of the warp ships its own chunk to the same peer — the shape SameQueue is for.

__global__ void ship(ccoWindowDevice* sendWin, ccoWindowDevice* recvWin, size_t chunk,
                     ccoDevComm devComm, int peer) {
  ccoSdma sdma{devComm};
  const int lane = threadIdx.x;                 // one queue per warp
  const int q = (threadIdx.x / 64) % devComm.sdma.sdmaNumQueue;

  sdma.put<ccoCoopThread, false, false, ccoSdmaOptFlagsDefault, ccoSdmaThreadSameQueue>(
      peer, reinterpret_cast<ccoWindow_t>(recvWin), lane * chunk,
      reinterpret_cast<ccoWindow_t>(sendWin), lane * chunk, chunk, q);

  sdma.quietQueue(peer, q);
}

Every lane of the warp has the same peer and q, so the promise holds: one reservation for the whole warp, one doorbell.

When the shape forces Independent

One lane per peer, one signal to the peer, no barrier on the receive side.

__global__ void alltoall(ccoWindowDevice* sendWin, ccoWindowDevice* recvWin,
                         size_t bytes, ccoDevComm devComm, uint64_t expected) {
  ccoSdma sdma{devComm};
  const int me = devComm.lsaRank, n = devComm.lsaSize;
  const int p = threadIdx.x;
  if (p >= n || p == me) return;

  // lanes target distinct peers -> distinct queues -> posted in parallel
  sdma.put<ccoCoopThread, false, /*remoteSignal=*/true>(
      p, reinterpret_cast<ccoWindow_t>(recvWin), me * bytes,
      reinterpret_cast<ccoWindow_t>(sendWin), p * bytes, bytes, /*queueId=*/0);

  sdma.quiet(p);                      // our sends are out
  sdma.waitSignal(p, 0, expected);    // peer p's data has landed
}