RunAnywhere opens Trainium2 engine for GLM-5.3-Flash
RunAnywhere developer said the company open-sourced an inference engine that serves the 320B-parameter, 1M-context GLM-5.3-Flash model on one AWS Trainium2 trn2.48xlarge. The project uses fused NKI layer megakernels to reduce a batch-8 decode step from 32.2 ms to 24.2 ms on the author's layout, with reproduction instructions and detailed profiles included. RunAnywhere also reported up to 28.9 agent requests per minute on the self-run SemiAnalysis AgentX benchmark, but the post provides no stock-AWS or GPU baseline.
RunAnywhere's open-source engine and profiling data provide a detailed reference for serving a large MoE model on AWS Trainium2, while its self-reported results still require comparison against stock AWS and GPU deployments.
We wrote an inference engine for AWS Trainium2 that serves GLM-5.3-Flash (320B MoE, 1M-token context) on one trn2.48xlarge, and open-sourced it. Most of the work went into decode, and the part I found most interesting is how little of a decode step was the model's actual math.
In an early build, in one KDA (linear attention) layer at batch 1, the recurrence was 2% of the attention half. The rest was the mHC pre-projection (36%), XLA glue (33%), transposes (10%), the stream mix (10%) and other small ops (8%). A batch-8 decode step took 24.2 ms against a 6.1 ms memory floor, and the profile showed why: long chains of small vector-engine instructions at about 153 ns of fixed cost each, a write back to HBM at every kernel boundary, and 125 cross-chip collectives per step.
So we rewrote each layer half as one NKI kernel (each whole layer, for the sparse-attention layers). The kernels follow three rules:
- Keep the working set on chip. Each physical NeuronCore has 28 MiB of SBUF, enough for a layer half's activations, its recurrent state and the weights in flight, so only the layer's inputs, outputs and streamed weights cross HBM.
- Use both physical cores of a logical core. One program streams 340–400 GB/s of weights, the pair 672–708 GB/s, so the weight stream has to be split between them.
- Move a collective into the kernel when it blocks the next one. Run by XLA, an all-reduce makes the next kernel wait for all of it. Inside the kernel, independent work can overlap it.
#RunAnywhere#GLM-5.3-Flash#AWS Trainium2#trn2.48xlarge#NKI#XLA#SemiAnalysis AgentX#vLLM Neuron#NxD Inference#AMD#Hazy Research#Mirage