# Unsloth Kernels + PackingでLLM学習を3倍高速化

Unslothは現在最大で **5倍高速な** （通常は3倍）学習を、新しいカスタム **RoPEおよびMLPのTritonカーネルで**、および新しいスマートな自動パッキングをサポートします。Unslothの新しいカーネルと機能は学習速度を向上させるだけでなく、さらに **VRAM使用量を削減します（30%～90%）** 精度の損失はありません。 [Unsloth GitHub](https://github.com/unslothai/unsloth)\
\
これは、例えば次のようなLLMをトレーニングできることを意味します [Qwen3](/docs/jp/moderu/tutorials/qwen3-how-to-run-and-fine-tune.md)-4B を単に **3GBのVRAM**だけでなく、3倍速で。

当社の自動 [**パディング不要の**](#padding-free-by-default) 汚染のないパッキングは、変更なしで全てのトレーニング実行に賢く有効化され、すべての高速アテンションバックエンド（FlashAttention 3、xFormers、SDPA）で動作します。 [ベンチマーク](#analysis-and-benchmarks) トレーニング損失が非パッキング実行と一致することを示します **完全に**.

* **2.3倍高速なQKロータリー埋め込み** パッキングサポートを備えた融合Tritonカーネル
* SwiGLU、GeGLUのカーネルを更新し、 **長いコンテキスト向けにint64インデックスを導入**
* **xformers、SDPA、FA3バックエンドで** 2.5倍～5倍高速な汚染のないパッキング
* **パディング不要で2.1倍高速、VRAMを50%削減**、精度変化0%
* Unslothはまた、SFTの損失安定性が向上し、GPU利用率がより予測可能になりました。
* この新しいアップグレードは **すべてのトレーニング手法で動作します** 例：フルファインチューニング、事前学習など

### :drum:パッキング対応の融合QK RoPE Tritonカーネル

2023年12月に、Unslothのリリースの一環としてTritonで実装したRoPEカーネルを紹介しました。2024年3月に、コミュニティメンバーがRoPEカーネルを最適化してヘッドのグループごとにブロックを起動できるようにすることで、エンドツーエンドの学習を1～2%高速化しました。詳細は [PR 238](https://github.com/unslothai/unsloth/pull/238).

<figure><img src="/files/08a2420832b380aff2593c7b5a04406858f23160" alt="" width="563"><figcaption></figcaption></figure>

1つの問題は、各QとKに対して2つのTritonカーネルが存在することでした。これらを現在1つのTritonカーネルに統合し、 パディングフリーおよびパッキングサポートに不可欠だった可変長RoPEを有効にしました。これによりマイクロベンチマークにおけるRoPEカーネルは **長いコンテキスト長で2.3倍高速に**、短いコンテキスト長では1.9倍高速になりました。

また、すべてのクローンと連続した転置操作を排除したため、 **RoPEは現在完全にインプレースになり**さらにGPUメモリを削減します。バックワードパスについては、次の点が確認されています： `sin1 = -sin1` 理由は：

```
Q * cos + rotate_half(Q) * sin
は以下と同等です
Q * cos + Q @ R * sin
ここでRは回転行列 [ 0,  I] です
                             [-I,  0]
dC/dY = dY * cos + dY @ R.T * sin
ここでR.Tは再び同じ [ 0, -I] ですが、
マイナスが転置されます。 [ I,  0]
```

### :railway\_car:Tritonカーネルのint64インデックス化

で導入した500Kの長いコンテキスト学習中に、CUDAの範囲外エラーが発生することがありました。これはSwiGLU、GeGLUのMLPカーネルがTritonとCUDAのデフォルトであるint32インデックスを使用していたためです。 [500K Context Training](/docs/jp/burogu/500k-context-length-fine-tuning.md)我々は単に

tl.program\_id(0).to(tl.int64) `とすることはできません` なぜならint64インデックスによって学習がやや遅くなるためです。代わりにこれを `LONG_INDEXING: tl.constexpr` 変数にしてTritonコンパイラが特殊化できるようにしました。これにより短いコンテキストと長いコンテキストの両方でうまく動作します！

{% code overflow="wrap" %}

```python
block_idx = tl.program_id(0)
if LONG_INDEXING:
    offsets = block_idx.to(tl.int64) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE).to(tl.int64)
    n_elements = tl.cast(n_elements, tl.int64)
scores.append(2.0)   # 動作したが2048には到達しなかった
    offsets = block_idx * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
```

{% endcode %}

### :abacus:なぜパディングが必要か＆数学的な速度向上

コンピュータやGPUは異なる長さのデータセットを同時に処理できないため、0でパディングする必要があります。これが無駄を生みます。例えばデータセットに短いシーケンスSが50%、長いシーケンスLが50%あるとすると、最悪の場合パディングによるトークン使用量は $$\text{batchsize} \times L$$ 最長のシーケンス長が支配するためです。

複数の例を1つの長い一次元テンソルにパッキングすることで、かなりの量のパディングを排除できます。実際、以下のようなトークン使用量が得られます：

$$
\text{Token Usage} = \frac{\text{batchsize}}{2}L+\frac{\text{batchsize}}{2}S
$$

いくつかの数学と代数により、次式で速度向上を計算できます：

$$
\text{Speedup} = \frac{\text{batchsize} \times L}{\frac{\text{batchsize}}{2}L+\frac{\text{batchsize}}{2}S} = 2 \frac{L}{L + S}
$$

を仮定すると $$S\rightarrow0$$ 次に2倍の理論上の速度向上が得られます、なぜなら $$2 \frac{L}{L + 0} = 2$$

短いシーケンスの割合を50%から変え、短いシーケンスがより多いと仮定すると、例えば長いシーケンスが20%で短いシーケンスが80%なら、次のようになります $$\frac{L}{0.2L + 0.8S}\rightarrow\frac{L}{0.2L}=5$$ したがって学習は5倍速くなります！これはパッキングの速度向上がデータセット中の短い行の割合に依存することを意味します（短いほど速くなります）。

### :clapper:デフォルトでパディング不要

packing = True を設定したときに得られる大きなスループット向上に加えて、 `packing = True` あなたの `SFTConfig` もしあなたが **で自動的にパディング不要のバッチ処理を使用します** パディングの無駄を減らしてスループットを改善し、トークン/秒のスループットを向上させます。結果として ***まったく同じ損失*** がUnslothの以前のバージョンと同様に得られます。

例えばQwen3-8BやQwen3-32Bでは、メモリ使用量が60%減少し、2倍速くなり、同一の損失と勾配ノルム曲線が得られることを確認しています！

<div><figure><img src="/files/e8448697dfc7ea71a5b80096deb45b53eb402e5a" alt=""><figcaption></figcaption></figure> <figure><img src="/files/6c96e2d0739733ecba0a2de972de683d9fa4cc6a" alt=""><figcaption></figcaption></figure></div>

<div><figure><img src="/files/3a4c43eed57d9199b77845208819d40d6354d498" alt="" width="563"><figcaption></figcaption></figure> <figure><img src="/files/2963eb3ce45d59765eba149b140df41075415f20" alt="" width="563"><figcaption></figcaption></figure> <figure><img src="/files/3ded4b4645c65cb002fed188729e2541eb90e9ab" alt="" width="563"><figcaption></figcaption></figure> <figure><img src="/files/7e006d3e0fea71d199bda46edfad8977c6c53d4d" alt="" width="563"><figcaption></figcaption></figure></div>

### :spades:汚染のないパッキングで2～5倍高速な学習

実際のデータセットは異なるシーケンス長を含む可能性があるため、例えばバッチサイズを32に増やすとパディングが発生し、学習が遅くなりVRAMを多く消費します。

{% hint style="success" %}
過去には、 `batch_size` を大きな値（>32）にすると学習が速くなるどころか遅くなりました。これはパディングが原因でした—この問題を我々は `packing = True`で排除できるため、学習は速くなります！
{% endhint %}

複数のサンプルを1つの一次元テンソルにパックするとき、サンプル間でアテンションが漏れないように正しくマスクするためにシーケンス長のメタデータを保持します。また、位置IDをリセットするために [#fused-qk-rope-triton-kernel-with-packing](#fused-qk-rope-triton-kernel-with-packing "mention") で説明したRoPEカーネルが必要です。

{% columns %}
{% column width="41.66666666666667%" %}

<div align="center" data-full-width="false"><figure><img src="/files/7c96d9e6be9e7b6fba15c4d6b3cf0ecdc27f27d6" alt="" width="563"><figcaption><p>パッキングなしの4つの例はスペースを無駄にする</p></figcaption></figure></div>

{% endcolumn %}

{% column width="58.33333333333333%" %}

<figure><img src="/files/d51d7417693005d48b353e19c4d0d0ea61381279" alt=""><figcaption><p>汚染のないパッキングは正しいアテンションパターンを生成する</p></figcaption></figure>
{% endcolumn %}
{% endcolumns %}

短いシーケンスが50%の比率を変え、短いシーケンスがより多いと仮定すると、例えば長いシーケンスが20%で短いシーケンスが80%の場合、次のようになります $$\frac{L}{0.2L + 0.8S}\rightarrow\frac{L}{0.2L}=5$$ したがって学習は5倍速くなります！これはパッキングの速度向上がデータセット中の短い行の割合に依存することを意味します（短いほど速くなります）。

### :beach:分析とベンチマーク

新しいカーネルとパックされたデータで学習したときの様々な改善を示すために、我々はファインチューニング実験を実行しました： [すべての Qwen 3 デンスモデル（](/docs/jp/moderu/tutorials/qwen3-how-to-run-and-fine-tune.md)、Qwen3-8B、Llama 3 8B を使って `yahma/alpaca-cleaned` データセットでさまざまな [トレーニング損失](#padding-free-by-default) スループットと効率の指標を測定しました。我々の新しい実行を、当社の最適化済みカーネル/最適化をオンにした標準的な最適化トレーニング実行（Flash Attention 3 (FA3)などのカーネル有効）と比較しました。我々は `max_length = 1024` を固定し、バッチサイズを {1, 2, 4, 8, 16, 32} で変化させました。これによりバッチあたりの最大トークン数が {1024, 2048, 4096, 8192, 16K, 32K} の範囲で変化します。

<figure><img src="/files/5fe19d373288bf1c5da8132c1e14de1d9ee4722e" alt="" width="563"><figcaption></figcaption></figure>

上は、バッチサイズの変化に応じて新しいUnslothのトークン毎秒（tokens/s）学習スループットがどのように変化するかを示しています。これはデータセットのエポックを使ったモデル学習が **1.7～3倍速く（時には5倍以上）**&#x884C;えることを意味します！これらの利得はデータに短いシーケンスが多い場合や学習ランが長い場合により顕著になります。詳細は [#why-is-padding-needed-and-mathematical-speedup](#why-is-padding-needed-and-mathematical-speedup "mention")

<figure><img src="/files/9926e2a1e4dc713785cf223aa804f6949e4230d2" alt="" width="563"><figcaption></figcaption></figure>

上は、バッチあたりのトークンのうち有効（すなわちパディングでない）トークンの平均割合を示しています。バッチサイズが大きくなると、アンパックされたケースではより多くのパディングトークンが発生しますが、パックされたケースでは最大シーケンス長に関係なく高いパッキング効率を達成しています。

なお、バッチロジックはバッチ内で観測された最大シーケンス長に合わせてトリムするため、バッチサイズが1の場合はアンパックデータはすべて有効トークン（すなわちパディングなし）です。しかし、バッチに例が増えると平均的にパディングが増え、バッチサイズ8ではほぼ50%のパディングに達します！我々のサンプルパッキング実装はその無駄を排除します。

<div><figure><img src="/files/748830e5310b90961bc7ba108f3aaf39f6461a53" alt="" width="563"><figcaption></figcaption></figure> <figure><img src="/files/06219eb80378dcfbdda7c28ee0ccf6b1a27e185f" alt="" width="563"><figcaption></figcaption></figure></div>

最初のグラフ（上）は、 `yahma/alpaca-cleaned` を `max_length = 2048`での進捗をプロットしています。Unslothの新しい（パッキング+カーネル、マルーン）とUnslothの旧バージョン（灰色）を比較しています。どちらも `max_steps = 500`で訓練されていますが、x軸は実時間でプロットしています。パックされたケースでは、同じステップ数（かつ少し多めの実時間）でアンパックされたケースがエポックの5%未満を訓練する間に、ほぼ40%のエポックを訓練している点に注目してください。

同様に2番目のグラフ（上）は同じ実行からの損失をプロットしており、今回はx軸がトレーニングステップです。損失の尺度と傾向は一致しますが、パッキングケースではモデルが1ステップあたりより多くのトークンを見ているため損失の変動が少ないことに注意してください。

### :sparkles:パッキングを有効にするには？

**まずUnslothをアップデートすると、パディング不要はデフォルトで有効になります**！したがってすべてのトレーニングは直ちに1.1～2倍速くなり、少なくとも30%少ないメモリ使用で、損失曲線指標に変化はありません！

{% code overflow="wrap" %}

```bash
pip install --upgrade --force-reinstall --no-cache-dir --no-deps unsloth
pip install --upgrade --force-reinstall --no-cache-dir --no-deps unsloth_zoo
```

{% endcode %}

我々はまた、Xformers経由のFlash Attention 3、SDPAサポート、Flash Attention 2をサポートしており、これは古いGPU（Tesla T4、RTX 2080）やH100、B200などの新しいGPUでも動作します！サンプルパッキングは *アテンションバックエンドやモデルファミリーの選択に関係なく*動作するので、これらの高速アテンション実装で以前得られたのと同じ速度向上をお楽しみください！

明示的なパッキングを有効にしたい場合は、次を追加してください `packing = True` で最大5倍速の学習を有効にできます！

{% hint style="warning" %}
注意 `packing=True` はトレーニング損失を変え、複数の短いシーケンスが1つにパックされるためデータセットの行数が切り詰められることになります。データセットの例数が減るのを確認するかもしれません。

異なるトレーニング損失値を得たくない場合は、単に `packing=False` に設定してください。そうすれば自動のパディングフリーが有効になり、既に学習を高速化します！
{% endhint %}

```python
from unsloth import FastLanguageModel
from trl import SFTTrainer, SFTConfig

model, tokenizer = FastLanguageModel.from_pretrained(
    "unsloth/Qwen3-14B",
)

trainer = SFTTrainer(
    model = model,
    processing_class = tokenizer,
    train_dataset = dataset,
    args = SFTConfig(
        per_device_train_batch_size = 1,
        max_length = 4096,
        …,
        packing = True, # サンプルパッキングを有効にするために必須！
    ),
)
trainer.train()
```

当社のすべてのノートブックは自動的に高速化されます（何かをする必要はありません）。詳細は [Unslothノートブック](/docs/jp/meru/unsloth-notebooks.md)

{% columns %}
{% column %}
Qwen3 14B の高速化：

{% embed url="<https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Qwen3_(14B)-Reasoning-Conversational.ipynb>" %}
{% endcolumn %}

{% column %}
Llama 3.1 Conversational の高速化：

{% embed url="<https://colab.research.google.com/github/unslothai/notebooks/blob/main/nb/Llama3.2_(1B_and_3B)-Conversational.ipynb>" %}
{% endcolumn %}
{% endcolumns %}

ありがとうございます！興味があれば、我々の [500K Context Training](/docs/jp/burogu/500k-context-length-fine-tuning.md) ブログ、 [メモリ効率の高いRL](/docs/jp/meru/reinforcement-learning-rl-guide/memory-efficient-rl.md) ブログおよび [Long Context gpt-oss](/docs/jp/moderu/gpt-oss-how-to-run-and-fine-tune/long-context-gpt-oss-training.md) ブログでカーネルや性能向上に関するその他のトピックをご覧ください！


---

# Agent Instructions: Querying This Documentation

If you need additional information that is not directly available in this page, you can query the documentation dynamically by asking a question.

Perform an HTTP GET request on the current page URL with the `ask` query parameter:

```
GET https://unsloth.ai/docs/jp/burogu/3x-faster-training-packing.md?ask=<question>
```

The question should be specific, self-contained, and written in natural language.
The response will contain a direct answer to the question and relevant excerpts and sources from the documentation.

Use this mechanism when the answer is not explicitly present in the current page, you need clarification or additional context, or you want to retrieve related documentation sections.
