Want to go deeper?

I kept the full technical walkthrough here so the main story could stay readable. You can read these notes in order, search for a topic, or jump straight to the glossary in note 60. The earlier experiments build toward 1.30×; note 63 covers the later official 1.3234× result.

I · The challenge and the model / 01

The run that crossed 1.30×#

Here's the result I was aiming for, then the rewind to the version that was slower. A speedup below one means the baseline won.

Start with the evidence

The number on the screen is 1.3012. It is the official speedup of my dense projection submission over SGLang, not an estimate from my own benchmark. The independent check came back at about 1.302. This article is the story of how I got there. I will follow the mathematics, the hardware, the experiments, and the decisions. I will also follow the mistakes, because the first candidate was slower than the baseline, and several sophisticated ideas stayed slower.

What this victory actually measures

A speedup is baseline time divided by candidate time. At 1.3012, the corresponding time is about seventy 6.85 percent of the baseline. That is approximately twenty 3.15 percent less time. It is not thirty percent less time, and it is not a claim that the entire model serves users thirty percent faster. The competition aggregates several projection workloads. Later I will reconstruct that aggregation and identify exactly where the gains appeared.

Rewind to a disappointing first result

My first recorded candidate passed correctness and checked at 0.7761× the baseline. Correct arithmetic had not made it a competitive kernel. One important diagnosis in the research notes was insufficient parallel work. Some GPU resources were waiting while the available blocks performed long reductions. The lesson was uncomfortable but useful: a more compact implementation can still use the machine less effectively. That finding set up the split router experiment, rather than a celebration of fusion for its own sake.

The path was a branching investigation

The progress line looks smooth only after I choose which results to draw. The actual work branched into vector kernels, split reductions, transposed weights, packed weights, library algorithms, custom CUDA, compression, and new reduction schemes. Some branches failed to execute. Others passed every arithmetic test and lost on time. I will keep those distinctions visible. When the available record supports a decision, I will explain it. Where I reconstruct the rationale from code and measurements, I will say so instead of inventing private deliberations.

The result that almost fooled me

Near the end, my local comparison against SGLang measured 1.3234. That looked as though the later 1.32 target had been reached. It had not. The official scored submission was 1.3012. Measurement conditions, paired ratios, and run variation matter when the remaining gap is small. This is why the article treats the local result as useful evidence for investigation and the official result as the competition outcome.

How to watch the investigation

I will repeatedly ask seven questions. What was slow? What did I think caused it? Why was the proposed change plausible? What changed in the implementation? Did correctness survive? What did the measurements establish? And what decision followed? The first part builds the vocabulary needed to answer those questions. The middle follows my experiments in order. The final part reads the winning design and shows how to reproduce and challenge the result. You do not need to arrive already knowing CUDA. The visual grids are schematic: small groups of cells illustrate a subset of a larger tensor or machine, rather than its literal unit count. Measured results are identified separately. The main story reconstructs my 1.30 milestone. A later official 1.3234 result arrived during production and is documented separately in the postscript.

← Previous Return to worklog Next →

I · The challenge and the model / 03

The complete competition map#

The challenge has several related problems. This is the map I used to keep a projection win separate from an attention or whole-layer win.

A map of the five problems

The competition separates a model layer into several opportunities. Dense projections cover matrix products around attention and the expert block. Attention decode covers the interaction between queries and the cached context. Routed experts cover the selected expert computations. The whole decoder layer joins those pieces with normalization, residual updates, and routing. Prefill runs a whole layer on larger prompt-processing shapes. These are related computations, not five unrelated puzzles. I will use one model diagram and repeatedly highlight the region each problem measures.

Dense projections: my main storyline

My 1.30 result belongs to dense projections. The operation is an input matrix multiplied by a transposed weight matrix. Four weight families are BF16, and the router weights are FP32. The workload contains single-token decode, larger decode batches, and speculative verification batches. Although the equation is the same, the geometry changes the best implementation. Much of the investigation was learning where a direct vector reduction, a custom tensor-core kernel, or a library matrix product deserved to run.

Attention decode: a growing history

Attention decode takes newly produced queries, keys, values, and gate information, then works with a paged key and value cache. Some layers attend over a sliding window, while others attend over the full available context. Preparation, cache writes, the attention calculation, merging partial results, normalization, and gating all contribute work. A long-context attention kernel can spend most of its time reading cached values. That changes the optimization question from merely making a matrix multiplication fast to organizing reuse, parallelism, and reductions across a large history.

Routed experts: irregular work

The expert problem selects eight of 256 experts per token. One GPU holds its local expert shard. The routing distribution is deliberately uneven, so some experts receive much more work than others. The computation includes two expert matrix products, the SwiGLU activation, and output normalization before weighted combination. Here the dimensions of each expert batch vary. Efficient grouping, padding, alignment, and dequantization can matter as much as a fast individual matrix multiplication. A square dense benchmark would hide those difficulties.

The whole decoder layer

The whole-layer problem includes projections, attention, expert routing, expert computation, normalizations, and residual updates. It exposes a common trap: independently faster components do not imply the product of their individual speedups. Their original shares of total time differ, integration can introduce extra materialization, and the actual execution path may differ from a standalone test. My whole-layer result improved, but it was smaller than the strongest component results. That makes it an important part of the story rather than an inconvenient number to omit.

Prefill and the separate board

Prefill handles many new prompt tokens, including large chunks and continued long conversations. Larger matrix dimensions offer much more reuse of weights than single-token decode, and grouped expert products become prominent. The supplied task states that prefill has its own board and does not contribute to the overall decode board. I will explain it so the competition map is complete, while keeping it separate from the projection accomplishment. This distinction also explains why optimizing the largest-looking kernel is not automatically the best route to a particular ranking goal.

← Previous Return to worklog Next →

I · The challenge and the model / 05

Inside the transformer layer#

A model layer is a sequence of calculations. Here I unpack the names so the projection kernels have somewhere to live in the bigger story.

The fused QKV and gate projection

The attention side starts with a projection that produces queries, keys, values, and gate information. These names describe different roles, even when a single matrix product computes them together. Queries ask what information a token needs. Keys provide features used to score relevance. Values provide the information to combine. The gate later modulates the attention output. Under the specified tensor-parallel partition, my local fused projection has a particular output width. The correct kernel must preserve that local arrangement so later operations read the right slices.

Normalization and position

Queries and keys are normalized according to the layer's definition. Sliding-window layers apply rotary positional information, often called RoPE. Then new keys and values are written into the cache in its specified format. These operations sit between the projection and the main attention calculation. They can be small individually yet costly as separate launches. That is why an attention preparation fusion can be useful even without changing the main attention algorithm. But the fusion must preserve scaling, rounding, positions, and cache addressing.

Attention combines the history

Attention compares the new queries with eligible cached keys, applies a softmax to the scores, and uses the resulting weights to combine cached values. The word eligible matters. A causal mask excludes future positions, and a sliding-window layer restricts the history to its window. A full-attention layer keeps the full specified history. An optimization cannot simply skip inconvenient keys or replace the softmax with an approximation outside the rules. The task is to organize all required work more efficiently.

Gating, output projection, and residuals

The attention result is normalized and multiplied by a sigmoid gate, then projected back through the output matrix. Residual updates combine the new contribution with a continuing representation, using the scales and normalizations in the model definition. These surrounding operations explain why an isolated projection result has a narrower meaning than a whole-layer result. They also create potential fusion opportunities, because an intermediate tensor written by one kernel may immediately be read by another. Whether fusion helps depends on dependencies and hardware resource use.

Router, pooler, and local experts

The router produces scores used to select the top eight experts. The expert input is formed through the pooler projection and normalization. Only routes to experts local to this GPU are computed in the isolated expert workload. Each selected expert applies its first learned projection, a SwiGLU transformation, and a second projection. The expert output is normalized and weighted by its route contribution. Uneven routing creates uneven batches, so the execution problem is an irregular collection of matrix products rather than one uniform large GEMM.

Unpooler and the next layer

The expert result is projected through the unpooler and merged into the residual path. A final normalization prepares the representation for subsequent work. At small token counts, the task specifies a special local unpooler slice, which is why a tiny matrix product becomes relevant in my optimization story. I now have the complete loop: attention changes what a token reads from history, and routed experts transform the resulting features. The dense projection challenge isolates five matrix products along that path.

← Previous Return to worklog Next →

I · The challenge and the model / 09

The mathematics of our projection#

A projection is a collection of weighted sums. I'll build up the matrix notation from that idea and explain what M, N, and K count.

One output is a dot product

Take one token's feature vector and one output feature's weight vector. Multiply matching entries and add the products. That sum is one output value. Repeat for every output feature and token row, and I obtain the projection matrix. If the input has M token rows and K input features, and the stored weights have N output rows and K columns, the output has M rows and N columns. This is the meaning of input times weight transpose. The transpose aligns the shared K dimension; it need not imply a physical copy.

The operation count

Each output contains roughly K multiplications and K additions. Across M times N outputs, the conventional floating-point operation count is approximately two M N K. This count helps estimate arithmetic cost, but it does not describe memory transactions, launch overhead, or reductions introduced by a particular kernel. Two implementations can perform almost the same useful arithmetic while moving very different numbers of bytes. Conversely, a method that performs some padded arithmetic can still win because it uses tensor-core hardware more effectively.

Strides tell me where values live

A tensor's shape says how I index it; its strides say how an index maps to memory. A contiguous row-major weight matrix stores adjacent K values together within each output row. A transposed view changes the indexing interpretation without moving values. A contiguous transpose physically rearranges values. A packed layout can go further by arranging tiles in the order a kernel wants to read them. These transformations preserve the mathematical weights when they are lossless, but they can change both memory coalescing and library dispatch.

