QLink - Distributed Communication for QwenGPUs

What is QwenGPU?

QwenGPU is a project that was recently developed by vixhal-baraiya inspired by Qwenputer. The original Qwenputer project shows how to compile and insert a 32-bit RISC-V CPU directly into the weights of Qwen3.5-4B, leveraging the Gated DeltaNet (GDN) layers of the model as the RAM for the CPU. The GDN layers carry a recurrent state that can be updated one row at a time via the delta rule, and this enables each row to function as a memory cell. QwenGPU extends this idea to GPUs and inserts a SIMT execution model directly into the Qwen weights.

Specifically, QwenGPU adds 45 extra GDN and SwiGLU layers in front of Qwen3.5-2B's own layers, and the weights of these layers are computed by a compiler instead of being learned. The GPU itself is pretty small. It has 4 SIMT lanes, 15 registers per lane, a 32-block program ROM, and 127 words of global memory, where each 32-bit word is one row of one head's state. To run a kernel, the host compiles the CUDA code to PTX, lowers it to machine blocks, and feeds the program to the model as tokens. After that, the model writes the launch on its own with greedy decoding, one record per warp per block. Each record is an EXEC token followed by the register values that the block produced (two 16-bit tokens per lane). Every value in the circuit can be represented in bf16, so the model reproduces a reference interpreter exactly, token for token.

QLink: Connecting the QwenGPUs

My idea with this project is relatively simple: how do we enable these QwenGPUs to talk with each other in the same fashion as how modern GPUs perform distributed operations for neural network training and inference. So, I went and asked my current favorite LLM (Opus 5.5) to help me build QLink, a distributed communication protocol that behaves like NCCL / NVLink. The rest of this blog will go over some of the interesting design decisions made by Claude to build QLink along with results on training real neural networks using multiple QwenGPUs.

Designing QLink

The goal with QLink is to build an interconnect that behaves like NVLink. The way this usually works is that a kernel on one GPU will store to an address in another GPU's memory, and the word shows up in another GPU's memory while that GPU's own kernel is still running. The one design principle that was maintained from the beginning is that the host can launch kernels and route packets, but it never computes, reduces, or decides a value. This means that every number has to come out of the models' weights directly.

Using Store Records as Packets

First, we need to decide what a message between two QwenGPUs looks like. In QwenGPU, the host writes a word of memory by feeding the model three tokens: gpu_mem a gpu=lo gpu=hi. When a kernel running on the model stores to global memory, the model also writes these three tokens itself as a write-back record, right after the warp's register values. Claude noticed that this makes a store from one model a valid memory write for any other model. If we feed that store into a second model's context, the second model's memory head writes the word the same way it does for a cudaMemcpy from the host, so we can use the store record as the packet format without making up a new one.

The other way to do this would be to add new tokens or layers for network messages. We didn't want to do that because any change to the compiled circuit means checking its exactness again from scratch. Since QLink only reuses the store record, it adds no tokens or weights to QwenGPU, nothing has to be trained, and every guarantee QwenGPU makes about exactness still holds for QLink.

Routing Stores with Apertures

Next, each device needs a way to decide which stores should leave. Claude based this on how NVLink maps a peer's memory into a GPU's address space. Each device keeps a table of address ranges (we call these apertures) that belong to its peers. When the model writes gpu_mem a for an address inside an aperture, the device forwards the record to the peer and feeds itself gpu_mem 127 instead. The router only ever looks at the header token, and the payload is passed through as is.

I think the sink word is one of the more subtle details in the design. By the time the router sees the header, the model has already predicted gpu_mem a, so the device has to feed some header back into the model's context. It can't feed back a itself, since then the store would also overwrite the local word. Word 127 is a row that the driver never uses, so sending the local copy there keeps every record well-formed without touching local memory. The one exception is a multicast aperture, which keeps the local write and also sends the store to every peer.

Loads are never routed, so reading an aperture address still reads local memory. This was an intentional choice, since a remote read would need the peer to compute a reply and send it back, which QLink doesn't support, and as we will see later, local loads are what make qccl's channels work. Atomics into an aperture are refused for the same reason, since a remote atomic would need the peer to do the addition.

When Can a Packet Arrive?

