NVIDIA CUDA 13.3 で暗号化を高速化

15 年以上にわたり、x86 CPU はキャリーレス乗算専用のハードウェア命令を搭載してきました。認証暗号化、誤り訂正符号、そして現代のゼロ知識証明の根底にある、小さくても堅牢なこのプリミティブです。

これまで NVIDIA GPU にはネイティブサポートがありませんでした。NVIDIA CUDA 13.3 でその欠落が埋まります。新設された PTX 命令「clmad」は、Ampere 以降のすべての NVIDIA GPU（SM 80 以上）で利用可能です。本稿では、キャリーレス乗算がようやくハードウェアアクセラレーションされた GPU プリミティブとなったことで何が可能になるかを示します。

これに依存する 2 つの暗号化ワークロードをベンチマークしました。1 つは AES-GCM（TLS や VPN、データセンター暗号化の基盤となる AEAD 暗号）内部の整合性ハッシュ「GHASH」です。もう 1 つは、高度なゼロ知識証明システムの根幹を成す内側ループである「sum-check プロトコル」です。NVIDIA B200 では GHASH のスループットが約 6.3 TB/s に達し、これは DRAM の読み取り帯域幅に匹敵します。また、従来のビットスライス方式の最先端技術と比較して最大 18.8 倍高速化されました。sum-check プロトコルも同様に、既存の最先端技術よりも 4〜13 倍高速化されています。

AES-GCM 以外でも、なぜこれが重要なのか。それは、 Carryless Multiplication（キャリーレス乗算）が、意外にも幅広い暗号処理や符号理論のワークロードにおいて共通する中核カーネルとなっているからです。

ストレージシステムや通信基地局のベースバンド処理で使われる CRC やリード・ソロモン符号、フラッシュメモリ向けの BCH 符号、量子スタビライザー符号、いくつかのポスト量子暗号方式、そして Binius に代表される現代のゼロ知識証明システムの基盤となる二進体上の算術演算。これらすべてのワークロードのコスト構造は、すでに導入されている Ampere アーキテクチャ以降の GPU 上でハードウェアアクセラレーションによって劇的に変化します。

この記事は、GPU パイプラインに暗号機能を統合する CUDA 開発者や、二進体プロトコルを基盤とするセキュリティ・プライバシー研究者向けに書かれています。CUDA 13.3 をダウンロードして、以下の PTX の例を実際のカーネルにそのまま組み込んで試してみてください。

2 進拡大体

最も小さな有限体 GF(2) は単なる 1 ビットの世界です。ここでは加算が XOR（排他的論理和）、乗算が AND（論理積）になります。暗号利用の文脈では、ビットは通常 *2 進拡大体* GF(2^m) に組み込まれます。ここで m ビットは、GF(2) 上の多項式の係数として表現され、演算は既約多項式を法とした剰余計算となります。

この体における多項式の加算は依然として単純な XOR（係数の要素ごとの和）ですが、乗算には各入力のビット間での長い乗算処理が必要になります。

具体的には、入力値の拡張体積乗算における「th」番目の結果係数は、個々のビット積の XOR（排他的論理和）を計算した後に、その体の既約多項式による簡約処理を行うことで得られます。

例えば、2 ビット拡張体積では、既約多項式は一意に定まり、特定の値となります。

この計算を AND 演算と XOR 演算の組み合わせでシミュレートすることは可能ですが、個々のビット値を取り出すために繰り返されるビットシフトやマスク処理が大きなオーバーヘッドとなり、性能を著しく低下させます。このような体積乗算を直接ハードウェアレベルでサポートする機能は x86 アーキテクチャでは以前から利用可能でしたが、主に AES-GCM の実装に特化していました。しかしこれまで、NVIDIA GPU ではキャリーレス乗算のネイティブサポートが提供されておらず、開発者はパフォーマンス向上のためにビットスライス回路などの代替手法を採用する必要がありました。

キャリーレス乗算のためのハードウェアサポート