GEMV and skinny GEMM

At one token, the input is effectively a vector, so the projection resembles a matrix-vector product. At a few tokens, it is a skinny matrix multiplication. At larger token counts, weights can be reused across more rows. A single-token kernel might stream weight rows and reduce them directly in registers. A tensor-core kernel groups many products into tile instructions. A library may choose yet another schedule. The equation does not tell me which implementation wins; the shape and machine constraints influence that choice.

Different reduction orders

In exact real arithmetic, splitting a dot product into partial sums and combining them gives the same answer. Floating-point arithmetic is finite, so changing the summation order can change rounding. Split-K, atomics, and different tensor-core tile schedules can therefore produce slightly different outputs even while retaining the required formats. The challenge's numerical tests define what is acceptable. I preserve the required precision and validate actual outputs, rather than assume algebraic equivalence guarantees identical machine results.

Equivalent products, different execution

The identity relating a product and its transpose lets me exchange how a matrix multiplication is presented to a library. Computing a transposed output and returning a view can sometimes expose a faster algorithm. But output strides, alignment, dimensions, and downstream expectations matter. I explored these orientations explicitly and later added alignment-based fallbacks. This is a recurring theme: a mathematically simple identity can open a useful implementation branch, while the practical details determine whether it is a robust optimization.

← Previous Return to worklog Next →

I · The challenge and the model / 10

The exact projection workload#

These are the actual sizes the projection benchmark tested. Think of a shape as the dimensions of one particular job.

QKV and the output projection

The local QKV-and-gate weight has 2,304 output rows and 8,192 input features. In BF16, its weight payload is about thirty 7.7 megabytes in decimal units. The output projection has 8,192 output rows and 1,024 input features, about sixteen point eight megabytes. These sizes explain why even a single token can require substantial weight streaming. The small activation does not make the weight matrix disappear.

Pooler and unpooler

The pooler has 4,096 output rows and 8,192 input features. The full unpooler has 8,192 output rows and 4,096 input features. Each contains about sixty 7.1 megabytes of BF16 weights. For token counts up to eight, the specified unpooler projection only returns its 1,024 output slice, substantially reducing the weight rows it needs. A kernel should read the requested slice rather than automatically stream the whole matrix.

The FP32 router

The router has 256 output features and 8,192 input features. Its stored weights are FP32, about 8.4 megabytes. The output is also FP32, unlike the dense projection outputs, which are BF16. Although the router matrix is smaller than the pooler, its stricter numerical contract and low output-feature count make its kernel design interesting. At small token counts, enough parallelism can be difficult to expose without splitting the reduction.

Thirty scored cases

Each of the six loads times all five projection families. Within a load, each shape has the same listed weight, 62, corresponding to repeated layer work in the model. That common factor cancels in the load's speedup ratio, but the individual projection times do not. A one-microsecond saving on a twenty-microsecond pooler and a one-microsecond saving on a five-microsecond router contribute the same absolute reduction to that load's time. Comparing percentage improvements alone can therefore mislead my priorities.

Nominal lower bounds are guides

The task publishes floors based on moving data at 4.8 terabytes per second or performing arithmetic at 989 tera floating-point operations per second, whichever takes longer. These are nominal hardware bounds. Actual kernels also pay for instruction scheduling, imperfect traffic patterns, synchronization, and launch-related costs. My practical read probes were slower than the nominal floor. I used both figures: the nominal bound explained what could not be beaten, while the practical measurement suggested what this machine and access pattern could realistically achieve.

The shape atlas becomes a decision map

Instead of remembering that one method was generally good, I built a table across layer and token range. That table becomes the dispatch plan. It captures where a vector reduction wins, where cluster split-K helps, where packed pooler weights help, and where a library should remain in control. The appendix exposes every official projection case and its measured result. This makes the final 1.30 score inspectable: you can see which cases improved, which stayed near baseline, and which still regressed.

← Previous Return to worklog Next →

I · The challenge and the model / 11

Numerical formats and error#

A computer has to choose how many bits to spend on a number. That affects storage, speed, and how close an answer can be to the reference.

Bits divide range and precision

A floating-point representation divides its bits among sign, exponent, and significand information. The exponent largely determines range; the significand determines how finely values are represented within that range. FP32 uses one sign bit, eight exponent bits, and 23 stored fraction bits. BF16 uses one sign bit, eight exponent bits, and seven stored fraction bits. Sharing the exponent width does not make their precision equal. BF16 can cover a broad range while rounding many distinct FP32 values to the same stored value.

Storage precision and accumulation precision

A matrix product can read BF16 operands and accumulate partial sums in FP32. That preserves the original operand values while making the running sum more precise than a BF16 accumulator would be. The final dense result is converted to the required BF16 output. These are separate choices: operand storage, product instruction behavior, accumulation, split reduction, and output format. A library configuration that changes a reduction to a lower-precision type can violate the intended numerical design even if its output tensor has the correct dtype.

Why summation order matters

Consider adding a small value to a much larger finite-precision value. The small contribution may fall below the larger value's representable spacing and disappear after rounding. If several small values are summed first, their combined contribution may survive. A split reduction changes grouping, and atomic contributions may arrive in different orders. This is why two legitimate implementations can differ slightly. The task does not ask for arbitrary bitwise identity, but it does impose concrete error tests, including row-level checks that can expose concentrated failures.

How the projection error is checked

The task compares outputs with a high-precision reference using relative L2 error. Dense outputs have a one-percent bound, while the router has a much tighter 0.01 percent bound. There are also per-row safeguards and special sparse probe rows. Roughly one row in eight has one nonzero power-of-two input, exposing a weight column directly. That check helps catch a method that quietly stores approximate weights. It is possible for a global norm to look acceptable while a particular row is unacceptable.

The other formats belong to other contracts

The attention cache uses the specified FP eight format, while the expert weights use NVFP four. Those facts do not authorize changing the dense projection weights to FP eight or four-bit weights. Each problem has its own representation and numerical rules. A tempting optimization can be scientifically interesting and still be outside this challenge's contract. My compression experiments therefore attempted lossless representations of the original BF16 values, rather than substituting a lower-precision approximation and hoping the aggregate error would hide it.

The high/low router exception

For the router, the rules explicitly permit a BF16 high part plus a BF16 low part, accumulated in FP32. First round a weight to the high part. Then compute the residual in FP32 and round that residual to the low part. The pair approximates the original FP32 weight much more accurately than the high part alone. It is not an arbitrary permission to reduce precision elsewhere, and it should not be described as a mathematically exact decomposition for every FP32 value.

← Previous Return to worklog Next →

I · The challenge and the model / 12

Why the router needs special care#

The router helps select which experts process a token. Its numerical requirements are stricter, so I couldn't treat it like every other projection.

A small matrix controls a discrete decision

The router's outputs help determine which experts process a token. That selection is discrete: an expert is included or excluded. If two scores are close, a small numerical change can alter their ordering. The projection task measures the router's numeric output directly with a tighter bound, and the full-layer task also cares about routing behavior. A relatively small matrix therefore carries important correctness consequences. It is not safe to assume its values can be approximated as loosely as a downstream dense activation.

Constructing the two weight parts

Let a router weight be w. Define the high part by converting w to BF16. Convert that high part back to FP32, subtract it from w, and convert the residual to BF16 to obtain the low part. The low part represents detail that was too small to fit in the first rounding. Preparing these static parts once avoids repeating the conversion for every activation. The evaluated matrix products still use fresh input activations and FP32 accumulation.

Two products and a sum

The intended computation can be organized as an input product with the high weights plus an input product with the low weights. A straightforward implementation might use two GEMMs and then an addition. That is easy to understand but can pay multiple launch and intermediate-memory costs. Other implementations place the two weight parts into one larger or interleaved operand so a tensor-core kernel produces both contributions in one path. The final sum must combine corresponding high and low outputs, not unrelated adjacent features.

Packing high and low values

I explored layouts that place high and low contributions in paired output columns. A kernel can load the packed weight region, accumulate the two contributions in FP32, and combine the pair in its epilogue. This changes memory access and the shape seen by the dot-product instruction. It also changes the amount of intermediate state the kernel holds. Packing is useful only if the saved launches or better transactions outweigh register pressure, repeated activation reads, and any reduction overhead introduced by the new design.

Why the router invited split-K

There are only 256 router outputs. If a block handles many outputs and several token rows, the output grid can contain too few blocks to occupy the GPU well. Splitting the long input-feature reduction creates additional blocks. Each block computes a partial contribution and must participate in a correct final sum. This is the central tradeoff behind my early router improvement: more parallelism can reduce the long per-block work, but the resulting communication and synchronization must remain cheap enough.

The router as a recurring test case

The router appears repeatedly in my story because it exposes several systems lessons at once. Insufficient parallelism can dominate. Extra launches can dominate. Packed layouts can help or hurt. Atomic reductions can change rounding order. Fresh graph replay can reveal counters that were not reset. Finally, a substantial router-only speedup may translate to a modest overall projection score because the larger weight matrices still consume much of the load's time. I will examine both the local router gain and its aggregate contribution.

← Previous Return to worklog Next →

II · Hardware, kernels, and measurement / 13

A guided tour of the H200#

The H200 was my experiment machine. Here's the useful hardware vocabulary, with each part connected to the work it does.

A GPU is many execution resources

A GPU is not one enormously fast CPU core. It contains many streaming multiprocessors, or SM units, that execute groups of threads. Each SM has instruction execution resources, registers, shared memory resources, and access to specialized tensor-core arithmetic. The device also has a shared last-level cache and high-bandwidth external memory. A kernel's design determines how much of that machine can work at once and how efficiently data reaches it. Owning a fast GPU is therefore only the beginning of obtaining a fast result.

