3x Faster LLM Training with Unsloth Kernels + Packing

TL;DR

Unsloth reports up to 5× (typically 3×) faster fine-tuning and 30–90% lower VRAM with no accuracy loss by combining (1) a fused QK Rotary Embedding Triton kernel that supports variable-length packing, (2) updated SwiGLU/GeGLU MLP kernels with int64 indexing for long-context safety, and (3) auto-enabled “uncontaminated” sequence packing that eliminates padding waste while preserving exact attention masking between packed samples. Headline: Qwen3-4B trainable on 3.9 GB VRAM; Qwen3-8B/32B see 60% lower memory and 2× speedup with identical loss/grad-norm curves. The auto padding-free mode is on by default after update across FA3/xFormers/SDPA backends and old (T4/2080) + new (H100/B200) GPUs; explicit packing=True unlocks the 5× ceiling.

Key claims

The fused QK-RoPE Triton kernel merges previously separate Q-side and K-side kernels into one and adds variable-length / reset-position-id support; in microbenchmarks it runs 2.3× faster at long context, 1.9× at short context, and is fully in-place (eliminating clones and contiguous transposes) [§Fused QK RoPE Triton Kernel with packing].
The SwiGLU/GeGLU kernels are upgraded from int32 to compile-time-specialized LONG_INDEXING: tl.constexpr int64 indexing to avoid the CUDA out-of-bounds errors seen during 500K-context training, without slowing short-context runs [§Int64 Indexing for Triton Kernels].
Uncontaminated packing concatenates multiple samples into a single 1D tensor while carrying per-sample length metadata so attention masking blocks cross-sample leakage; packing+fused-RoPE delivers 2.5×–5× faster training depending on dataset sequence-length distribution [§Uncontaminated Packing 2-5x faster training].
Theoretical packing speedup is (short_frac · S + long_frac · L) / L so very short-tailed datasets get larger gains; a 20%-long / 80%-short mix yields ~5× [§Why is padding needed & mathematical speedup].
Padding-free batching is auto-enabled after the update for all training runs (full fine-tuning, pretraining, SFT) with no config changes, across FA3, xFormers, SDPA, FA2 backends [§Padding-Free by Default; §How to enable packing?].
On Qwen3-8B and Qwen3-32B with padding-free default, memory drops 60% and step time is 2× faster with exactly identical loss and grad-norm curves vs. the prior Unsloth version [§Padding-Free by Default].
Benchmark methodology: Qwen3-32B / Qwen3-8B / Llama-3-8B SFT on yahma/alpaca-cleaned with max_length=1024, batch sizes ∈ {1,2,4,8,16,32}; reported throughput is 1.7–3× higher and reaches 5× on batches with high length variance, while packed-batch token-validity stays near 100% vs. unpacked dropping to ~50% padding at batch size 8 [§Analysis and Benchmarks].
packing=True changes the visible loss curve and shrinks the apparent dataset row count because multiple short rows are concatenated into one; the unchanged-loss path is the default padding-free mode (packing=False) [§How to enable packing?].
All Unsloth notebooks become automatically faster on update; sample packing is independent of attention-backend choice and works on old (Tesla T4, RTX 2080) and new (H100, B200) GPUs [§How to enable packing?].

Method

Three composable optimizations layered onto the existing Unsloth Triton-kernel stack:

Fused variable-length QK RoPE. Previous Q-side and K-side RoPE kernels are merged; the kernel accepts per-sample length metadata so RoPE position ids reset at each packed-sample boundary, which is the missing piece that lets packed sequences share an attention kernel without leaking positional information across boundaries. In-place computation eliminates intermediate tensor allocations; the backward exploits sin1 = -sin1 symmetry.
Long-indexing-aware MLP kernels. SwiGLU and GeGLU kernels use a Triton tl.constexpr flag to switch between int32 and int64 indexing at compile time, so 500K-context training stops crashing with CUDA out-of-bounds errors but short-context runs keep int32 speed.
Uncontaminated packing. Multiple short samples are packed into one 1D tensor; sequence-length metadata is preserved and combined with the variable-length RoPE kernel + attention mask structure so the packed tensor is mathematically equivalent to a batch of separate forward passes, with no cross-sample attention.

Padding-free batching (a weaker form of the same idea — trim batches to the longest actual sample in the batch) is auto-enabled even with packing=False, giving the 1.1–2× “free” speedup; explicit packing=True unlocks the 5× ceiling at the cost of a changed visible loss curve (since per-step token counts change).

Results

Headline numbers, all from Unsloth’s benchmarks unless noted:

Qwen3-4B fine-tunable on 3.9 GB VRAM, 3× faster [§intro].
Qwen3-8B / Qwen3-32B with padding-free default: 60% memory reduction, 2× faster, exactly identical loss + grad-norm curves [§Padding-Free by Default].
Per-kernel speedups: QK-RoPE 2.3× (long ctx) / 1.9× (short ctx); SwiGLU/GeGLU updated for 500K-context safety; uncontaminated packing 2.5–5×; padding-free 2.1× with 50% less VRAM at 0% accuracy change [§intro list].
Throughput across batch sizes: 1.7–3× tokens/s on Qwen3-32B / 8B / Llama-3-8B with max_length=1024, climbing to 5× on length-variant batches [§Analysis and Benchmarks].
Packing efficiency: packed runs maintain near-100% valid-token ratio across all batch sizes; unpacked runs drop to ~50% valid tokens at batch size 8 due to padding [§Analysis and Benchmarks].
Loss-curve fidelity: at fixed step count, packed runs cover ~40% of an epoch vs. <5% for unpacked on yahma/alpaca-cleaned max_length=2048; loss-vs-step curves match in scale and trend, with the packed curve slightly less noisy because each step sees more tokens [§Analysis and Benchmarks].
Compatibility: works on T4 / RTX 2080 → H100 / B200, with FA3 (via xFormers), SDPA, FA2 backends, regardless of model family [§How to enable packing?].

Why it’s interesting

The fused variable-length QK-RoPE + uncontaminated packing fits squarely in the IO-Aware Kernel Design template — find a per-batch coordination cost that ought to be O(1) (here: rebuilding RoPE position state per sample inside the kernel, plus per-batch padding that scales with the longest sample) and rewrite it so the kernel absorbs the variable-length metadata directly. It is the kernel-side complement to the How to Make LLM Training Faster with Unsloth and NVIDIA systems-side gains, which cached packed-sequence attention metadata across layers; this post caches position-id reset inside the RoPE kernel itself. The two posts together cover both ends of the same packed-training stack — the attention kernel boundary and the per-layer attention-metadata rebuild — and Luma already runs the layer-cache piece, so the kernel piece here is the remaining systems lift on the packed-SFT path.

The “5× when datasets have many short rows” math is the relevant guidance for Parameter-Efficient Finetuning workflows on conversational / alpaca-style data where most rows are 100–1000 tokens but a few are 8k+; on Luma’s typical SFT mixes the realistic ceiling is closer to 3×, but the auto-enabled padding-free default still gives the 1.1–2× free.