CUDA 13.3 では、キャリーレス乗算・加算積算（multiply-accumulate）を行う新しい PTX 命令 clmad が追加されました。

x86 アーキテクチャの PCLMULQDQ に相当する clmad 命令は、2 つの 64 ビット入力に対して演算を行ない、128 ビットの結果を返します。.hi と .lo のバリアントは、それぞれ 128 ビット出力の上側と下側の半分を計算し、そこに 64 ビットの累算値（アキュムレータ）を加算します。実装例は以下の通りです。

__device__ inline uint128_t clmad_mul_128(uint64_t a, uint64_t b, uint128_t acc) {

uint64_t acc_lo = (uint64_t)acc;

uint64_t acc_hi = (uint64_t)(acc >> 64);

uint64_t lo, hi;

asm("clmad.lo.u64 %0, %1, %2, %3;" : "=l"(lo) : "l"(a), "l"(b), "l"(acc_lo));

asm("clmad.hi.u64 %0, %1, %2, %3;" : "=l"(hi) : "l"(a), "l"(b), "l"(acc_hi));

return ((uint128_t)hi << 64) | lo;

}

この命令は、Ampere アーキテクチャ（sm_80 以降）を搭載した NVIDIA GPU でハードウェアレベルで加速されており、現代の NVIDIA GPU における高速なバイナリ拡張体積乗算を可能にしています。

本稿の後半では、clmad がどのようにして 2 つの暗号化ユースケースである GHASH とサムチェックプロトコルの処理を加速するかを見ていきます。これらはどちらもより大きな体積上で動作します。この体積内では、要素同士の乗算に カラツバアルゴリズム を適用することで、clmad 命令を 6 回使用して計算を実行できます。

__device__ inline uint256_t clmad_mul_256(uint128_t a, uint128_t b) {

uint64_t a0 = (a & 0xFFFFFFFFFFFFFFFFull), a1 = (a >> 64);

uint64_t b0 = (b & 0xFFFFFFFFFFFFFFFFull), b1 = (b >> 64);

uint128_t z0 = clmad_mul_128(a0, b0, 0);

uint128_t z1 = clmad_mul_128(a0 ^ a1, b0 ^ b1, 0);

uint128_t z2 = clmad_mul_128(a1, b1, 0);

z1 ^= z0 ^ z2;

z0 ^= z1 << 64;

z2 ^= z1 >> 64;

return {z2, z0};

}

256 ビットの結果を再び 128 ビットのフィールドに戻すには、追加の CLMAD 呼び出しを用いた乗算ベースの Barrett reduction を計算するか、多項式の筆算除法を模倣するシフトと XOR のループを直接実行します。

GHASH の高速化

GHASH は clmad（キャリーレス乗算）の適用に最適な候補です。これは、既知の不可約多項式で割った後の 2 進数乗算に強く依存しているためです。

AES-GCM の文脈全体において、GHASH は入力データ全体（暗号化テキスト、追加認証データ、および最終的な長さブロック）に対する核心的な認証ハッシュを計算します。GHASH は入力を 128 ビットブロックに分割し、各ブロックに対して、それをランニングアキュムレータに XOR して、AES でゼロブロックを暗号化して導出されたハッシュ鍵（H）と乗算します。

この乗算演算は、clmad に続いてモジュラ減算を行うことで実現できます。すべてのブロックの処理が完了すると、128 ビットの認証ハッシュが得られます。その後、このハッシュに初期カウンタブロックの AES 暗号化結果を XOR することで、最終的な認証タグが生成されます。

image*図 1. 入力ペイロード長最大 2GB の場合の GHASH スループット測定結果。ALU 演算を用いたビットスライス回路と、B200 および NVIDIA GeForce RTX 5090 で実装した CLMAD ベース版を比較*。

NVIDIA GeForce RTX 5090 では、CLMAD を用いた GHASH のピークスループットが約 1,300 GB/s に達し、ビットスライス方式の実装と比較して 2 倍の性能を発揮しました。