HBM supplies the large tensors

High bandwidth memory holds the model weights and large activation or cache arrays. Its capacity allows these arrays to reside on the GPU, but capacity and bandwidth are different properties. Even when the weights fit comfortably, reading them still takes time. At single-token projection shapes, a large fraction of useful work involves streaming weights with little reuse across token rows. This is why the payload sizes in the shape atlas matter. A memory-bound kernel can leave arithmetic hardware waiting even though the GPU has plenty of free memory.

L2 changes what reaches HBM

The L2 cache sits between execution resources and external memory. If data is reused while resident, later accesses can avoid another external fetch. That is valuable in production, but it complicates isolated benchmarking. A candidate repeatedly timed on the same small weight set may appear fast because the benchmark warms the cache. The judge deliberately begins each replay with cold L2 conditions and rotates weight copies. I copied that principle into the lab so a method would not win solely because my test accidentally favored cached weights.

Shared memory and registers

Registers are fast storage private to executing threads. Shared memory is explicitly managed storage accessible to threads in a block. Kernels use these resources to hold tiles and reuse data rather than repeatedly fetch from external memory. But both resources are finite. A larger accumulator consumes registers. A deeper load pipeline consumes shared memory. Those choices can reduce the number of resident blocks or warps, making it harder to hide latency. Resource use is therefore part of the algorithm's performance cost, not just a compiler detail.

Specifications and measured floors

The task's nominal bandwidth and arithmetic rates give useful lower bounds. They do not imply that every access pattern reaches those rates. My practical probes measured how quickly this machine streamed selected payload sizes under the lab's conditions. For example, the handoff recorded about ten point 43 microseconds for a thirty 7.7 megabyte read, slower than the task's nominal bandwidth-only floor. I used such measurements to ask whether a kernel was leaving obvious room, while remembering that a matrix kernel must also compute and synchronize.

Match the evaluation machine

The remote lab used an H200 S X M and the specified SGLang image so that architecture, libraries, and toolchain resembled the judge. I also read clocks and power settings rather than assume any machine named H200 behaved identically. This reduces avoidable differences but cannot eliminate all run variation. The judge and pod had different driver versions in the recorded machine metadata. A useful local result is evidence for a promising change, not a replacement for evaluating it on the official system.

← Previous Return to worklog Next →

II · Hardware, kernels, and measurement / 15

The journey of a byte#

Before the GPU can multiply a value, it has to get that value. I'll follow the memory path and explain why moving data can dominate the calculation.

Movement can dominate arithmetic

Consider a single-token projection with a large BF16 weight matrix. Each weight value participates in a small amount of useful arithmetic before the kernel moves on. The tensor cores may be capable of tremendous arithmetic throughput, but they cannot use values that have not arrived. If the required byte movement takes longer than the minimum arithmetic time, memory is the limiting resource. Optimizing a few arithmetic instructions in that regime may barely change the measured time. Improving access patterns or avoiding extra traffic can matter much more.

Reuse changes the cost per output

When multiple token rows use the same weight tile, a kernel can keep that tile in a faster level of memory and apply it to several rows. This amortizes weight movement over more useful operations. But reuse requires storage and a schedule that makes the reused values available at the right time. Large tiles can capture more reuse while increasing shared-memory and register demands. A design therefore trades movement against resource pressure. The ideal tile for one token need not resemble the ideal tile for 256 tokens.

Scratch traffic is real traffic

Split algorithms often write partial outputs into scratch memory and read them back during a reduction. That traffic is additional to the original input, weights, and final output. Scratch also introduces a dependency: the final answer cannot be complete until the needed partials exist. A second reduction kernel adds another launch. A fused last-arriving-block design replaces that launch with counters and synchronization. Neither arrangement is free. The correct comparison counts all of these costs, not just the shortened matrix-product loop.

Cache policies are implementation choices

A load can be configured to interact with cache levels differently. In my single-token experiments, streaming weight loads with the selected .cg policy helped specific kernels. The input vector had a different reuse pattern from the large weight stream, so treating both identically was not obviously optimal. Cache policy syntax is architecture and compiler dependent, and invalid combinations can fail during assembly. I encountered that distinction during the search. A failed policy configuration is not a measured performance result.

Kernel launches create boundaries

A kernel boundary can mean another launch, another synchronization dependency, and another set of intermediate memory accesses. Fusion can remove some of those costs. But fusion also combines resource needs and can reduce parallelism. My first slow router and later failed fused reductions demonstrate both sides of that tradeoff. It is better to ask which boundary is expensive and what replacing it costs than to assume fewer launches always wins. The complete captured graph is the object I ultimately time.

Account for the whole byte budget

For any proposed kernel, start by listing weight bytes, activation bytes, output bytes, and scratch bytes. Then ask which reads are reused through cache or shared memory and which values are redundantly fetched by multiple blocks. This accounting is approximate, but it often explains why a seemingly clever method loses. A compressed representation may read fewer weight bytes yet add metadata traffic and decoding instructions. A split kernel may expose more blocks yet reread activations. The byte budget turns intuition into a testable prediction.

← Previous Return to worklog Next →

II · Hardware, kernels, and measurement / 16

Coalescing, tiling, and weight layout#

The same numbers can be arranged in memory in different ways. Some arrangements make neighboring threads' reads much easier to serve together.

Adjacent lanes should help each other

A warp's memory accesses can be combined efficiently when nearby lanes request nearby addresses in a suitable alignment pattern. This is the intuition behind coalescing. If lanes access widely separated locations, the hardware may need more transactions to deliver the same useful values. The precise transaction behavior depends on the access size, alignment, and architecture, but the design principle is simple: arrange work so the machine transfers useful contiguous regions rather than many poorly utilized fragments.

Why a tile is more than a rectangle

A tile describes the subset of matrix values a block loads and computes. Its shape determines which dimension is contiguous, how data is reused, which tensor-core instructions fit, and how much shared memory is required. The tile can also change the number of independent blocks. A wider output tile may improve reuse but reduce the grid size. A narrower output tile may expose more blocks but repeat activation loads. Tuning is a search through these interacting effects, not a search for the largest rectangle that fits.

A physical transpose

A contiguous transpose rearranges weights so a different matrix dimension becomes contiguous in memory. That can change the memory path seen by a custom kernel or the layout seen by a library. Because the weights are static, I can perform the rearrangement during preparation and reuse it. The tradeoff includes preparation time and additional resident memory. A transposed view without a copy is different: it changes stride metadata while keeping the same physical values in the same places. I distinguish those operations carefully in the code walkthrough.

The packed pooler layout

My packed pooler groups output features into tiles and stores each tile as a contiguous sequence across K and the tile's output columns. A block then streams the weight region it owns in the layout expected by its dot product. This is a lossless reordering: every original BF16 value is preserved. The selected tile width was 32 output features for the included pooler paths. The benefit came from the complete access and execution pattern, not from changing the mathematical projection or quantizing its weights.

Shared-memory bank conflicts and swizzling

Shared memory is organized into banks. Certain access patterns make multiple lanes contend for bank resources, which can serialize work. The custom cluster kernel rearranged shared-memory positions with a swizzle so its matrix-load instruction pattern encountered a more favorable bank mapping. The swizzle changes addresses within the staging tile, not the mathematical values. Such details are where a high-level matrix equation becomes a hardware-specific implementation. They are also a reason to validate both address mapping and final outputs after changing a layout.

Alignment and nearby shapes

An output layout that is well aligned for a token count divisible by eight may behave badly for an odd token count. My alternate library orientations encountered this risk, so the final design retained a fallback when the swapped output's leading dimension did not meet the chosen alignment condition. Nearby-shape measurements were valuable here. Without them, a method could look excellent on the scored values while collapsing on a legitimate neighboring shape. Robust specialization uses meaningful shape properties and preserves a safe general route.

← Previous Return to worklog Next →

II · Hardware, kernels, and measurement / 19

Split-K and the cost of combining work#

Split-K means dividing a long sum into partial sums. It creates more parallel work, then charges me for putting the answers back together.

Start from one long reduction

An output tile is a set of token rows and output features. Each output value sums products across the input-feature dimension K. In an unsplit design, one block can own that whole reduction for its tile. This avoids communicating partial results to other blocks, but it leaves the grid size determined by output geometry. When that grid is too small, the block can have plenty of work internally while the overall device has too little independent work to schedule.

Turn K into another grid dimension

Split-K assigns different portions of K to separate blocks for the same output tile. With four splits, each block performs approximately a quarter of that reduction. There are now four times as many contributing blocks, potentially improving utilization and hiding latency. The useful multiplication count is broadly similar, although padding and implementation details can change it. The final output is incomplete until all four contributions are combined. Split-K exchanges a long local reduction for a parallel reduction followed by a combination problem.

A separate reduction kernel

The straightforward design writes FP32 partials to a scratch tensor. A second kernel reads the partials, sums across splits, and writes the required output. This makes dependencies easy to understand: the second launch follows completion of the first on the stream. But it adds scratch traffic, another launch, and final reduction work. For a large matrix, the new parallelism may easily repay that cost. For a tiny router or a kernel already close to a latency floor, the extra cost can consume the entire benefit.

Atomic addition into an output

Another approach initializes the output to zero, then lets each split atomically add its FP32 partial contribution. An atomic operation protects the update from conflicting simultaneous writers. This can avoid a separate partial-sum reduction kernel, but the output-zeroing operation remains part of the captured call path. Contention and arrival order also matter. FP32 atomic addition preserves the intended accumulator type while potentially changing reduction order, so correctness and run behavior must still be checked on fresh inputs.

