← All notes and experiments

Making a ternary reasoning model fast in vLLM

Latest published comparison. The opening tables report the current checkpoint; later sections explain the earlier stages that led to it.

By the CUDA checkpoint in chapter 9, four hand-written kernels handled the fastest PQ2 paths: one-token generation, small batches and prompt processing. Together they contained about 930 lines of CUDA with inline GPU assembly.

I wanted to find out whether I could keep that speed in code that was easier to work with. Plain Triton handled the single-token algorithm. Small batches and prefill needed more control over registers and memory; Gluon, Triton’s layout-explicit dialect, provided it.

All four paths now have faster Triton or Gluon replacements, enabled by default for this model on the tested RTX 3090-series GPUs. CUDA remains available for other shapes and hardware. The final comparison below measures the complete serving path, including a Triton replacement for one small unquantised layer.

Measurement note. The paired CUDA/Triton comparison uses one RTX 3090 Ti at 300 W, PQ2_0 Ternary Bonsai 2 27B, FP32 recurrent state, an INT8 KV cache, plain decoding and no prompt caching. The two servers run one after the other on the same card, with --max-num-seqs 16 and fresh compilation. These are recorded GPU measurements; the website build does not rerun them.

For a short read, start with the final results and what I took from the work. The sections between them explain the kernel changes; expand the code when you want the implementation detail.

Final results

Plain generation

Request performance

Rates above one request are aggregate output across users, not the rate seen by each user.

Decode uses 512-token prompts and 1,024 generated tokens per request, with medians of two or three runs. The percentages follow the recorded results; displayed values are rounded. The old defaults here are the paired CUDA baseline, rather than the separate llama.cpp comparison in chapter 9.

Prompt processing

Request performance

The 128-token row measures time to first token (TTFT); lower is better. The other rows report prompt-processing throughput; higher is better.

Answer checks

Quality check: GSM8K

Scores count correct answers out of 128. The columns compare the previous CUDA defaults with the final Triton/Gluon defaults.

Both runs differ on the same question, 942, which also changed in the earlier single-token A/B. I traced this to a near-tie affected by rounding. Every other answer is identical. These 128 questions provide a useful check, but they do not establish equivalent quality across all reasoning tasks.

Nsight traces with the final defaults show no plugin CUDA kernel at 1–16 token rows per matrix multiplication or during prompt processing. This describes the plugin’s kernel choices. The runtime still uses NVIDIA’s CUDA platform and other CUDA-backed libraries.

RTX A4000: deployment results

The A4000 brings the measurements back to the card I wanted to use in the first place: serving the model from 16 GB of GPU memory.

Request performance on the RTX A4000

Prefill measures input tokens processed per second. Decode measures output tokens per second; concurrent rates are the total across all requests. Additional tuning made little difference in these runs.

Three findings shaped the work:

• The single-token gain came from the algorithm. I could express the same calculation in plain Triton without keeping the hand-written CUDA kernel.
• Small batches needed more control over the register layout. Gluon let me place weight bytes directly in the registers used by the matrix instructions. Extending the kernels to 9–16 tokens then produced the largest serving gains.
• The measurements needed closer checks. Replays ran at a GPU clock roughly 30% below serving, and one A/B script failed to set the switch it was meant to test.

Where the CUDA gains came from

I started by comparing each CUDA kernel with the Triton version it had replaced. The earlier measurements helped separate changes to the algorithm from changes to how data reached the instructions:

Component test: earlier CUDA paths and their predecessors

Two gains came mainly from different algorithms, which I could express in Triton. The other two depended on which thread held each group of bytes. Reproducing those arrangements needed more control than the plain Triton path provided.

One token at a time: plain Triton is enough

PQ2_0 stores each weight as a 2-bit code (0, 1, 2 for -1, 0, +1) packed four to a byte. The old Triton kernel turned each byte into four signed bytes with a PRMT lookup and then ran DP4A, about ten integer instructions per weight byte, which made it ALU-bound.