The last piece of the link is deciding when a packet is allowed to enter a model's context. We can't feed a packet in the middle of a record, because the circuit counts the VAL tokens to decide which register each one commits to. However, there is a window between two records that works. After the model has predicted the next EXEC but before that token is fed back, the record counter is already past every commit slot. A packet fed in at this point can only write a memory row, and the next record reads memory once its EXEC arrives. This gives a kernel the memory semantics that it needs. For example, a kernel spinning on a flag will see a peer's store show up on its next iteration, which is the same thing a CUDA kernel sees with an NVLink store.

This holds in the compiled circuit (evaluated exactly) and also in the bf16 model running on a GB200. In both cases, the token streams with packets spliced in are identical to a reference interpreter that applies the same writes between records.

Running in Lockstep

Real GPUs run asynchronously, but Claude decided to run the QwenGPUs in lockstep, one record per tick. On every tick, each busy device writes one record, the packets from those records are exchanged, and each device feeds in what it received at its next record boundary. Since the schedule only depends on what the devices write, every run is deterministic, and the reference interpreter, the exact circuit, and the bf16 model all have to produce the same token streams. This is what makes it possible to check every single token. Between processes, the packets are sent with torch.distributed over a CPU (gloo) group, so an exchange never has to wait behind the model's work on the GPU.

Note that lockstep is only a property of the simulation and not of the protocols. qccl's kernels only ever wait on flags and never on timing, which is why the same kernels also run correctly on real GPUs that are asynchronous. For the longer runs, Claude later added a burst exchange mode where each device writes several records between exchanges. This adds a little latency but needs far fewer exchanges, and the schedule is still deterministic.

qccl: NCCL for QwenGPUs

Once stores can move between models, the rest of the problem is basically implementing NCCL. For this, Claude wrote qccl, a small communicator library that follows NCCL's design closely.

Channels on a Symmetric Heap

Each rank gets a ring channel to the next rank, and the channel is laid out at the same addresses on every rank (i.e. a symmetric heap, like in NVSHMEM). A channel has a FIFO of slots, a tail word that counts how many chunks have arrived, and a head word that counts how many of this rank's chunks the next rank has consumed.

The decision to only route stores is what makes this layout work. Since loads always read local memory, the send side of a channel can sit on the same addresses as the receive side. A kernel loads from a FIFO slot to get what the previous rank sent, and its stores to that same slot land at the next rank, while the sink keeps them from overwriting the local copy. So one set of addresses works as both the receive buffer and the send window, which matters a lot when the whole GPU only has 127 words.

NCCL's Simple Protocol

For flow control, Claude went with NCCL's Simple protocol. In this protocol the sender pushes data directly to the receiver and all of the synchronization happens through two flags, which works well for a link that can only carry stores. The sender waits for a free slot (head), writes its chunk into the receiver's FIFO, issues a fence, and then posts the receiver's tail. The receiver waits for tail, consumes the slot, and posts the sender's head. The whole protocol is one CUDA kernel that gets launched with a different mode at every step:

// kernels/qccl/prim.cu, compiled by clang to PTX, run by the model
if (mode & RECV) while (*rtail < rwait) {}          // the chunk has arrived
if ((mode & SEND) && swait > 0) while (*rhead < swait) {}  // a slot is free
int r = 0, v = 0;
if (mode & RECV)  r = rslot[i];
if (mode & LOCAL) v = x[i];
if (!(mode & SPLIT)) v += r;
if (mode & COPY)  y[i] = (mode & SPLIT) ? r : v;
if (mode & SEND)  sslot[i] = v;                     // these stores cross QLink
__threadfence_system(); __syncthreads();
if (i != 0) return;
if (mode & SEND) *stail = spost;                   // -> next rank: data is there
if (mode & RECV) *shead = rpost;                   // -> previous rank: slot is free

Each collective is a sequence of these steps, in the same order that NCCL's runRing uses. For example, AllReduce on N ranks is a send, then N−2 recvReduceSend steps, a recvReduceCopySend, N−2 recvCopySend steps, and finally a recv. The host only enqueues the launches and keeps track of the step counters. It never touches the data, and there is no host code that makes one rank wait for another. The kernels handle that on their own by polling the flags that their peers' stores write.

Why Does the FIFO Need Two Slots?