The last-arriving-block design

A fused design can write each partial, increment a per-tile counter with suitable memory ordering, and let the last arriving block perform the final sum. This removes one launch but introduces global synchronization logic, partial reads, and a correctness-sensitive reset. A counter that is not returned to the initial state can fail on the next graph replay. Memory ordering must ensure the last block sees completed partial writes. This is a good example of complexity that must earn its place through measured savings, not merely fewer launches.

How I choose a split count

Too few splits can leave the GPU underused. Too many can increase scratch, atomics, repeated activation reads, and synchronization. The best count depends on M, N, K, tile dimensions, and hardware residency. I therefore sweep candidate counts and retest finalists with paired timing and nearby shapes. The early split-router result showed the value of added parallelism. Later failed fused reductions showed that combining the work is often the harder half of the optimization. Neither lesson makes split-K universally good or bad.

← Previous Return to worklog Next →

II · Hardware, kernels, and measurement / 20

Hopper pipelines and block clusters#

The next hardware ideas are about overlap and cooperation: fetch the next tile while computing this one, and let selected blocks combine their work.

Overlap loading and computation

A simple tiled kernel loads data, waits, computes, and repeats. A pipelined kernel tries to load a future tile while computing a current tile. Multiple staging buffers keep these tiles distinct. The goal is to hide movement latency behind useful arithmetic. But more pipeline stages consume more shared memory and may reduce resident blocks. They can also be wasted if each block has only a few K tiles. A deeper pipeline is therefore a tunable tradeoff, not a free source of throughput.

What cp.async supplies

The custom QKV kernel uses asynchronous global-to-shared copies. Groups of copies are committed, and the kernel waits for an appropriate number of groups before using a tile. Shared-memory addresses are swizzled for the later matrix loads. Threads then synchronize where needed before consuming staged values. These operations let the design organize movement and computation explicitly. They also create obligations: a consumer must not read unfinished data, and a staging slot must not be overwritten while another operation still needs its previous contents.

ldmatrix and mma.sync

After staging, the kernel loads matrix fragments from shared memory into registers using ldmatrix instructions. Warp-level MMA updates FP32 accumulators. The loops walk through the K tile, then advance the pipeline. Different warps own output-feature rows, and the token tile determines how many input rows the block reuses. The implementation is specialized to the matrix instruction's fragment layout. My diagram follows one subset of weights and activations so the assembly-level names stay connected to actual data and arithmetic.

Distributed shared-memory reduction

Hopper thread-block clusters let cooperating blocks access distributed shared memory. My QKV design splits K among blocks in a cluster, places each block's FP32 partial result in shared memory, synchronizes the cluster, and reads peer partials to form the final answer. A second cluster synchronization protects the lifetime of that shared state. This avoids a separate global scratch-reduction launch, but cluster barriers and peer accesses still have costs. The method was useful in a particular small-token QKV range, not automatically across the whole workload.

TMA, multicast, and WGMMA

Hopper also offers tensor memory transfers and warp-group matrix operations. A well-designed kernel can use these mechanisms to arrange bulk tile movement, reduce repeated transfers, and feed larger cooperative arithmetic structures. However, these features require compatible layouts, synchronization, and scheduling. My separate TMA experiment did not demonstrate a win, and my selected QKV kernel did not become a TMA-plus-WGMMA implementation. I explain the capabilities because they inform future hypotheses, while keeping the actual implemented and measured designs clearly identified.

Why deeper cluster pipelines lost

The later cluster sweep tried different stage counts, warp counts, K widths, and token tile variants. The deeper designs did not provide a useful improvement over the selected path. That outcome is consistent with several possible limits, including shared-memory pressure and the amount of per-block work available to overlap, but the exact cause needs more profiling than the available counters supplied. I keep the measured rejection separate from a proven microarchitectural diagnosis. Timing tells me the design lost; it does not by itself identify every reason.

← Previous Return to worklog Next →

II · Hardware, kernels, and measurement / 23

How the judge times a kernel#

A timing depends on how it was taken. This is why I used replay, fresh inputs, cold-cache treatment, and paired baseline/candidate measurements.

Eager time is not replay time

A Python call can include dispatch overhead, allocation behavior, compilation, and synchronization effects. The challenge instead captures the required call as a CUDA graph and measures graph replay, resembling the production decode path. A candidate that looks good in eager execution may lose under captured replay, and the reverse can happen. I therefore compare the execution mode the judge actually uses. Timing a different mode is useful for debugging, but it cannot establish the competition result.

Start with cold L2

Before each timed replay, the judge reads a buffer larger than the cache to remove an accidental warm-cache advantage. The local timer uses the same basic principle. It also rotates independent weight states so repeatedly using one convenient resident copy does not become the basis of the result. This does not reproduce every detail of the official evaluator, but it aligns the lab with a central property of the benchmark. I am trying to improve the required data path, not exploit a cache state the real test deliberately excludes.

Alternate the execution order

Suppose the first of two measurements tends to experience a different condition from the second. Always running the baseline first could bias the ratio. The judge pairs repetitions with the order swapped: baseline then candidate, and candidate then baseline. It computes paired ratios and uses a median summary. This helps cancel simple order effects and reduces sensitivity to isolated outliers. Pairing does not eliminate every possible difference, but it is much stronger evidence than comparing two unrelated best-case timing numbers.

Subtract the fixed replay cost

For a short kernel, the graph replay's fixed cost can be a substantial fraction of the raw event interval. The evaluator estimates that fixed overhead and subtracts it from both sides before calculating the kernel comparison. My local runs used an explicitly recorded overhead value, around four microseconds. Subtracting a nearly fixed number from very small timings increases sensitivity to measurement differences. That is one reason I avoid interpreting a tiny local percentage change as a guaranteed official gain.

Profiling and scoring are different tools

A profiler tells me which kernels ran and helps identify where time went. The official paired benchmark establishes the score. Profiling can introduce overhead or show durations affected by waits and overlap, so adding profiler rows is not always the same as measuring a complete call's wall time. My investigation used profiles to form hypotheses and paired timing to choose implementations. This separation prevents a beautiful profile breakdown from being mistaken for an independently reproduced speedup.

Reproduce the method, then use the judge

The remote lab matched important conditions: graph replay, fresh activations, cold cache, rotated weights, and swapped-order pairs. It still differed from the judge in machine details and parts of the surrounding harness. The local 1.3234 measurement was therefore a reason to submit, not proof of an official 1.32 score. This distinction is central to the story. A good local benchmark narrows the search; the official evaluation tells me what the competition recorded.

← Previous Return to worklog Next →

III · The experiment story / 27

A correct kernel that lost#

The first candidate got the answers right and still lost on time. That was the starting point for looking at how much independent work it exposed.

Candidate 001: a plausible starting point

The first recorded projection candidate used direct vector reductions for single-token work, a custom high-and-low router kernel for other token counts, and the baseline dispatch for the remaining dense products. The router held FP32 accumulators and looped through K tiles. On paper, it was straightforward: prepare static router parts, perform the required products, and combine them. It passed the reference checks. That established a usable starting implementation, but it said nothing yet about whether the GPU was being used efficiently.

The first check: 0.7761×

The official check reported 0.7761. That is a slowdown: a ratio below one means the candidate takes longer than the baseline on the aggregated workload. The result was not a failed arithmetic test. It was a correct implementation whose performance was worse. This distinction matters because the next action should address execution structure rather than relax numerical requirements or declare the reference too strict. I needed to understand what the kernel asked the GPU to do.

Too little independent router work

The unsplit router output grid was determined by token tiles and output-feature tiles. For a small batch and only 256 router outputs, that can expose few blocks. Each block still traverses the long K dimension, so a large amount of work remains serial within the available blocks. The research notes identify insufficient parallelism as an important diagnosis. This is a concrete mechanism: idle SM capacity cannot help a block unless the computation is partitioned into schedulable independent work.

Fusion had not solved the real limit

The candidate combined useful work, but reducing source-level steps is not the same as reducing elapsed time. If the resulting kernel exposes too little parallelism or uses resources inefficiently, a library with more launches can still win. This is the first major change in the investigation's direction. Rather than ask only how to fuse more, I asked how to create enough useful work across the device. The mathematics allowed splitting the reduction, provided the partial results were combined correctly.

A falsifiable next hypothesis

The proposed change was to distribute the long router reduction across multiple blocks. Each block would compute a subset of K and write a FP32 partial. A final kernel would sum those partials. This adds work, traffic, and a launch, so it could have failed. The prediction was that better utilization would save more time than the extra reduction cost. That is a useful hypothesis because a complete correctness check and paired performance run can test it directly.

What this failure contributed

Candidate 001 did not become the final submission, but it contributed a valid reference-compatible custom path and a clear negative result. The source and check remain in the ledger. Keeping them lets me explain the next improvement as a response to evidence rather than pretend the winning architecture was obvious from the beginning. The lesson is not that custom kernels are bad. It is that correct arithmetic and a plausible implementation are starting points, while measured execution determines the next engineering decision.

← Previous Return to worklog Next →

III · The experiment story / 28

Split the router: our first useful win#

The first useful change split the router's long reduction into more blocks. Here's the extra work it added and why it could still finish sooner.

Candidate 002 changes the work grid

Candidate 002 added a split dimension to the router launch. For the small token ranges, it selected sixteen splits; larger ranges used a different count. Each block traversed its assigned interleaved K tiles and accumulated high and low products in FP32. The output grid now contained many more independent blocks than the unsplit version. This directly addressed the earlier diagnosis while preserving the mathematical operation and the router's stricter output requirements.

