In 2018 I hand-wrote a C++ deep learning framework so I'd never pad a batch. In 2023 LLM serving landed on the same structure.

A developer who built InsNet, a ~21,000-line C++14 deep learning library with 4,828 lines of hand-written CUDA kernels, starting in 2018, describes how its padding-free dynamic batching design anticipated the flat, padding-free token streams with per-sequence offsets now used by modern LLM serving engines like vLLM. The library represents each value as a flat buffer with only element count and width, eliminating batch dimensions, max lengths and mask tensors, a bet the author says was early but aimed at the wrong layer of the stack.

There's a sentence in the README of a library I wrote that I've been thinking about lately: "To summarize, we believe that Padding-free Dynamic Batching is the feature that NLPers will dive into but is surprisingly not supported by today's deep learning libraries ." I wrote that around 2021, about InsNet https://github.com/chncwang/InsNet , a C++14 deep learning library I'd been building since 2018. Then transformers ate the field, everyone padded their batches to rectangles like they always had, and I moved on. Two years later, vLLM launched, "continuous batching" became the load-bearing idea of the entire LLM serving industry, and the input to a modern inference engine became — a flat, padding-free token stream with per-sequence offsets riding alongside as data. I wasn't wrong. I was early, and I was aiming at the wrong layer of the stack. This post is the story of that bet: how a C++ library made padding disappear, how the canonical prior art DyNet made a subtly different choice at the same fork, and how the modern serving stack vLLM + FlashAttention ended up rediscovering the same trick — with receipts from all three codebases, because last time I compared engines from memory people rightly asked for sources, and reading the actual code is where all the good surprises live anyway. In 2018 I was a master's student doing NLP research — the last years before transformers swallowed the field. The workhorses were still RNNs and LSTMs, and the frontier I found interesting was the models whose computation graph changed shape with every single input : tree-LSTMs folded along a sentence's parse tree, transition-based parsers emitting a different sequence of stack operations for each sentence, hierarchical encoders running one sub-model per sentence and another over the document. Two examples in a batch almost never had the same shape — one had 7 tokens, its neighbor had 212, and their tree structures didn't line up at all. Tree-LSTM, transition-based parser, hierarchical encoder — three inputs, three graph shapes, no shared rectangle. The standard answer was padding: extend everything to the longest sequence in the batch, add a mask tensor, and burn FLOPs computing values you'd immediately multiply by zero. For flat same-ish-length batches this is a mild tax. For instance-dependent structures it's obscene — and worse than the wasted compute was the wasted thinking : every model became two models, the one you meant and the one that handles the mask. I wanted to write the model for one instance and have the library figure out the batching. So I wrote a library: InsNet. It grew out of N3LDG — an earlier dynamic-computation-graph NLP library I helped build and first-authored the 2019 paper for — reworked into about 21,000 lines of first-party C++, of which 4,828 lines are one file of hand-written CUDA kernels. InsNet's core representation makes padding impossible rather than optional. A value's data lives in one flat buffer, and its shape is just two integers: the total element count and the width. In an NLP model every value is a 2-D matrix — d rows the hidden size by some number of columns, and only the column count ever changes. A transformer holds a whole sentence at once, so a value is a d×L matrix, one column per token; an RNN steps through the sentence one token at a time, so each value is a single column — a d×1 vector. Same d rows either way; the width is the only thing that moves, which is exactly why two integers pin the whole shape. There isn't even a stored length — divide the buffer's size by the hidden dimension and the token count falls out. A value is never a batch, L max, d slice with a mask; it's a matrix as wide as the work in front of it. No batch dimension anywhere in the type system, no max length, no mask tensor. You cannot pad because there is nothing to pad to . Padding stretches every sentence to the longest and burns compute on the gaps 29% here ; InsNet keeps each value exactly as wide as its sentence. The interesting part is how batching happens with no batch dimension. InsNet uses lazy execution: your model code builds a graph of small nodes, and nothing runs until you call forward . At that point the executor repeatedly takes the current wave of ready nodes Kahn's algorithm — every node whose inputs are all computed and buckets the wave by a type signature : // graph.h — the ready set, bucketed by signature typedef std::unordered map<std::string, std::vector<NodeAbs NodeMap; NodeMap free nodes; Each bucket becomes one batch, executed by one operator call. The ready wave, grouped by signature: one signature per op call — the two Linear·W₁ nodes together, the two Linear·W₂ next, while the pending Adds wait. And the whole design lives or dies on one question: what goes into the signature? Here's InsNet's signature for a linear layer: js // operator/linear.cc string typeSignature const override { return Node::getNodeType + "-" + addressToString param ; } Operator type, plus the address of the weight matrix . That's it. Not the input width, not the number of columns. Two linear nodes batch together if they apply the same weights — even if one is transforming 3 columns and the other 300. The column count, the part that varies per instance, is deliberately left out of the key, so a single "batch" in InsNet genuinely contains matrices of different shapes. Why width doesn't matter to a linear layer: W hits each column on its own, so a 3-column and a 300-column input concatenate into one GEMM. This is the fork in the road, and to see why it matters you have to look at what the grown-ups did. Dynamic batching was not my idea. DyNet — the library from CMU that powered a lot of 2016–2018 NLP research — shipped on-the-fly autobatching in 2017 Neubig et al., "On-the-fly Operation Batching in Dynamic Computation Graphs" , and its implementation is beautiful. Same skeleton as InsNet: build the graph lazily, compute a signature per node, batch same-signature ready nodes. But look at the signature for the workhorse op, affine transform b + W x : // dynet/nodes-affinetransform.cc Sig s nt::affine ; s.add node args 0 ; // identity of b for size t i = 1; i < args.size ; i += 2 { s.add node args i ; // identity of W s.add dim cg.nodes args i+1 - dim ; // SHAPE of x <-- the fork } The fork: DyNet's key includes shape → same-shape buckets; InsNet's key omits width → one batch of mixed shapes, and the raggedness moves into the kernel. The shape of the data operand is part of the key . Only identically-shaped operations ever share a batch. That single decision shapes everything downstream: // dynet/exec.cc — check contiguity, else memcpy into fresh memory if contig { // use current mem, zero copy my xsi- v = min node; my batch.concat i = 2; } else { // gather into new mem combine tensors my batch.ids, i, my xsi ; } DyNet even ships a wonderful piece of engineering honesty: pass --dynet-autobatch 100 and it benchmarks its own three scheduling strategies on your graph and keeps the fastest — a built-in admission that the scheduler itself has a cost worth measuring. So: same problem, same lazy-graph skeleton, and one bit of difference in the key. DyNet says a batch is same-shaped things , and keeps its kernels boring. InsNet says a batch is same-operation things , and pays for it one level down — because now the kernels have to handle a batch of differently-shaped matrices. Why the two defaults differ comes down to where each era put the variable length . DyNet grew up in the RNN age, where a sentence's length becomes a variable number of steps — but each step's value is still a fixed d×1 vector. The length lives in the node count, not the node shape, so shape stays constant and keying on it costs nothing. Transformers move that length into the width of a value: one d×L matrix per sentence, L varying by input. Now the length lives in the shape, so keying on shape fragments the batch by length — two sentences of different length can't share a sequence-level op. InsNet drops the column count to undo exactly that: it keys each op on the fixed part — the weight matrix, or the attention head dimension — and lets the width vary, so a length-7 and a length-12 sentence batch together right through the attention matmuls, not just the linear layers. Same fork, opposite defaults — each matched to where its era hid the raggedness. Which brings us to the part of InsNet I'm simultaneously proudest and most embarrassed of: the CUDA. Here's the batched matmul kernel's signature and how it handles the ragged batch: // cuda/impl.cu — one launch, a whole ragged batch global void KernelMatMul dtype a, bool transpose a, dtype b, bool transpose b, int a rows, int b cols, int ks, dtype vals, bool acc, bool use lower triangle mask = false { int count i = blockIdx.x; // which instance am I? int b col = b cols count i ; // THIS instance's width int b col i = blockDim.x blockIdx.z + threadIdx.x; if b col i = b col return; // past my instance's edge: do nothing // ... then accumulate this cell's dot product over ks count i ... } Look at the parameter list. Not tensors — arrays of per-instance pointers dtype a and arrays of per-instance sizes int a rows, int b cols, int ks . The launch grid is rectangular, sized to the largest matrix in the batch; every thread first looks up which instance it belongs to and its instance's true dimensions, and threads that fall past their instance's real edge just return. One kernel launch, one batch, N different shapes — the raggedness rides in the size arrays, not in padded buffers, so the kernel masks threads instead of data. No thread ever computes a pad cell, and the real payoff isn't the memory saved so much as the compute never spent on values you'd only multiply by zero. One rectangular launch over a ragged batch — the b cols size array masks the threads past each instance's real width, no padded memory, only skipped threads. The causal attention mask is baked into the same kernel: use lower triangle mask writes −1e30 — its finite stand-in for −∞ — above the diagonal. In 2018 that felt like a hack. In 2024 fusing the mask into the attention kernel is called "writing an attention kernel." The implementation is rough — a scalar accumulation loop, no tiling, no cuBLAS on this path cuBLAS wants uniform shapes, and a ragged, masked batch isn't that . One person, nights and weekends, hand-rolling ragged-batch GEMMs because the alternative was padding. You make trade-offs. One thing I want to defend properly: this wasn't an inference trick. InsNet trains. The executor records batches in the order it ran them, and backprop just replays the recording backwards: // graph.cc — backward = the forward tape, reversed for int idx = count - 1; idx = 0; --idx execs.at idx - backwardFully ; Because forward executed in topological waves, the reversed tape is a valid reverse-topological schedule, and each backward step is batched exactly like its forward step was — the linear layer's backward is two batched GEMMs weight grad and input grad over the same concatenated columns. The padding-free property propagates through training for free: no masked loss terms, no gradient contributions from pad tokens, because pad tokens don't exist. Now put InsNet down and look at what LLM serving converged on. vLLM's scheduler — the "continuous batching" everyone talks about — recomposes its batch on every single step . The v1 scheduler's own docstring is admirably blunt: vllm/v1/core/sched/scheduler.py NOTE woosuk on the scheduling algorithm: There's no "decoding phase" nor "prefill phase" in the scheduler. Each request just has the num computed tokens and num tokens with spec. At each step, the scheduler tries to assign tokens to the requests so that each request's num computed tokens can catch up ... Requests join the running batch mid-stream, advance by whatever token budget allows, and leave the moment they finish. And the tensor that actually enters the model? A comment in the input-preparation code draws you the picture: php vllm/v1/worker/gpu model runner.py E.g., 2, 5, 3 - 0, 0, 1, 1, 1, 1, 1, 2, 2, 2 req indices = np.repeat self.arange np :num reqs , num scheduled tokens cu num tokens: 2, 5, 3 - 2, 7, 10 Three requests with 2, 5, and 3 tokens this step become one flat array of 10 tokens . Not a 3, 5 matrix with padding — a length-10 stream, with the boundaries 2, 7, 10 carried alongside as data. Those cumulative offsets are handed to FlashAttention's variable-length entry point as cu seqlens q , and the attention kernel does per-sequence masking internally — one kernel launch spanning differently-sized sequences, boundaries as arrays, not as tensor shape. Which is to say: the modern stack's answer to raggedness is put the sizes in a side-array and make the kernel consult it — the same structural move as int b cols in a 2018 impl.cu , executed by professionals with tiling, TMA, and a few billion dollars of demand behind them. Flat memory, sizes carried as a side-array, a kernel that reads them — int b cols in 2018, cu seqlens q in 2023. An old friend in much better clothes. Even the "sequences just leave the batch when they end" idea — DyNet's per-time-step LSTM batching from 2017 — is recognizable as the ancestor of continuous batching: batch membership defined per step by who still has work, rather than per batch by who arrived together. I want to be precise here, because "I invented vLLM in 2018" is not the claim and would be false in at least three ways. I didn't invent dynamic batching. DyNet's autobatching 2017 and TensorFlow Fold 2017 are prior art, published and cited; InsNet was a contemporary of that line making a different trade at the signature fork — raggedness in the kernels rather than homogeneity in the batches. I'd been working in that line since N3LDG 2019 , which auto-batched dynamic graphs — tree-LSTMs included — faster than PyTorch; the padding-free reframing came later, with InsNet. The idea was in the air; my bet was on which layer should absorb the raggedness. vLLM batches requests, not graphs. Its scheduler moves token-chunks of many requests through one static model , under @torch.inference mode — no autograd, no operator-level graph surgery. InsNet batched arbitrary training graphs , operators-first: it would happily batch two transformers from the same instance, something entirely outside vLLM's problem statement. Same enemy padding , same weapon flat memory + offset bookkeeping , different battlefield serving scheduler vs. training executor . And they solved the problem I never had. Serving needs the KV cache of every in-flight request to survive while the batch churns around it — that's PagedAttention, block tables decoupling logical sequence length from physical memory, and it has no analogue in InsNet because training graphs don't have persistent per-request state. The part with no analogue in InsNet is on the memory side, not the batching side. Right: padding is not a law of nature. It's an artifact of insisting that a batch be a rectangle, and if you're willing to carry sizes as data and write kernels that read them, the rectangle dissolves — along with the mask logic that infects every model built on top of it. The 2021 README sentence holds up; "NLPers will dive into it" just turned out to mean the inference-serving industry rather than the tree-LSTM researchers I was writing for. Wrong: the layer. I bet that padding-free batching mattered for training arbitrary structures , and built a general graph executor with autograd. The world's raggedness problem turned out to be concentrated in one place — transformer decoding, where sequences in a serving batch naturally diverge in length — and the winning implementations attached themselves to that single, ferociously-optimized case. Generality was the wrong axis. Specificity, plus paged memory, plus a scheduler that treats the batch as a per-step decision, is what shipped. Also wrong, in a smaller way: nearly everything about how I wrote it. The erased early history. The scalar GEMM loop. A hyperparameter I'd tune differently today on basically every line. But the shape of the thing — flat buffers, size arrays, kernels that mask threads instead of data — is the shape the field landed on, and there's something quietly vindicating about opening flash attn.py , seeing cu seqlens q , and recognizing an old friend wearing much better clothes. The library is at github.com/chncwang/InsNet https://github.com/chncwang/InsNet , erased early history and all. Code excerpts are from the actual sources of InsNet, DyNet, and vLLM v1 engine , read for this post. If you find a misreading, tell me — my previous source-diving post survived three rounds of adversarial review and I'd like this one to earn the same.