The depth of the FIFO turned out to be important. With only one slot, every rank waits for its successor's free-slot flag, but that flag only gets posted at the end of the successor's own receive-and-send launch, so the whole ring deadlocks. This was caught by a test on the reference interpreter, and it is the reason qccl uses two slots. The order of operations matters in a similar way. If every send is issued before any receive, a 2-slot FIFO also deadlocks, and this is true on QwenGPUs just like it is with NCCL. To avoid this, qccl software-pipelines SendRecv like NCCL's grouped calls do, and there is a test that checks that the naive ordering really does deadlock.

Changes to the Driver

To compile these kernels, QwenGPU's driver needed a few extensions. First, pointer arguments can now be fixed device addresses that stay resident across launches, so parameters and FIFOs live in memory like they would in HBM. Second, loads from warp-uniform addresses are treated as uniform, so a spin loop like while (*flag < k) compiles to a plain loop instead of forcing the whole kernel to be predicated. Third, an early exit after a uniform loop is if-converted, which lets a single thread post a flag. Finally, the reference interpreter now produces one record at a time so that packets from peers can land between records.

I think the most interesting change is how memory ordering is handled. ld.volatile and acquire/release accesses are lowered to plain accesses, and fences are lowered to no-ops. On a normal GPU this would not be safe, but on this machine it is fine because only one warp runs at a time and stores leave in program order over FIFO links. So the ordering that a fence would enforce is already guaranteed.

Running the Same CUDA on Real GPUs

Since the kernels are just ordinary CUDA, Claude also ran them on ordinary GPUs. The same PTX and the same host schedule run on four GB200s, where an aperture address becomes the peer's own pointer and the stores go over NVLink, similar to NCCL's P2P transport. All of the launches are asynchronous, and the only synchronization is through the flags.

Interestingly, this version found a bug that the lockstep simulation never could have found. Loading a CUDA module allocates device memory, and with peer access enabled this can end up waiting for the peers. Pinning host memory for a copy also synchronizes the device. If either of these happens while a peer's kernel is spinning on a flag that this GPU hasn't written yet, the whole system hangs. The fix was to load every kernel before anything runs, and to issue the data loader's copies before the kernels that spin.

Does It Match NCCL?

For the full test, we run four merged Qwen3.5-2B models, one per GB200 with one process each. Each model decodes greedily, one token per forward pass, and every token is checked against the reference as it gets written. We then compare each collective's result with what real NCCL computes on the same four GPUs from the same int32 inputs.

The design also isn't limited to four ranks. With two merged models on each GB200, eight QwenGPUs run the same six collectives correctly, and an 8-rank AllReduce takes 15 launches and 1,480 written tokens per rank. NCCL can't place two ranks on one GPU, so numpy is used as the reference for this case. On the real GB200s over NVLink, all six collectives are also correct after 88 asynchronous launches.

Checking Every Token

Claude designed the system so that every result could be reproduced exactly, and every number in this post is verified in three ways. First, a reference interpreter runs the machine blocks record by record. Next, the compiled circuit, evaluated exactly in float64, has to write the same tokens. Finally, the bf16 model running on a real GPU also has to write them. All of the training arithmetic is done in integers, so each training run also has an exact numpy replay, and the parameters held by the models have to match it bit for bit after every step.

For the longer runs, Claude added a verify mode. In verify mode, the reference drives the system, with real packets moving between processes. At regular intervals, each model runs the exact context it would have seen in one teacher-forced pass, and its argmax has to equal the token at every position it owns. If this holds for every model at every position, then the greedy system would have produced exactly the same streams. The argument is an induction over ticks, since the first place where any model could deviate is a position that was already checked. QwenGPU's own merged_check relies on the same argument. In the first training run, verify mode checked about 410,000 positions per model, and a two-step greedy run (with every token written autoregressively) reproduced it exactly.

Long runs also have one more problem, which is that a context keeps growing as the model writes. The fix for this is a context rebase between launches. The host reads the 127 words back, starts a fresh context, and writes them back in. This is just a state save and restore, so the host copies words but still doesn't compute anything. The point-cloud diffusion run below used about 260 rebases per rank.

Training with Only 127 Words

Training adds its own constraint, which is that the parameters, gradients, minibatch, and communication channel all have to fit in 127 words per GPU. Every model in this section uses Q8 fixed point (1.0 = 256) with round-to-nearest shifts, so every run can be replayed exactly in numpy.