一方、B200 では約 6,335 GB/s というピーク値を記録（DRAM 読み出し帯域幅にほぼ匹敵）し、ビットスライス版よりも最大で 18.8 倍もの高速化を実現しています。ビットスライスのベースラインでは B200 は SM 数やクロック周波数が低いこともあり RTX 5090 よりも遅いですが、CLMAD のハードウェアアクセラレーションにより GHASH の処理が劇的に高速化されました。

ゼロ知識証明：Sum-check の加速

ゼロ知識（ZK）証明プロトコル 証明方式 では、二進拡張体が利用されています。これは、ある計算を実行した当事者が、検証者に再計算を求めずに結果の正当性を暗号学的に証明することを可能にする技術です。ハッシュ処理や通常の整数演算のエミュレーションなど、ビット単位の操作として自然に評価されるアルゴリズムの結果を証明する際、二進体は特に有効です。当然ながら、ゼロ知識プロトコルでは多数のキャリーレス乗算が必要となります。

多くの証明システム（many proving systems）の基盤となっている基本的な要素が、sum-check プロトコル です。このプロトコルは、ある当事者がブール入力に対する n 変数多項式の和を証明する仕組みを提供します。これにより、検証者の作業量は n に比例するだけで済み、n+1 点目のランダムな点における多項式の値を評価すれば十分となります。

Sum-check プロトコルは、変数の数だけラウンド（反復）を行います。ここでは、多項式の合成（P = P_1 ∘ ... ∘ P_k）を処理するシンプルな実装に焦点を当てます。各多項式 P_i は、F 上の n 点における評価値のリストとして表現されます。

ラウンド i では、証明者が以下の手順を実行します：

• 各評価行において多項式間の積を合成し、評価値の積和（claimed product sum）を計算する
• n 点で多項式を補間してその積和を求め、このラウンド用の単変数多項式を生成する
• 検証者が選んだランダムな点で多項式を補間し、評価値の数が半分になった新しい多項式のセットを用意して次のラウンドへ進む

各補間と積計算には、多項式の評価値間の非常に多くの二元拡張体乗算が必要となり、入力サイズや合成の規模に依存してスケーリングします。しかし、これらの処理はほぼ完全に並列化可能です！

image*図 2. ポリノミアルサイズと合成サイズの範囲における Sum-check プロトコルのレイテンシ改善。ALU 演算を用いたビットスライス回路と、CLMAD ベースのバージョンを比較*。

ここでは、線形時間・線形空間で動作する Sum-check プロトコルにおいて、clmad を用いたフィールド演算を採用した手法と、Irreducible が同様のフィールド演算向けに開発したビットスライス回路を用いた手法を比較しました。バイナリフィールドに対するハードウェアサポートを活用することで、RTX 5090 では性能が 3〜4 倍、B200 GPU では最大 13 倍向上します。また、ポリノミアルサイズや合成サイズが大きくなるにつれて、その性能向上効果はわずかに高まる傾向にあります。

結論

CUDA 13.3 は、NVIDIA GPU に欠けていた重要な機能である「ハードウェア加速されたキャリーレス乗算」を実装しました。B200 において、GHASH と和チェックプロトコルの 2 つのワークロードをベンチマークした結果、それぞれ最大で 18.8 倍、13 倍の高速化が実現されました。最も劇的な性能向上が見られるのは、現代の暗号技術が進むべき方向性、すなわち大規模入力に対応する認証付き暗号や、二進体におけるゼロ知識証明の分野です。

活用方法：

• CUDA 13.3 をダウンロードしてください
• 上記のインライン PTX サンプルを、ご自身のカーネル（Ampere アーキテクチャ以降が必要）で試してみてください
• 暗号化ワークロードを対象にしている場合は、cuPQC SDK をご覧ください
• テレコム向けアプリケーションを開発中の方は、NVIDIA Aerial のドキュメントをご参照ください