The CUDA kernel instead stores the activations in plane order: the activation for weight 4·byte + s goes to position s·32 + byte. Then plane s of a 32-bit word of weights, (w >> 2s) & 0x03030303, is four codes that line up with four consecutive activation bytes, and one DP4A multiplies them with no decoding.

The ternary offset of -1 becomes one subtraction per 128-weight group: dot(codes, q) - sum(q).

@triton.jit
def _dp4a(a, b, c):
    return tl.inline_asm_elementwise("dp4a.s32.s32 $0, $1, $2, $3;", constraints="=r,r,r,r",
                                     args=[a, b, c], dtype=tl.int32, is_pure=True, pack=1)

@triton.jit
def _single_planes(Q, S, T, W, WS, Y, N: tl.constexpr, K: tl.constexpr,
                   BN: tl.constexpr, BC: tl.constexpr, STAGES: tl.constexpr):
    G: tl.constexpr = K // 128                       # 128-weight groups per row
    rows = tl.program_id(0) * BN + tl.arange(0, BN)
    offs = tl.arange(0, BC)
    j = tl.arange(0, 8)                              # 8 int32 words = 32 code bytes per group
    acc = tl.full((BN, BC), 0, tl.float32)
    for base in tl.range(0, tl.cdiv(G, BC), num_stages=STAGES):
        g = base * BC + offs
        w = tl.load(W + rows[:, None, None] * (K // 16) + g[None, :, None] * 8 + j[None, None, :], ...)
        ws = tl.load(WS + rows[:, None] * G + g[None, :], ...).to(tl.float32)
        qb = Q + g[:, None] * 32 + j[None, :]        # plane-ordered activations as int32 words
        d = _dp4a(w & 0x03030303,        tl.load(qb)[None],      tl.zeros((BN, BC, 8), tl.int32))
        d = _dp4a((w >> 2) & 0x03030303, tl.load(qb + 8)[None],  d)
        d = _dp4a((w >> 4) & 0x03030303, tl.load(qb + 16)[None], d)
        d = _dp4a((w >> 6) & 0x03030303, tl.load(qb + 24)[None], d)
        dot = tl.sum(d, 2) - tl.load(T + g)[None, :]  # the ternary -1, once per group
        acc += dot.to(tl.float32) * tl.load(S + g)[None, :] * ws
    tl.store(Y + rows, tl.sum(acc, 1), rows < N)

On all 257 projections of a real decode step this ran in 8.02 ms against 8.72 ms for the CUDA kernel and 10.94 ms for the old Triton one. In serving it was +2.9% tok/s at one request, with GSM8K answers identical on all 128 questions. The integer products are the same as CUDA's; only the order of the FP32 group sums differs (relative RMS at most 1.6e-4).

The same approach became less efficient with several tokens. Adding a token dimension costs four DP4A instructions per weight word for each token, making arithmetic throughput the limit from three tokens. It was 1.48x slower than the tensor-core kernel at M = 3 and 2.5x slower at M = 8. For these batches, the tensor-core path remained the better choice.

Small batches: make the register layout explicit

In the Triton path I was replacing, tl.dot chose the arrangement of its operands. It wrote the weight tile to shared memory, then read it back with ldmatrix. The two-bit codes first had to be expanded into four INT8 planes.

That required four times the bytes and instructions of the CUDA path. CUDA loaded 8 code bytes per lane directly into the A-fragment registers of mma.sync.m16n8k32, then extracted each plane with a shift and a mask. This staging overhead had limited the old Triton small-batch kernel to about 500 GB/s.

Gluon lets me specify the distributed layout of a tensor, manage shared memory and call mma_v2 and async_copy directly. Triton 3.7.1 includes the Ampere support needed here. I wanted to check whether those tools could reproduce the CUDA kernel’s loads.

reg_bases  = [[0, 1], [0, 2], [8, 0], [0, 16]]       # k+1, k+2 (one 32-bit register), row+8, k+16
lane_bases = [[0, 4], [0, 8], [1, 0], [2, 0], [4, 0]] # lane = (gid, tid): tid -> k+4, k+8; gid -> rows
warp_bases = [[16, 0], [32, 0]]

def layouts(warps, rt):
    bm = 16 * warps * rt
    mma = ttgl.NVMMADistributedLayout(version=[2, 0], warps_per_cta=[warps, 1], instr_shape=[16, 8])
    words = ttgl.DistributedLinearLayout(
        reg_bases=[[0, 1, 0], [8, 0, 0]] + [[16 * warps << i, 0, 0] for i in range(int(math.log2(rt)))],
        lane_bases=[[0, 0, 1], [0, 0, 2], [1, 0, 0], [2, 0, 0], [4, 0, 0]],   # tid -> t, gid -> rows
        warp_bases=[[1 << i, 0, 0] for i in (4, 5, 6, 7) if (1 << i) < 16 * warps],
        block_bases=[], shape=[bm, 2, 4])                                       # [rows, h, t] int32
    return bm, mma, words

@gluon.jit
def _bytes(p):                       # int32 words [rows, 2, 4] -> int8 bytes [rows, 32]
    b0 = (p & 0xFF).to(ttgl.int8)
    b1 = ((p >> 8) & 0xFF).to(ttgl.int8)
    b2 = ((p >> 16) & 0xFF).to(ttgl.int8)
    b3 = ((p >> 24) & 0xFF).to(ttgl.int8)
    x = ttgl.join(ttgl.join(b0, b2), ttgl.join(b1, b3))
    return ttgl.reshape(x, [x.shape[0], 32])

# in the kernel:
w = ttgl.load(wp + g * 8)                                                  # LDG.64 per lane
wb = ttgl.convert_layout(_bytes(w), A_L, assert_trivial=True)             # no data movement

With prefetching in place, handling two row tiles per warp became useful, as it had in CUDA. On captured model inputs, the Gluon kernel improved every tested shape. Replaying a complete decode step gave a 5-15% gain over the production kernels.

Measure at the clock speeds used in serving

The 5-15% replay gains became 2% in serving at four requests, with no measurable gain at eight. I first checked launch overhead. Nsight showed the same 0.08 us idle gap before each matrix multiplication, and split-K reduction launches added at most 0.1-0.3 ms per step.

I then measured the same comparison in three ways:

Component timing: replay and serving traces

Two differences deserved attention. Repeating all layers of one shape could favour kernels suited to that repetition. The GPU could also run at a different clock. I sampled nvidia-smi every 100 ms during each workload:

GPU profile

The replay kept the GPU at the 300 W cap with repeated memory-bound matrix operations. Serving interleaved those operations with normalisation, DeltaNet updates, attention and sampling. Average power fell and the clock rose. The CUDA kernels benefited more from that higher clock than the Gluon kernels did.

For gate/up at four requests, CUDA went from 81 us in the replay to 64.6 us in serving, and Gluon only from 72.4 to 64.3 us. At serving clocks both sat near the same ~720 GB/s, where the access pattern rather than the instruction count is the limit.

The tuning tables I had selected so far reflected the replay’s clock rather than the clock reached in serving.

A power proxy for the rest of the model

Testing every candidate in the full server would be slow: startup takes minutes, and serving rates vary at the 1-2% level that separates candidates. I needed a replay that better represented the pauses between projections.

I added a low-power spin kernel after every projection. torch.cuda._sleep(cycles) keeps one thread waiting for a fixed number of cycles. It occupies part of the GPU timeline while drawing very little power, standing in for the lower-power work between matrix operations.

I estimated the spin length from a serving trace. In each decode step, I counted the time outside PQ2 matrix multiplication and input quantisation as “other work”:

GPU profile and proxy calibration

I calibrated the cycle counts at 1 and 4 requests and interpolated the others, using 25-55 thousand cycles from 2 to 16 tokens. This is about 1.5 cycles per nanosecond of measured other work, less than the ~2 GHz clock alone would suggest. A spinning thread draws less power than the kernels it represents, so I calibrated against the power state rather than trying to match the timeline exactly.

# real captured decode rows: 257 projections in model order, real prepared weights, real rotated activations
layers, wf = load(CAPTURE)
xs = [l["x"][:M].contiguous() for l in layers]

fns = []
for l, x in zip(layers, xs):
    fns.append(lambda l=l, x=x: prism.pq2_matmul(x, l["w"], prepared=True))   # the kernel under test
    fns.append(lambda: torch.cuda._sleep(spin))                               # stands in for norms, DeltaNet, ...
g = graph(fns)                 # one CUDA graph: the whole step, replayed sustained for several seconds
sustained(g, 8.0)

Measure the matrix kernels separately. The spin lasts a fixed number of cycles, so its wall time changes with the clock. Timing the whole graph would mix matrix-kernel time with the changing duration of the spin. I used nsys profile --cuda-graph-trace=node and extracted the individual matrix-kernel durations from the trace’s SQLite export.

Does it reproduce serving?

At four requests, I compared the CUDA kernel times in the proxy with their times in serving. I also measured the Gluon change in both environments:

Component timing: proxy versus serving traces

Without the spin, replay predicted -3.0% overall and -5.5% for gate/up. At one request, the proxy measured the CUDA lm_head at 351 us against 349 us in serving, and the Triton M = 1 change at -4.4% against -4.3%.

In these comparisons, the proxy estimated the overall change to within a couple of percentage points and recovered its direction. The individual shapes still differed, as the table shows.

Choose the configuration with the proxy

order = [None] + list(range(10)) + [None]        # production, 10 candidates, production again (drift check)
graphs = [build(i) for i in order]               # candidate i applies config i of each shape's list to every
                                                 # layer of that shape: six shapes are tuned in one graph
for g in graphs:
    g.replay()                                   # compile and warm everything before the capture starts
torch.cuda.synchronize()
torch.cuda.profiler.start()                      # nsys --capture-range=cudaProfilerApi
for g in graphs:
    sustained(g, 2.5)
    time.sleep(0.3)                              # idle gap: separates the segments in the trace
torch.cuda.profiler.stop()

def per_position(seg):
    t = collections.defaultdict(float); pend = 0.0; last = None; idx = 0
    for start, end, name in seg:
        d = (end - start) / 1e3
        if "_quantize" in name:
            pend += d; continue                     # charged to the next matmul
        if "_reduce_splits" in name:
            if last is not None: t[last] += d       # charged to the previous matmul
            continue
        if not re.search(MATMUL_KERNELS, name):
            continue                                # spin kernels and everything else
        last = idx % 257
        t[last] += d + pend; pend = 0.0; idx += 1
    steps = idx // 257
    return {p: v / steps for p, v in t.items()}, steps

Re-tuning with the proxy raised the serving gain at four requests from +1.8% to +3.3%, and at eight from -0.7% to +1.5%. It also helped explain the earlier configuration choice. Gate/up at four tokens had used a K split; the proxy rated that configuration at 1.02x production, consistent with the -0.6% serving result. The best proxy configuration used no K split and was also bit-identical.

Limits of the proxy

• The size of the gain still needed a serving check. For the ring kernel below, the proxy predicted -10% matmul time at four requests and -13% at eight; serving measured -12.5% and -20.2%. Its baseline for gate/up at eight tokens ran 12% faster than in serving. In these comparisons it ranked candidates correctly. I still used serving A/Bs for the final figures.
• It does not apply to prompt processing. Prefill is compute-bound at the cap, and the work between its matmuls (chunked attention and DeltaNet) is heavy too, so a low-power spin models the wrong thing. There, a kernel's in-serving speed depends on what its neighbours draw: switching only the wide prefill matmuls to Triton made them 3.6-4.7% slower in serving, while the unchanged CUDA out and down kernels got 10% faster. I therefore compared complete prefill chunks using serving A/Bs and NVML energy readings.
• It is calibrated for one card and one model. The spin lengths come from this model's non-matmul time on a 3090 Ti at 300 W. Another GPU or power limit needs a fresh serving trace to size them.

Use a ring buffer for longer rows

The direct kernel still lost on down (5120 x 17408). Its long rows benefited from loading whole 128-byte segments, as the CUDA ring kernel did. That kernel staged SG groups from each row through shared memory using cp.async.

Gluon exposed the operations needed to reproduce it: allocate_shared_memory, async_copy_global_to_shared, commit_group, wait_group and barrier. I kept the same register layouts and read the fragments from shared memory:

for it in range(ITERS):
    cp.wait_group(ST - 2)                 # stage `it` has landed for this thread ...
    ttgl.barrier()                        # ... and for every thread
    nit = it + ST - 1
    if nit < ITERS:                       # refill the slot that was consumed last iteration
        _issue(W, WS, Q, S, T, w_s, ws_s, q_s, xs_s, qs_s, nit % ST, first + nit * SG, ...)
    cp.commit_group()
    wst, qst = w_s.index(it % ST), q_s.index(it % ST)
    for gl in ttgl.static_range(SG):
        w = wst.slice(gl * 8, 8, dim=1).load(WORDS)                 # [BM, 8] words, LDS.64 per lane
        w = ttgl.permute(ttgl.reshape(w, [BM, 4, 2]), [0, 2, 1])    # register view [BM, h, t]
        wb = ttgl.convert_layout(_bytes(w), A_L)
        c = ttgl.zeros([BM, BT], ttgl.int32, layout=MMA)
        for s in ttgl.static_range(4):
            b = qst.slice(gl * 128 + s * 32, 32, dim=1).permute([1, 0]).load(B_L)
            c = mma_v2(_plane(wb, s), b, c)
        ws = ttgl.sum(ws_s.index(it % ST).slice(gl, 1, dim=1).load(MMA).to(ttgl.float32), 1)
        xs = ttgl.sum(xs_s.index(it % ST).slice(gl, 1, dim=0).load(MMA), 0)
        qs = ttgl.sum(qs_s.index(it % ST).slice(gl, 1, dim=0).load(MMA), 0)
        acc += (c - qs[None, :]).to(ttgl.float32) * ws[:, None] * xs[None, :]

With the same K split, the Gluon ring kernel was bit-identical to CUDA. In the proxy it improved every tested shape from three tokens upwards, including shapes beyond down: candidate/reference time was 0.73-0.82 for output, 0.84-0.91 for gate/up and 0.78-0.93 for down.

Serving gains reached +10.8% at four requests and +16.4% at eight. At eight concurrent requests, GSM8K scored 121 with the switch off and 120 with it on, with 127 of 128 answers identical.

Prompt processing: check configuration and launch costs

The plain-Triton prefill port used four masked-plane tl.dot operations per group and the -sum(q) correction. It was bit-exact, but took 4.5% more matrix-kernel time than CUDA in serving. The shared-memory staging described above still added work. Under the power cap, that instruction cost affected elapsed time.

I adapted the Gluon ring kernel to a 2-D prefill tile: 128 x 128 or 256 x 64 on an eight-warp grid. Keeping CUDA’s order for the final output calculation made the result bit-exact. In sustained loops it took 0.85-0.98 of the CUDA kernel’s time across the tested shapes and prompt lengths.

The first serving A/B showed +0%. Two checks explained why that result did not settle the comparison:

• The script had not set the intended switch. I had derived the A/B script from another one with a sed rule that renamed a log label GLUON= to PFGLUON=, which also turned GGUF_PQ2_MMQ_GLUON=$flag into GGUF_PQ2_MMQ_PFGLUON=$flag. Both servers ran CUDA, so the apparent difference was variation between identical configurations. I checked the exact switch string in the generated script and read it back from /proc/<pid>/environ in the running server.

• Short prompts were launch-bound. With the correct switch set, long prompts were 7-9% faster but a 128-token prompt was 12% slower. Prefill runs outside CUDA graphs, and every Gluon launch went through Triton's Python launcher, which binds the arguments and hashes the layout constants: 108 us of CPU per projection, against 65 us for the CUDA path. The fix caches a launch plan per weight and calls the compiled kernels' launchers directly (pq2_gluon_prefill, abridged):

  key = (weight.data_ptr(), n, k, cfg, x.dtype, m % 16 == 0, (m * G) % 16 == 0)   # Triton specializes on M % 16
  plan = _PLANS.get(key)
  if plan is None:                                                                # compiles every tile once
      plan = _PLANS[key] = _make_plan(x, weight, m, n, k, tuple(cfg), warm_cfgs)
  kern, qkern, codes, scales, bm, bt, consts = plan
  q = torch.empty((m, k), dtype=torch.int8, device=x.device)                      # per-call buffers
  s = torch.empty((m, G), dtype=torch.float32, device=x.device)
  t = torch.empty((m, G), dtype=torch.int32, device=x.device)
  y = torch.empty((m, n), dtype=x.dtype, device=x.device)
  qkern[((m * G + 7) // 8, 1, 1)](x, q, s, t, m * G, 8)                          # quantize
  kern[((n + bm - 1) // bm, (m + bt - 1) // bt, 1)](q, s, t, codes, scales, y, m, n, k, *consts)
  

  That is 36 us of CPU per call, half of what the CUDA path costs, bringing the 128-token prompt back to parity.

Matched serving runs measured with NVML's energy counter show that the Gluon prefill is +6.5 to +9.5% faster at 4096-token prompts and uses 7% less energy per token (139-151 mJ against 150-162 mJ).

Extend the kernels to sixteen token rows

The serving configuration had used eight sequences, so I had tuned up to eight tokens per matrix multiplication. With --max-num-seqs 16, vLLM pads 9 to 15 tokens to a 16-token CUDA graph. That shape still used the old Triton kernel and a floating-point output projection taking 52 us per call.

The baseline was slower at 12 requests than at 8, which made this path worth checking.

I extended both Gluon kernels to a 16-token tile by making the token dimensions of the accumulator, B operand and scale vectors parameters. With an unroll of 4, the direct kernel spilled temporary values into slower memory. The ring kernel was faster on every tested shape: 0.33-0.84 of production time, with bit-identical outputs on four of the six shapes.

Request performance: intermediate sixteen-token checkpoint

At this intermediate checkpoint, GSM8K at 16 concurrent requests scored 121 versus 121, with all 128 answers identical. The opening tables report the separate final comparison after all default changes.

Speed up a small unquantised projection

After replacing the PQ2 matrix kernels, cuBLAS still accounted for 10.7% of GPU time at four requests, compared with 2% at one and eight. The remaining cost came from in_proj_ba, a 96 x 5120 BF16 projection for ssm_alpha and ssm_beta, used by each of the 48 Gated DeltaNet layers.

At 2 to 4 tokens, cuBLAS selected eight thread blocks with no split-K. Each call took 30 us and moved data at about 30 GB/s; the bandwidth-only estimate was about 1 us. Swapping operands and pre-transposing the weight did not change its choice at four tokens (ba_gemm.py).

A split-K Triton kernel with a fixed-order reduction takes 4-5 us at every size up to 16:

@triton.jit
def _small_gemm_splitk(X, W, P, M, N: tl.constexpr, K: tl.constexpr,
                       BM: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr):
    pn, pk = tl.program_id(0), tl.program_id(1)      # one program per (N tile, K split)
    m = tl.arange(0, BM)
    n = pn * BN + tl.arange(0, BN)
    k = pk * BK + tl.arange(0, BK)
    x = tl.load(X + m[:, None] * K + k[None, :], mask=m[:, None] < M, other=0.0)
    w = tl.load(W + n[:, None] * K + k[None, :], mask=n[:, None] < N, other=0.0)
    acc = tl.dot(x, tl.trans(w), out_dtype=tl.float32)
    tl.store(P + pk * M * N + m[:, None] * N + n[None, :], acc, mask=(m[:, None] < M) & (n[None, :] < N))

# then: y = sum over splits in order 0, 1, 2, ... (deterministic), rounded once to bf16

My first version made no measurable difference in serving, and all 128 GSM8K answers were identical. I checked whether the new kernel was being called. Unquantised layers went directly to vLLM’s UnquantizedLinearMethod, bypassing the plugin’s GEMM.

I moved the routing into a subclass of that method and used a custom op so torch.compile would not specialise on token count. Serving improved by +9.5% at four requests, with no measurable change elsewhere.

Defaults and switches

The new paths belong to the default GGUF_PQ2_KERNELS=tuned set. The exact set enables only paths that are bit-identical to the older kernels: Gluon prefill and the small-batch cells marked exact. The legacy set disables all the new paths. Each path also has an override for paired before/after runs:

With these defaults on an RTX 3090 or 3090 Ti, no plugin CUDA kernel runs for this model at 1 to 16 tokens per matmul or in prompt processing; Nsight shows only cuBLAS for the remaining unquantised layers and PyTorch elementwise kernels. The CUDA kernels remain as fallbacks for other GPUs, other shapes and larger batches.

What did not work

• Plane DP4A for small batches in plain Triton: ALU-bound from three tokens, 2.5x slower at eight.
• The single-dot prefill variant (planes reshaped into one K = 128 dot) was no faster than the old _mmq; four masked K = 32 dots were.
• Unroll 4 at 16 tokens in the direct Gluon kernel spilled and ran 8-14x slower; unroll 2 is near production, and the ring kernel is better still.
• Moving loads earlier in the source without carrying them across loop iterations let ptxas reschedule them alongside the arithmetic.

Limits of the port

• The kernels are NVIDIA-specific. Gluon's mma_v2 and async_copy and the Triton DP4A inline PTX are Ampere-style NVIDIA instructions, so this is a port to Triton, not to ROCm. The CUDA extension is still built for the other GGUF formats.
• The tables are tuned for SM86 (RTX 3090 / 3090 Ti, 82-84 multiprocessors). The deployment card, the A4000, has 48 and a 140 W budget. The separate A4000 serving results above show little benefit from additional tuning; they do not calibrate the power proxy for that card.
• Above 16 tokens per matmul (more than 16 concurrent requests, or larger speculative verify batches) the old Triton kernel still runs; the Gluon kernels would need 32- and 64-token tiles.
• Gluon is experimental. The compiler limitations described above depend on the Triton version and may change in later releases.

What I took from this

The algorithms transferred more easily than the layouts. Plain Triton handled one-token generation, while Gluon gave me the control needed for the matrix instructions. I could express the work at a higher level, but choosing an efficient data layout still required understanding the GPU.

The measurement work mattered just as much. A replay at the wrong power state could select the wrong configuration. A script that failed to set its switch could compare the same configuration twice. I used the final serving comparison to check the combined changes under the workload they were meant to improve.

The A4000 runs complete the original serving test: the model runs on the 16 GB target card and reaches 349.4 aggregate output tokens per second at sixteen concurrent requests. Tuning made little difference there, so the development-card tuning gains should not be assumed to carry over.

Code and measurement sources

The CUDA/Triton comparison and kernel investigations refer to the pq2-m1-triton-planes-20261017 development record. The A4000 throughput figures were recorded separately. Raw logs and replay harnesses are not published with this website; selected code is included above. The public kernels are on the pq2-bundled-clean branch. Source links point to the latest published snapshot; the port was developed in these three commits:

• 6b13df3 triton and gluon: the M = 1 Triton kernel (pq2_single_gemv.py), the Gluon direct-fragment kernel (pq2_gluon_small.py) and the Triton plane prefill (_mmq_planes), with their tables and tests.
• 6072473 triton port of pq2: the Gluon ring and prefill kernels (pq2_gluon_ring.py, pq2_gluon_prefill.py), the 16-token tiles, the small_gemm.py GEMM with its linear method, the dispatch tables, and the switch to default-on in the kernel sets.
• 4fc90f9 update for triton: the README's switch documentation.