The partial tensor and final sum

The partial tensor had a split dimension followed by token and output dimensions. Each contributing block stored its fresh result into the slice it owned. A reduction kernel then walked the split dimension and wrote the final router output. The implementation had a simple stream dependency between launches, making completion easy to reason about. The performance question was whether the extra global reads and writes were cheaper than the underutilized long reduction they replaced.

The official evidence changed direction

The new check returned about 1.1014, and the independently scored submission returned 1.1037. Both passed. The aggregate result had moved from a slowdown to a useful improvement over SGLang. It was important to obtain that independent scored run rather than report the check as though it were already on the board. The improvement supported the parallelism hypothesis for the selected workload, even though it did not establish the same benefit for every individual shape.

The profile revealed uneven gains

The recorded analysis of the split-router result showed strong router improvements at several small token counts, while the 256 token router case regressed. This uneven behavior is not surprising once I consider the output grid. Larger M already creates more independent tiles, so the benefit of splitting can shrink while reduction overhead remains. The same implementation choice can therefore be useful at one end of the workload and harmful at the other. That observation became the next dispatch question.

Why I did not stop at the aggregate

An aggregate win can hide a serious regression. If I only kept the final number, I would know the candidate was better overall but miss an opportunity to improve it further. The saved per-shape records let me identify where custom work helped and where baseline dispatch deserved to remain. This moved the strategy from replacing a whole operation family to selecting complementary paths. It also made the investigation more auditable: the next change had a specific target rather than an unspecified hope for another few percent.

The first decision pattern

The sequence now had a recognizable structure. I observed a slowdown, identified a plausible utilization problem, introduced split parallelism, validated the arithmetic, and measured a better result. Then the measurements exposed shape dependence. This pattern repeats throughout the project. A successful experiment rarely answers every question. It often reveals the next constraint more clearly, which is why a useful result should include both its achievement and the conditions under which it stops helping.

← Previous Return to worklog Next →

III · The experiment story / 30

Packed router: reaching 1.1555×#

The router's high and low weight parts could be packed into one library operand. This was an arrangement change with different methods retained for different sizes.

Candidate 009 prepares a paired operand

Candidate 009 formed a packed router matrix by concatenating the transposed high and low BF16 weight parts along output columns. The original FP32 weights were still retained, and the decomposition followed the permitted high-plus-low scheme. Preparation performed this static rearrangement once. The goal was to let a library matrix multiplication compute both product families in a single call for the larger token range, rather than force that range through the same small-batch custom split kernel.

One product produces two output regions

The packed product has twice as many output columns as the logical router. The first region contains the high-weight product, and the second contains the low-weight product. Adding corresponding columns returns the FP32 router result. This is a different packing from the later interleaved-column kernel, where each high and low pair sits adjacent. Both encode the same two contributions, but their memory layouts and epilogues differ. Keeping those distinctions clear is necessary to understand which experiment actually changed what.

Keep the small-batch custom path

For smaller token counts, candidate 009 retained the split-router implementation and final reduction. For the larger range, it used the packed library product. This was not a claim that the library route had replaced custom kernels everywhere. It was an explicit combination of two regimes. Selected single-token dense layers also kept direct vector reductions. By this point, the submission was becoming a policy over several methods, each chosen for a different geometry and measured tradeoff.

The check and scored run

The packed-router candidate checked at about 1.1372. Its independently scored result was 1.1555. The difference is another reminder that checks and scored executions are separate noisy evaluations. The scored result became the early projection incumbent. I do not rewrite the check to match the better number, or treat the difference as a new code improvement. Both measurements belong to the same experiment history and should remain visible.

Why this still left a large gap

The router improved substantially, but many dense projections still used baseline dispatch or modest one-token custom paths. The larger weight matrices occupied much of the load time. Further progress therefore required work beyond the router. The per-shape breakdown suggested transposed weights, optional Hopper GEMV routes, wider K tiles, and alternative library kernels. I would test those ideas rather than assume the packed-router gain could simply be repeated on unrelated matrix families.

A useful principle emerges

The packed-router result established a practical pattern: transform static weights to present an equivalent computation to a better execution path, then dispatch by the shape regime where that path helps. This principle later reappeared in the transposed dense experiments, packed pooler tiles, and cuBLAS Lt orientation searches. The implementations were different, but the reasoning connected them. I was using freedom in representation and scheduling while preserving the operation the challenge actually required.

← Previous Return to worklog Next →

III · The experiment story / 35

Systematic vector and library searches#

Instead of guessing one setting at a time, I searched related configurations. The hard part was keeping the search valid and retesting promising results.

A search needs an explicit hypothesis space

Instead of changing one number repeatedly without a record, the next candidates explored defined families. Vector kernels varied output rows per block and warp counts. Router designs varied token tiles, output tiles, K tiles, splits, and reduction strategies. cuBLAS Lt searches varied operand layouts and validated algorithm choices. Each family represented a mechanism I could explain. Bounding the space made the search reproducible and prevented compilation or preparation from consuming the whole run.

Candidate 022: reusable library plans

The cuBLAS Lt plan search explored how an equivalent product could be presented to the library and which suggested algorithms were valid. Plans were reused rather than performing expensive heuristic selection on every timed invocation. FP32 compute and acceptable reduction schemes were retained. The complete checked result was about 1.1590. That alone did not establish the target, but per-shape results could identify useful regions for a later combined policy.

Candidate 023: direct vector reductions

The vector search explored direct reductions for the smallest-token regime. A block's row grouping and warp count influence memory streams, register distribution, and the final reduction. The checked aggregate reached about 1.1815, stronger than several earlier dense branches. I still needed to examine which layer shapes produced that benefit and avoid assuming the same settings would help larger token counts. A direct vector method is a regime-specific tool, not a substitute for all matrix multiplication.

Candidates 024, 025, and 027: router alternatives

Several router searches examined tile choices, paired high/low layouts, and fused reduction ideas. Their complete checked results stayed around 1.14 to 1.15. A design that reduced a launch count did not automatically beat the best composed route. Counter logic, global partial traffic, and repeated activation loads could outweigh the proposed benefit. The exact microarchitectural cause was not always established, but the complete timings were sufficient to avoid promoting those candidates as aggregate winners.

Candidate 028 broadens algorithm choices

The expanded library search considered validated split, tile, and stage variants beyond a narrow default choice. It checked at about 1.1783. This broadened the evidence without proving that exhaustive search would be efficient or that one algorithm index would generalize to every neighboring shape. The later C2 investigation would revisit library orientations more explicitly. At this point, the important outcome was a collection of candidate paths with measured strengths that could be composed carefully.

Screen, retest, then compose

A broad search should spend fewer repetitions screening clearly poor configurations and more repetitions on finalists. The finalists also need fresh-input correctness, outliers, and neighboring shapes. Only after that should complementary methods enter a combined source. This process cannot remove all selection noise, but it is stronger than choosing the fastest single replay from thousands of trials. It also gives a clear explanation of a decision: a method was kept because repeated measurements supported a useful region, not because its best isolated number looked exciting.

← Previous Return to worklog Next →

III · The experiment story / 38

Atomic router: reaching 1.2785×#

Could each router split add directly to the output? Atomic addition helped, as long as I also included the fresh zero-fill and checked the arithmetic.

The dedicated reduction was a target

A split-router path can spend useful time computing partials and then additional time writing, reading, and reducing them. Candidate 030 explored accumulating split contributions directly into the final FP32 output with atomics. This targeted the combination path rather than changing the high-and-low representation or relaxing its precision. It was a plausible next step because the router's small overall computation made launch and reduction overhead a meaningful part of the total.

Initialize a fresh output

Before atomic accumulation, the output must contain zeros. Reusing an output with old sums would make the next invocation incorrect. The selected function therefore allocates a zeroed FP32 output and launches the split matrix kernel against it. This is an important qualification to the phrase single-launch router: there is no dedicated final reduction launch, but output initialization still belongs to the complete captured path. The benchmark must include it, as my complete-call timing does.

Interleave high and low contributions

The router weights are arranged so corresponding high and low columns form pairs. The matrix kernel accumulates those products in FP32, reshapes the paired accumulator, combines each pair, and atomically contributes the split result. Token and output masks preserve the logical shape. This layout differs from the earlier concatenated packed library route. The difference matters because it changes which values a block loads together and how the epilogue forms each router output.

What atomics guarantee and what they do not

Atomic addition coordinates conflicting updates so one contribution does not simply overwrite another. It does not promise the same summation order every time. FP32 rounding can therefore differ slightly across arrivals. The task's strict router checks and fresh-input replay tests remain necessary. The implementation must also ensure each split contributes once to each valid output location. Correctness comes from valid indexing, the intended arithmetic, initialization, and synchronization semantics together, not from the word atomic alone.

The independently scored milestone

Candidate 030 independently scored 1.2785. This was a substantial step beyond the earlier 1.1993 bundle. It established a stronger projection incumbent while retaining selected dense paths. I keep the exact source and result linked rather than infer a score from the router-only improvement. The aggregate milestone reflects the full dispatch and all six loads, including cases where the baseline library remained the selected route.

What remained after this win

The router improvement did not remove the dense weight-traffic floor. QKV, pooler, output, and unpooler still occupied important shares of the loads. The next useful changes targeted a small unpooler slice and custom cluster reduction for small-token QKV. This is the repeated pattern of the investigation: after a successful change, inspect the remaining shape profile rather than apply the same technique everywhere. The best next hypothesis often belongs to a different part of the computation.

← Previous Return to worklog Next →

III · The experiment story / 40

Custom cluster QKV: reaching 1.2888×#

This QKV kernel combined partial results through cooperating blocks' shared memory. The barriers mattered as much as the matrix instructions.