Data Parallel Training on XOR

A data-parallel training step is a forward pass, a backward pass, an AllReduce, and a parameter update, and the QwenGPUs run all four of these as CUDA kernels. The first network we trained is a 2-8-1 ReLU MLP with 32 parameters on XOR, which is a classic problem that a linear model can't learn. Each of the four ranks holds a replica of the parameters and a quarter of the training set, and samples 8 points per step. Before step 0, rank 0 initializes the parameters and a Broadcast gives every rank the same replica, like DDP does at construction.

H2D        the rank's minibatch                       // its data loader
forward    mlp_forward   o[s] += W2[k] relu(W1[k] x[s] + b1[k])   // atomicAdd
backward   mlp_backward  per-sample gradients, summed by atomicAdd
comm       qccl ring AllReduce(g)                     // over QLink
update     sgd           p -= round(g / 2^6)

When the same program is compiled for the GB200s and run over NVLink, it ends up with the same 32 parameters, bit for bit.

Multicast and ZeRO

On a QwenGPU, every token the model writes costs a forward pass, but every token that is fed in (the host's records or a peer's packets) is prefill and is nearly free. So the cost of a collective is really the number of tokens each rank has to write. With this cost model, a ring isn't the best choice, since it moves every word through every rank. GB200's NVLink switch can multicast a store to every GPU (NVLink SHARP's multimem.st), and QLink can do the same thing with a multicast aperture. With multicast, each word only needs to be written once.

With the parameters in a multicast window, the training step turns into ZeRO-1. Each rank ReduceScatters the gradients, updates only its quarter of the parameters, and lets those stores land in every replica. This gives the same parameters at every step while cutting the communication tokens in half. This version of the XOR run uses 119 of the 127 words: 18 for the channel, 32 each for the parameters, gradients, and batch, 4 for flags, and 1 for the loss.

SHARP can also reduce inside the network (multimem.ld_reduce), but that would mean the fabric is adding numbers, which goes against the rule that nothing outside the models computes. So all of the reductions stay in the GPUs' weights.

Tensor Parallelism for MNIST

MNIST doesn't fit in 127 words, since even a linear classifier on a 5×5 thumbnail needs 260 weights. Instead, Claude split the model across the four QwenGPUs the same way Megatron splits a RowParallelLinear layer. Each image becomes 32 input features, and each rank holds 8 of them along with the 80 weights that read them. In a training step with one image, every rank computes its partial logits, a qccl AllReduce sums them (this is the only communication in the layer), and then each rank updates its own 80 weights using only local data.

The first version fed the classifier the central 20×20 pixels, average-pooled to 5×5, and trained it with least squares. This reached 80.6% test accuracy, which is close to the best that least squares can do on those 25 pixels (80.9%). I then asked if we could get the accuracy to be even higher with some modeling changes, and Claude was able to make two changes that improve accuracy without using any extra memory. First, the data loader's fixed encoder now keeps the top 31 whitened principal components of the image, which preserve much more of the digit than 25 pooled pixels. Second, the update uses a one-vs-rest squared hinge loss instead of least squares, so that once a class is past its margin, it stops contributing to the update. This allows the classifier to spend its 320 weights on the images it gets wrong. On the QwenGPUs this is one small kernel, and everything else stays the same.

Diffusion on a Point Cloud

After MNIST, I wanted to see if we could train a diffusion model. Claude started with a point cloud as a warm-up, where the training data are points sampled from the glyph "7". The model is a DDPM with 20 noise levels, and the denoiser predicts the noise in a point given the point and the noise level. The denoiser is an MLP with 32 hidden units, split with Megatron style TP. Each rank holds 8 hidden units, and one AllReduce combines their outputs.

There are two small choices here that make this split cheap. First, rank 0 starts its partial output at the negative of the true noise, so the output of the AllReduce is already the error on every rank. Second, because of this the backward pass doesn't need any communication, since every rank has the full error and each rank's gradient only involves its own hidden units. The data loader samples a point, a noise level, and the noise, and everything after that (the forward pass, backward pass, and updates) runs on the QwenGPUs in Q8 fixed point. Sampling runs on the QwenGPUs too. Each of the 20 DDIM steps is a forward pass of the split denoiser, an AllReduce, and a fixed-point update kernel, applied to eight points at a time.

The result is a blurry 7, but most of the blur comes from how small the network is. If we train the same 32-unit denoiser in float64 on the same batches with the same schedule, it is only slightly sharper, with a two-sided average distance to the glyph of 0.090 compared to 0.101 for the fixed-point model. So fixed point costs a little bit of sample quality, but running on the QwenGPUs doesn't add any error on top of that, since their result matches the fixed-point arithmetic bit for bit.

Latent Diffusion for Digits

I also wanted to try training a real image model on this setup, even if it was a bad one. A 28×28 digit is 784 numbers, which is six times QwenGPU's entire memory, so this model follows the approach of Stable Diffusion and runs diffusion in a latent space. An encoder compresses each image into 8 numbers, a decoder maps 8 numbers back to an image, and the diffusion model only ever sees the 8 numbers. The encoder computes the top 8 principal components of the training images, whitened, with a sign convention so that any linear algebra library computes the same encoder and the numpy replay stays exact. The decoder is a small MLP (8 → 128 → 784) trained on a normal GPU. Neither the encoder nor the decoder runs on the QwenGPUs.

The QwenGPUs learn the diffusion model itself. The denoiser is a 9 → 16 → 8 MLP (8 latent numbers and the noise level as input, and the predicted noise as output). It is split over the four QwenGPUs like before, with four hidden units and 80 words of weights on each one. To save on kernels, each bias is stored as an extra row of its weight matrix, so one update kernel can handle both. A training step processes one image and consists of a forward pass, an AllReduce, and three backward and update kernels, which takes about 4,000 tokens per QwenGPU. Sampling is 20 DDIM steps of the split denoiser, one digit at a time.

One nice side effect of the exact numpy replay is that it predicts the QwenGPU run ahead of time. Claude used this to pick the training seed, by scoring the replays of a few candidate seeds before launching the real run. With the MLP decoder, an MNIST classifier recognizes 76% of the samples.

Try It Yourself

The trained denoiser is only 296 integers, so it ships with this page along with the decoder, and you can run it directly in your browser. It uses the same int32 arithmetic as the QwenGPUs' kernels, and runs as a WebGPU compute shader if your browser supports WebGPU (and in JavaScript otherwise). The slider picks the weights at any checkpoint, from untrained to final. Replaying the QwenGPUs' noise also works as a check, since the integers that come out have to match the ones the QwenGPUs wrote.

As you can see, the generated digits are far from perfect. A lot of them are recognizable, but there are also samples that look like a mix of two digits or that don't look like a digit at all. This is expected given how small the model is, since both the denoiser (16 hidden units) and the latent (8 numbers) are tiny and the decoder can only sharpen what the latent keeps. It's still pretty cool that these digits come from a diffusion model that was trained by four language models talking to each other over QLink. Perhaps language is all we need 😉.

Limitations and Conclusion

This system is ridiculously slow to train any legit model. A 16-word AllReduce takes 664 model-written tokens per rank, and one training step of the XOR network takes about ten thousand. However, most of this cost comes from the GPU's narrow design (4 lanes and 4 stages per record) and not from the interconnect. Sending a word costs the sender exactly what storing it locally would (three tokens), and the receiver reads them as prefill. Greedy decoding is also dominated by host overhead, and CUDA graphs would remove most of it, so there's definitely a lot of room for optimization here.

Something I found interesting about this project is that most of the pieces QLink needed were already in QwenGPU. The packet format, the memory write, and the exactness guarantees all came with the original GPU, and QLink itself is mostly a routing table and a rule for when tokens can be spliced into a context. qccl also follows NCCL's design almost line for line.

I'm also surprised (yet again) at how effective Opus 5.5 is at doing this kind of exploratory work extremely fast and efficiently. I was able to just point Opus to the original QwenGPU repository, have a short discussion on what I intended the project to look like, and it helped build QLink and run all the experiments within a day. It's very exciting to see the crazy projects and ideas that are now within reach as the marginal cost of research and engineering keeps dropping.

@software{agrawal2026qwenputer,
  author  = {Agrawal, Divyansh},
  title   = {{Qwenputer}: a computer compiled into the weights of a pretrained language model},
  year    = {2026}, version = {0.1.0},
  url     = {https://github.com/divagr18/qwenputer}, license = {MIT}
}