Why QKV remained a target

The small-token QKV projection streams about thirty 7.7 megabytes of weights. Its execution remained above the practical read floor, and its skinny token dimension made ordinary large-GEMM strategies awkward. The custom cluster design attacked the long K reduction and data movement together. It used tensor-core arithmetic within blocks and split contributions across a cooperating block cluster. The goal was not to eliminate required work but to schedule it so more of the machine contributed effectively.

Stage swizzled weight and activation tiles

The CUDA kernel asynchronously copies tiles into shared memory. The shared tile stores weight rows and token rows, with a swizzle designed for its matrix-load pattern. Several stages allow a future tile to be in flight while the current tile computes. Threads then load register fragments and issue BF16 matrix products with FP32 accumulation. Each block owns a K subset and a set of output rows. The staging layout and fragment ownership are explicit in the native source.

Reduce through the block cluster

After finishing its local reduction, each block writes its FP32 accumulator into shared memory. The cluster synchronizes so peer partials are available. Blocks then read distributed shared-memory partials in a defined rank loop and write their assigned final output elements in BF16. A final cluster barrier protects the shared memory's lifetime. This replaces a separate global reduction path with cluster cooperation, whose own synchronization cost has to be repaid by better execution.

Candidates 032 and 033

Candidate 032 checked at about 1.2712. Candidate 033, with the selected deeper cluster QKV configuration, checked at about 1.2771 and independently scored 1.2888. These are complete projection results. The cluster path was selected for the appropriate small-token QKV range, while the other projection families retained their chosen methods. The source therefore joined a specialized native kernel with vector, Triton, and library routes.

The next official incumbent

The 1.2888 result became the incumbent for the C2 round. It was archived with source and result evidence. I wanted at least 1.30 and then 1.32, so this result was progress rather than completion. It also came with a measured map of remaining gaps: QKV at other token ranges, small pooler behavior, larger router activation rereads, and dense paths already close to library performance.

What this design did and did not prove

The selected kernel showed that a carefully matched custom CUDA design could help this skinny QKV regime. It did not prove that custom CUDA beat cuBLAS at larger token counts, or that adding more stages would continue to improve it. Later sweeps tested deeper variants and did not find another useful win. This is the boundary I carry forward: preserve the successful region, keep its mechanism understandable, and make the next experiment earn a separate region rather than extrapolate from one victory.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 41

Building our remote H200 laboratory#

The GPU ran remotely, so a useful experiment needed more than code: a matching environment, a working benchmark, recovered logs, and a machine that wasn't left billing idly.

A local Mac controls a remote GPU

The editing workspace was on the Mac, while the performance experiments needed an H200. The lab synchronized source and the supplied problem files to a Runpod machine, then ran correctness and timing through SSH. Results were copied back into the workspace. This separated authoring from execution without changing the submitted computation. The submission itself did not contact the network. Remote orchestration happened outside the evaluated source, as ordinary experiment infrastructure.

Match the judge image and verify readiness

The pod used the specified SGLang development image and a CUDA-compatible H200 host. A running status alone was not enough to establish a useful lab. I checked SSH access, the visible GPU, framework versions, and machine settings before trusting performance work. The recorded environment used PyTorch two point fourteen point one with CUDA thirteen and Triton 3.8. Those details helped keep the kernel and compiler behavior close to the official environment, while still requiring official checks.

The lab benchmark and correctness runner

The benchmark captures candidate and baseline calls as graphs, rotates independent weight states, supplies fresh inputs, flushes L2, swaps execution order, and records paired ratios. The correctness runner uses the supplied reference, difficult official cases, scored shapes, outliers, neighboring token counts, and fresh graph replays. Keeping these tools stable matters. Changing the benchmark to make a candidate look better would invalidate the investigation. I changed candidate methods while preserving the measurement and reference contracts.

The profiler limitation

Nsight Compute hardware counters were unavailable on the Runpod host because the required performance-counter permission was not exposed. The handoff records the corresponding permission error. I could still use available timing and activity profiling, but I could not pretend to have counter evidence for occupancy stalls or memory behavior that had not been measured. This limited some causal diagnoses. I could establish that a candidate lost and offer a plausible explanation, while labeling the explanation as a hypothesis requiring deeper profiling.

The cost of leaving allocation idle

The pod's quoted compute price was five dollars and 29 cents per hour. The recorded bill for the first session reached about 37 dollars and fifty cents, including allocation after productive searches had ended. The pod was already stopped when work resumed later. That was an operational cost, not a hidden amount of successful optimization. A better workflow needs evidence synchronization and verified compute shutdown, especially when a turn can be interrupted or a long job completes unattended.

Resume with the evidence intact

The stopped disk retained exploratory logs. A later restart was initially rejected because the account lacked credit; I added funds, and the pod became reachable again with a new SSH port. The recovered broad library-search log showed a time limit rather than a completed breakthrough. This sequence belongs in the story because operational state affects what can honestly be claimed. Prepared source, a scheduled job, a completed run, and a recovered result are four different states.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 42

The packed pooler experiment#

For the pooler, I tried keeping each output tile's weights together. The original values stayed the same; their addresses changed.

The C2 hypothesis

The pooler had a large weight stream and a small output-tile grid at low token counts. The C2 packed design rearranged the original BF16 weights into contiguous regions for each output-feature tile. The hypothesis was that a block could stream its owned K-by-output region efficiently and accumulate the output directly, avoiding a split reduction for this regime. Preparation paid the lossless repacking cost once, leaving fresh activations for the timed kernel.

What the packed tensor means

Start from a stored weight with N rows and K columns. Group N into blocks of BN output features, then reorder each block so K steps contain adjacent output columns. The packed representation has a conceptual shape of N divided by BN, then K, then BN. Every original BF16 value has a new address but the same value. The kernel selects its output tile, streams along K, and uses a dot product to update a FP32 accumulator for its token tile.

Tune a balanced tile

The search varied token tiles, output tiles, K widths, warps, and pipeline stages. The included pooler choices used output tiles of 32 features. For the smaller token range, the selected configuration used a sixteen-row token tile and K width 512; for the next range it used a 32-row token tile. The stage count differed. These values balanced streaming, grid size, shared-memory use, and dot-product efficiency. The search results, rather than a single theoretical rule, selected the final combinations.

Repeated pooler gains

Local paired screening and retesting found useful gains in the small-token pooler cases. The exact ratios varied across runs, but the selected regions repeatedly improved by several percent rather than rely on one isolated replay. Finalists were checked on fresh inputs and neighboring shapes. That made the packed path worth including in a combined candidate. It did not by itself establish a 1.32 complete score, because the pooler occupies only part of each of several loads.

Candidate 034 did not beat the official incumbent

The first combined packed-and-library candidate, 034, passed its official check at about 1.2852. That was below the 1.2888 scored incumbent. I did not turn the local pooler win into a claim of official aggregate progress. Instead, the per-shape evidence remained useful while the combined design needed further work. This is exactly why complete official evaluation follows local selection: a promising part can coexist with an unimproved whole result.

What I carried forward

The packed pooler route remained a measured candidate for selected small-token regions, while the larger pooler and QKV ranges invited more explicit library-layout choices. The next round expanded cuBLAS Lt orientations and algorithms. I preserved the original incumbent and created new immutable snapshots, so a losing combined check could not overwrite the verified best. This allowed exploration to continue without confusing the current experiment with the current proven result.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 43

Exploring cuBLASLt orientations#

An equivalent matrix product can look different to a library. I tried layouts and algorithm choices, then checked the neighboring sizes too.

The library sees a layout, not just an equation

cuBLAS Lt chooses among implementations based on matrix descriptors, dimensions, types, strides, workspace, and algorithm settings. Two mathematically equivalent products can therefore lead it to different execution paths. My C3 and C7 work explored the original weight orientation, a lossless transpose, and swapped products returning a transposed output view. The objective was to expose a faster valid algorithm for a shape family, while preserving the original values and the required accumulation behavior.

Constrain the reduction scheme

The extension requested FP32 compute and restricted split reductions to the compute-type scheme. Lower-precision output-type reductions were excluded from the intended search. This matters because a BF16 output tensor does not tell me how intermediate split results were combined. Workspace was allocated for valid algorithms, and algorithm checks screened configurations before execution. Those safeguards did not replace numerical tests, but they aligned the search with the required precision rather than let a cheaper reduction silently change the problem.

Heuristics, splits, and retesting

The search enumerated library heuristics and additional validated split choices for several orientations. It timed captured calls against the incumbent, then retested finalists with more repetitions and neighboring shapes. Algorithm positions in a heuristic list can differ with shape, so a configuration that looks strong at one M can choose an unsuitable path nearby. This was one reason to inspect neighboring behavior rather than treat an index as a universal algorithm identity.

Useful regions emerged

The expanded orientation search found useful local improvements for the larger pooler, QKV around 32 tokens, and the large unpooler range. Some nearby configurations looked good at a scored shape but regressed badly at a neighbor and were rejected. The selected combined source retained different orientations and indices for the supported regions. These were library dispatch improvements, not a new approximation to the matrix product. Their per-shape contribution is visible in the later official result and the shape atlas.

Add alignment-based fallbacks

Swapped-output products make M the leading dimension of the allocated transposed output. For the selected routes, token counts not divisible by eight retained the original path instead of forcing an unfavorably aligned swapped layout. The pooler also retained the original transposed path through an intermediate token range where the newer plan regressed. These choices made the dispatch more robust across neighboring shapes. They expressed layout and range conditions rather than pretending the faster orientation was universal.

Candidate 035 stayed around the incumbent

The layout-focused candidate 035 passed its official check at about 1.2860. The local bundle had looked promising, but this official result still did not beat the earlier scored incumbent. I preserved the useful shape-level evidence and continued looking for another complementary gain. This is an important narrative point: the last step over 1.30 did not come from declaring the layout search successful enough. It required the later single-token streaming improvement and another independent evaluation.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 44

Lossless compression: fewer bytes, slower execution#

Reading fewer bytes sounded useful. The lossless compression attempts also had to decode those bytes, and that extra work overwhelmed the savings in the tested versions.

A tempting response to the bandwidth limit

The large dense weights made a lossless compression experiment plausible. If fewer stored bytes could be read and reconstructed cheaply inside the kernel, a memory-bound projection might improve without changing required values. The word lossless is essential: the original BF16 bit patterns had to remain recoverable. These experiments were not attempts to replace the dense weights with lower-precision approximations. They explored representations whose decompression cost might be repaid by a smaller external-memory stream.

Blocks, metadata, and escapes

The compression branches explored different ways to exploit structure in BF16 values. Some packed portions of the exponent or mantissa information, used block metadata, or kept escape paths for values that did not fit a compact local representation. Such schemes add addressing logic and integer instructions. An escape path can require another original-weight read. A block's metadata must also be fetched and interpreted. The compressed byte count therefore describes only part of the new execution path.

Register pressure can erase a bandwidth saving

A direct projection already needs input values, weight values, products, and reduction state. A decoder may add bit masks, shifts, offsets, metadata, and intermediate representations. That can increase registers per thread and reduce residency. It may also serialize instructions before the arithmetic can proceed. These are plausible mechanisms to inspect when compression loses, but timing alone does not prove their exact share. The measured fact in my branch was simpler: the tested exact encodings were too slow for the selected regime.

The C4 through C6 outcomes

The lossless branches preserved the required arithmetic in tested cases but performed far below the incumbent in the local searches. Some configurations were less than half as fast as the existing route. That is not a small tuning loss to hide within measurement noise. The byte reduction was overwhelmed by the complete execution cost. I therefore rejected the compressed methods for the winning bundle and kept the reports as evidence of a tested hypothesis that did not work in these implementations.

The model needs a new cost term

A bandwidth-only prediction compares original bytes with compressed bytes. The actual compressed kernel also pays decode instructions, metadata movement, possible escapes, and any resource-efficiency loss. If that added cost exceeds the saved memory time, total execution gets worse. This is a useful refinement of roofline reasoning: reducing one resource demand can introduce another limiting demand. I do not abandon modeling after a failed prediction. I update the model to include the work the implementation actually performs.

The decision was to move on

There may be other lossless encodings or architectures where this idea helps. My evidence does not disprove compression in general. It does establish that these branches did not earn a place in the current dispatch. The rational next decision was to preserve the negative results and focus on better-supported opportunities, including launch sizes, cache policies, and library orientation choices. An honest optimization story can end a technically interesting chapter with a rejection rather than force every idea into the final design.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 45

Router redesigns that did not help#

These router redesigns moved tile, read, and combine costs around. None of these tested branches earned a place over the atomic incumbent.

A smaller matrix can still be difficult

After the atomic-router improvement, the remaining router calls were short enough that overhead and repeated input reads mattered. The C2 round tried several reorganizations rather than assume the existing method was optimal. These included changing where high and low weights were interleaved, narrowing output tiles, and packing partial results for a fused final reduction. Each design preserved the intended high-plus-low computation. The open question was whether a better access or synchronization pattern would beat the complete atomic path.

C8 interleaves along the reduction dimension

One design placed high and low contributions along an expanded K-like dimension and loaded the corresponding input value for each pair. The attraction was a unified dot-product path. The cost included a new operand layout and duplicated input use. Local screening found the tested configurations slower, increasingly so in some larger token cases. The result showed that simplifying the source expression did not simplify the generated execution enough to help. I retained the report and excluded the method from the final bundle.

C11 explores smaller output tiles

Smaller output tiles can create more independent blocks from the same router width. That sounds attractive when the output grid is small. But each tile may reread activations, and a narrower dot-product shape can be less efficient. The C11 sweep varied token, output, K, split, and launch choices. Its best tested finalists still did not provide a useful improvement over the incumbent. The search narrowed an opportunity rather than support the simple rule that more blocks must be faster.

C14 packs partials for an in-kernel reduction

Another branch wrote split partials into a tile-local packed buffer and let the last completing block reduce them. It explored different partial layouts, including which dimension made the split contributions contiguous. This targeted the combination path and memory layout together. But it still paid partial writes, counter synchronization, last-block reads, and a final sum. Sampled results were clearly slower than the atomic incumbent, so the search was stopped rather than spend more GPU time on an unpromising branch.

What the rejection establishes

The measurements establish that these implementations did not improve the tested regimes. They do not prove every possible version of the ideas must lose, or isolate every hardware reason without counters. That distinction lets the evidence remain useful. A future design could remove a specific cost and deserve another test. For the current result, however, none of these router branches entered the winning submission. Keeping a correct, measured incumbent was better than promoting a more elaborate but slower path.

A new router idea remained only a proposal

Later source preparation explored a block-cluster router using distributed shared-memory reduction, adapting ideas from the useful QKV kernel. In the record available to this written companion, that later design was prepared but not validated as a winning experiment. It therefore belongs on a proposed branch, not on the path that achieved 1.30. Showing this boundary is part of an honest story: writing code and having a plausible mechanism are not the same as completing correctness and performance evaluation.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 46

Advanced features that still lost#

New hardware features and deeper pipelines were worth trying, but they still had to win a complete timing. Here are the versions that didn't.

C9 tests tensor-memory transfers

The C9 branch used Triton tensor descriptors to test Hopper tensor-memory transfers for dense matrix products. It explored tile sizes and split choices, with FP32 accumulation and the required BF16 output. This was a concrete TMA experiment, not merely a suggestion in a strategy document. The full design still needed a good schedule, sufficient reuse, efficient arithmetic, and an acceptable reduction path. A more capable load mechanism could not guarantee those other requirements were satisfied.

The TMA result did not beat the incumbent

The tested TMA configurations failed to provide a useful local win for the target shapes. Some larger pooler configurations came close, while other shapes were substantially slower. The search did not establish that tensor-memory transfer hardware is ineffective. It established that this particular implementation and parameter family did not outperform the existing route. A more complete persistent or warp-specialized design would be a different experiment and should not be credited with a result that had not been measured.

C12 deepens the cluster search

The deeper QKV cluster sweep varied stage counts, warps, and K widths across additional compiled instances. The chosen incumbent configuration remained difficult to beat in its useful small-token regime. Some variants looked close enough to be noise, while many were worse. Adding pipeline depth can consume shared memory, change residency, and exceed the useful overlap available in a short per-block loop. The result supported keeping the existing path rather than assume the next level of staging would produce another milestone.

C13 refines the packed pooler

The pooler refinement revisited output tiles, K widths, warp counts, stages, and load cache policy. The best choices were broadly consistent with the already selected packed path. A .cg variant could look slightly better in one retest without establishing a clear incremental aggregate gain. This is where search discipline matters: replacing a validated configuration based on a tiny noisy difference can create churn without progress. I kept the clearer wins and treated marginal alternatives as uncertain.

Why advanced designs can disappoint

A feature can improve one part of a kernel while another part becomes limiting. TMA may improve movement but leave too few useful tiles. More stages may improve overlap but reduce resident work. A wider tile may improve reuse but lower grid parallelism. These examples are mechanisms to examine, not universal diagnoses of every failed variant. The useful habit is to compare the complete resource tradeoff with the practical limit suggested by the workload, then let correctness and paired timing decide.

The search returns to a simpler opportunity

After these complex branches, the next useful gain came from the single-token vector path: launch width and cache behavior. That is not a moral that simple code always wins. It is a reminder that the best remaining opportunity can be a concrete mismatch in an existing method, even after exploring sophisticated hardware mechanisms. I returned to the measured path, changed a small set of meaningful parameters, and found a complementary gain that the combined candidate could use.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 47

The final single-token streaming gain#

The next gain came from revisiting an existing one-token path: how many warps it used and how it treated the large weight stream in cache.

One token still streams millions of weights

The direct single-token kernel loads an input vector, reads weight rows, multiplies in FP32, and reduces each row. It does not reuse a weight across many token rows because there is only one token. Efficient streaming and lane distribution are therefore central. The C10 sweep revisited output rows per block, warps per block, and cache policy. This was a focused hypothesis about how the existing reduction used the memory system, not a new approximation or a larger model change.

Change the warp distribution

A block with more warps distributes the row's values across more lanes. That can change instruction count per thread, register use, reduction structure, and residency. The selected sixteen-warp choice helped the QKV and pooler single-token cases in repeated local measurements. Other layers did not show the same clear opportunity, so the final plan did not apply the change everywhere. The later 32-warp check is a useful counterexample to the idea that more parallel lanes must continue to improve performance.

Stream weights with the selected cache policy

The winning single-token variants used .cg weight loads. The input vector's reuse differed from the much larger weight stream, so the kernel could treat those access patterns differently. Earlier policy combinations also exposed assembly constraints: a configuration that combines unsupported load and eviction attributes can fail before performance is measured. The corrected sweep separated valid policies and launch choices. That preserved the distinction between an invalid generated instruction sequence and a valid kernel that happened to be slow.

Retest the useful gains

The single-token sweep found around six to seven percent local QKV improvement and around five percent local pooler improvement in promising configurations. A later full paired comparison showed roughly six percent and just under five percent for those selected shapes. The exact ratios varied, as expected, but the effect was useful enough to combine with the packed pooler and library-layout work. Fresh-input correctness and graph replay checks remained in place. The arithmetic still used the original BF16 weights with FP32 reductions.

Only two selected dispatch entries change

The combined source changed the first QKV and pooler dispatch entries to the streaming vector method with one output row per block, sixteen warps, and the selected cache policy. Other token ranges retained the measured cluster, packed, or library routes. This small source change mattered because it affected the single-token load, which counts equally with the other loads in the final geometric mean. A modest gain in that load could complement gains already present at larger token counts.

Candidate 036 becomes the next complete test

The new vector path was added to the combined layout candidate, producing experiment 036. The full local correctness runner passed across its tested seeds, official difficult cases, scored shapes, outliers, neighbors, and fresh graph replays. The complete paired benchmark then compared this bundle with the immutable 033 incumbent. This sequence matters: the final milestone came from a freshly validated combination, not from adding two component percentages on paper and assuming the resulting source was correct.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 48

Read the winning implementation end to end#

The winner was a collection of selected routes. I'll follow preparation, dispatch, fallback, and output so the final source reads as one program.

Preparation builds static representations

The winning source starts from contiguous original weights. It forms the router's high and low BF16 parts, transposes and packs the representations required by its methods, prepares dense transposes, and creates packed dense tiles. It also compiles the cluster and library extensions and reserves library workspace. No evaluated activation is available in this static phase. The later lean-preparation experiment removed unused layouts, but the officially scored 1.3012 source is the exact original 036 snapshot.

QKV chooses three regimes

At one token, QKV uses the streaming vector reduction. Through sixteen tokens, it uses the selected custom CUDA cluster path. Through 32, it uses an explicit cuBLAS Lt orientation and algorithm choice, with the alignment fallback where required. Larger token counts use the retained library dispatch. Each region corresponds to a different geometry and measured comparison. The final source is therefore a compact policy over specialized methods rather than an assertion that one kernel is best for the whole QKV matrix family.

Pooler uses packing and library layouts

At one token, the pooler uses the new streaming vector path. The small-batch ranges use the direct packed-weight Triton kernel with different token tiles and stage choices. An intermediate range retains the original transpose route, protecting nearby-shape performance. The larger selected ranges use explicit library orientations. This is where the C2 investigation contributes most of its layout work. Every route preserves the original BF16 values and produces the required dense output format.

Output and unpooler keep complementary paths

The output projection retains a direct vector method at one token and library dispatch otherwise. The unpooler uses a vector at one token, tiny GEMM for the small local slice above that, a transposed product through the medium range, and a selected library orientation for the large range. Some official output-projection cases still regress slightly relative to baseline. I keep those visible. The complete score is the result of all selected paths, not a promise that each individual shape is faster.

Router keeps exact FP32 output behavior

At one token, the router uses a direct reduction over its original FP32 weights. Larger token ranges use paired high-and-low BF16 products with FP32 accumulation and atomic combination. The method allocates a freshly zeroed FP32 output, so a new invocation does not inherit previous sums. The split and tile parameters change by token range. The complete call is captured and timed, including the initialization that supports its correctness.

Dispatch, fallback, and return

The timed entry point reads the actual input shape and the supplied layer configuration, chooses the first matching plan region, checks method applicability, and invokes the selected function. It returns a GPU tensor of the required shape and dtype. The source does not inspect reference outputs or cache evaluated answers. The exact winner's hash and official job identifier are saved in the snapshot. That identity matters more than a filename saying best, because a future edit may no longer be the source that earned the recorded score.

← Previous Return to worklog Next →

IV · The H200 investigation and the finish / 49

Local 1.3234×, official 1.3012×#

My local 036 benchmark looked like 1.32×. The judge recorded about 1.30×. Both measurements stayed in the notes with their actual labels.

The paired incumbent comparison

The complete local bundle was compared against the immutable 033 incumbent with eight weight copies and 120 repetitions under the recorded benchmark conditions. The aggregate local gain was about 1.0241× the incumbent. That was a useful controlled comparison: both paths ran on the same machine under paired conditions. It suggested the bundle was worth official evaluation. It did not imply the same multiplicative gain would transfer exactly onto an earlier official run taken under different conditions.

The direct local SGLang comparison

The same candidate also measured about 1.3234 against the local SGLang baseline. That was an encouraging number because my later target was 1.32. But the local benchmark was not the official result. It matched important features of the judge while still differing in machine metadata and surrounding implementation details. The right response was to submit and inspect the official evidence, not declare the target guaranteed from a nearby local point estimate.

The official check

The official check passed and returned about 1.3020. This validated the source under the competition's own correctness and timing path. It also showed that the local 1.32 result did not transfer at that magnitude. The candidate still exceeded the original 1.30 milestone. I preserved the check result as its own artifact and kept the independently scored submission separate, rather than round one result into the other.

The independently scored outcome

The scored job also passed and returned 1.301203, rounded here to 1.3012. This is the result that entered the projection leaderboard. The close check and scored values supported a result near 1.30, while neither established the later 1.32 target. The official point estimate and its interval are shown together. My story ends the accomplishment at the number the evaluation actually recorded.

What could explain the local-to-official difference?

Possible contributors include driver and machine differences, measurement overhead estimates, exact baseline dispatch behavior, memory placement, and ordinary run variation. The saved records show that local and official environments were not identical. They do not isolate a single proven cause for the entire gap. I therefore present the discrepancy as a measurement lesson rather than invent a precise diagnosis. A model that explains why differences can occur is useful, but a causal claim requires an experiment that controls or measures the suspected factor.

The correct decision

The decision was to promote the exact officially scored source as the new best, preserve the earlier source, archive all evidence, and state the remaining gap honestly. The source contained real layout, dispatch, and launch improvements. The competition result crossed 1.30, and the next target still needed more work. That is a stronger outcome than a larger unsupported claim, because another person can inspect and reproduce the exact achievement and identify what remains to be improved.

← Previous Return to worklog Next →

Postscript · A later result arrived during production / 63

Postscript · The official 1.3234× result#

The later candidate really did score 1.3234× officially. It used different code from the earlier local run with the same rounded number. Here's what changed.

A new official result arrives

There is a later development. While this written companion was being produced, the project recorded a newer projection submission that independently scored 1.32343. Its official check had returned 1.3218. Both passed, and the live board showed the new result at rank 66 at the time of my read. This clears the 1.32 point-estimate target discussed in the earlier chapters. The main film reconstructs the preserved 1.30 milestone; this postscript records the subsequent step with separate evidence.

Same rounded number, different evidence

The rounded number 1.3234 may look familiar. Earlier in the article, that was a local measurement of candidate 036, whose official score was lower. The newer 1.3234 is an official result for candidate 039. They are different candidates and different evaluations. The earlier local result did not retroactively become official. A new scored job supplies new evidence. This distinction is exactly why I preserved source identities, measurement labels, and job records throughout the story.

Three additions in the newer source

The saved build script describes three measured additions on top of the lean-preparation base. It uses programmatic dependent launch for the router's initialization and compute path. It selects expanded library configurations for QKV at 128 tokens and pooler at 256. It also adds a different skinny QKV kernel using tensor memory transfers, warp-group matrix operations, and cluster reduction for the selected nine-through-32-token range. This is not the same TMA implementation that lost in the earlier round. A failed earlier design did not prove every more complete later design would fail.

PDL preserves a correctness boundary

The router still needs a fresh zero output before atomic additions. Programmatic dependent launch gives the following kernel an opportunity to begin independent work before the zero-fill predecessor has fully completed. The router can perform loads and product work, then waits before its output atomics depend on the initialization. This does not make the dependency disappear, and concurrent execution is opportunistic rather than guaranteed. The source explicitly signals dependent launch and waits at the dependent-write boundary. The relevant NVIDIA documentation is linked with the postscript sources.

The later skinny kernel is a new design

The new skinny path uses a producer and consumer organization, tensor-map loads, shared staging barriers, warp-group matrix accumulation, and cluster cooperation to finish split contributions. Its instantiated configurations use three-way split work and an eight-stage pipeline for the selected token tiles. The complete design differs from the earlier simple TMA descriptor experiment and from the old warp-level MMA QKV path. The source supports the explanation of its mechanism, while the bundled official result establishes the new submission's aggregate outcome. Assigning an exact isolated causal percentage to each addition would require the corresponding controlled comparisons.

Library indices and neighboring routes

The source also selects expanded cuBLAS Lt enumerations at particular matrix shapes where those indices were measured. Nearby shapes retain earlier routes, and swapped-output alignment has a fallback. The index is an entry in a shape-specific enumeration, not an architecture-independent algorithm name. This explains why the newer source guards its selected choices rather than applying one integer across an arbitrary token range. The implementation remains subject to the competition's final review and neighboring-shape expectations, just like the earlier submissions.

The new result and its uncertainty

The newer scored point estimate is 1.323430, with a reported interval of roughly 1.3036 to 1.3433. Its independent check is also above 1.32. Relative to the earlier 1.3012 scored point, the new ratio is about 1.0171. These are recorded outcomes from separate runs, so I retain the uncertainty rather than claim every last decimal is a precisely isolated causal effect. The target crossing itself is visible in the passing official result and the board.

An updated ending, with the history intact

The main investigation still explains how the project reached 1.30. This later result extends that history rather than rewriting the earlier local measurement or erasing rejected experiments. It reinforces the central lesson: preserve what an experiment actually tested, and let new implementations create new evidence. The documentary includes the later source, both result files, and the dated board read in a separate postscript folder. You can audit the original milestone and the subsequent 1.32 target crossing without confusing their code or measurements.

← Previous Return to worklog Next →