はじめに
近年、準同型暗号（FHE）などの暗号の実装において、巨大な行列演算と多倍長整数演算の高速化は避けて通れない課題となっています。
これらをCPUで処理するには限界があり、GPUによる並列化が定石ですが、GPU開発の主流は依然としてC++（CUDA）です。一方で、（個人的に）開発が難しいC++ではなくメモリ安全性やエコシステムの利便性があり（個人的に）扱いやすいRustを使いたい。
欲を言えばRustでGPUプログラミングを行い、かつC++並みのパフォーマンスを出したい。
そこで今回は、RustにおけるGPU開発の実用性を確認した後、下記のライブラリの性能調査をおこないました。
- GMP: CPUにおける多倍長整数のデファクトスタンダード
- Rust-CUDA: Rustで直接CUDAカーネルを書くアプローチ
- CGBN: NVIDIA公式の研究者が関わるGPU用多倍長整数ライブラリ
本記事では、これらを用いた実装比較と、その結果を紹介します。
実験に用いた基盤ライブラリ
検証に入る前に、今回実験に使用した3つの技術要素について簡単に解説します。
1. GMP (GNU Multiple Precision Arithmetic Library)
CPU上で任意精度の算術演算を行うためのオープンソースライブラリです。
数十年以上の歴史があり、CPUにおける多倍長演算の「基準（ベースライン）」として最も信頼性が高いものです。GPU上では使用できず、Rustで使用するにはC/C++からRustへのラッパーを実装する必要があります。
2. Rust-CUDA (Rust for CUDA)
RustのコードをGPUが理解できる中間言語（NVPTX）にコンパイルするプロジェクト、およびそのエコシステムです。
従来はC++で書いていた __global__ 関数（カーネル）をRustで記述でき、#[kernel] や #[address_space(shared)] といった属性マクロを使ってGPUリソースを制御します。型安全性というRustの恩恵を受けながらGPUコードが書けるのが最大の魅力です。しかしながら、多倍長演算をするには新たにカーネル関数で多倍長演算の実装をする必要があります。
3. CGBN (CUDA Generic Big Number)
NVIDIAの研究者らが開発に関わっている、CUDA向けの多倍長整数ライブラリ（C++ヘッダオンリー）です。
GPUの各スレッドを協調させて一つの巨大な整数を扱うなど、高度な最適化が施されています。GMPと同様に、Rustで使用するにはC/C++からRustへのラッパーを実装する必要があります。
4. C++-CUDA（CUDA C++）
NVIDIA GPU上で高速な並列計算を行うための、標準言語／開発環境です。
__global__ カーネル関数，__device__ 関数，__shared__ メモリ，warp/shuffle命令など，GPUハードウェアの性能を最も直接的に引き出せます。GMPのようなCPU向け任意精度ライブラリはそのまま使えませんが，多倍長演算は「自前実装」か「CGBNなどのCUDA向けライブラリ」をそのまま利用する形で実現できます。Rust-CUDAよりエコシステムが成熟しており，ドキュメント・最適化ノウハウ・ツール（Nsightなど）も揃っています。
実験内容
ベンチマークの概要
今回は、計算負荷の高いタスクとして 正方行列の乗算 ($C = A \times B$) を実施しました。
環境依存の問題を排除するため、Windows 11 (WSL2) 上に Docker + NVIDIA Container Toolkit で構築した統一環境を使用しています。
検証環境:
- GPU: NVIDIA GeForce RTX 3060 Laptop
- Rust: Nightly toolchain
- CUDA: 12.8.1
検証フェーズ
検証は以下の2段階で行いました。
-
RustにおけるGPU開発の実用性の検証: 単精度浮動小数点数（
f32）を用いて、Rustで書いたカーネルがC++公式実装にどこまで迫れるか。 - CPUとGPUでの多倍長行列演算のパフォーマンスの検証: 多倍長整数（64bit, 128bit, 256bit）において、CGBN（GPU）はCUDA（GPU）とGMP（CPU）に対してそれぞれ，どの程度の性能差があるか。
実験結果と考察
実験1. Rust-CUDA vs C++-CUDA
まずは「RustでGPUコードを書いた場合、C++と比べて遅くなるのか？」を確認しました。
ここでは、RustでGPUコードを書くために、Rust-CUDAを用いました。Rust-CUDAでは、GPU計算の高速化に不可欠な共有メモリ（Shared Memory）を利用したタイリング最適化をする場合、としない場合があります。
実験では、タイリング最適化をしないRust-CUDA (Rust Naive)と、タイリング最適化をするRust-CUDA(Rust Tiled) また、C++のCUDAを用いた( C++ CUDA ) それぞれ実験しました。
【計測条件】
- 行列サイズ: $2048 \times 2048$
-
データ型:
f32(単精度浮動小数点数)
| 実装手法 | GFLOPS | C++公式比 | 備考 |
|---|---|---|---|
| Rust Naive | 225.7 | 25.3% | グローバルメモリへ直接アクセス |
| Rust Tiled | 498.5 | 55.9% | 共有メモリ ($16 \times 16$) 利用 |
| C++ CUDA | 891.9 | 100% | 共有メモリ ($32 \times 32$) 利用 |
結果
Rust Tiled版は、Naive版（225.7 GFLOPS）に比べて約2.21倍となる498.5 GFLOPSを達成しました。しかしながら、C++版（891.9 GFLOPS）には及びませんでした。
考察
-
タイルサイズによる性能差
C++版は $32 \times 32$ のタイルサイズを採用していますが、今回のRust版では $16 \times 16$ に留めています。これがスループットの差に直結しています。 -
なぜRustで $32 \times 32$ が動かなかったのか？
下記がRustのカーネルから見れる1スレッドあたりのレジスタ消費量です。- Rust ($16 \times 16$): 約70 レジスタ/スレッド
- Rust ($32 \times 32$): 約150 レジスタ/スレッド
使用したRTX 3060の1SM（ストリーミングマルチプロセッサ）あたりのレジスタ容量は 65,536 です。
$32 \times 32$ タイルを実行するには 1,024スレッド/ブロック が必要ですが、Rust版の消費量で計算すると：
$$65,536 \div 150 \approx 436 \text{ Threads}$$
となり、物理的に必要なスレッド数を確保できず、LaunchOutOfResourcesエラーが発生しました。 -
結論：最適化の余地
現在のRust-CUDA（NVPTXバックエンド）が出力するPTXコードは、C++（NVCC）ほどレジスタ割り当てが最適化されていません。
$32 \times 32$ タイルが動作すれば、RustでもC++と同等の性能が狙えるんじゃないでしょうか。（レジスタ使用量をどう減らすかが問題ですが）
実験2.GMP vs C++-CUDA vs CGBN
次に、多倍長整数の行列演算比較です。
多倍長演算をGPU上で行うCGBNのパフォーマンスを、GPU上で演算を行うデファクトスタンダードであるC++ CUDAと、多倍長演算をCPU上で行うデファクトスタンダードであるGMPと比較しました。今回，言語はC++で統一しています。
実験 2-1: 64bit整数
【計測条件】
- 行列サイズ: $2048 \times 2048$
-
データ型:
u64(64bit 符号なし整数)
| 実装手法 | 実行時間 | GFLOPS | GMP比 |
|---|---|---|---|
| CUDA (GPU) | 0.033 秒 | 517.3 | 約 19,960倍 |
| CGBN (GPU) | 3.00 秒 | 5.72 | 約 220倍 |
| GMP (CPU) | 662.8 秒 | 0.026 | 1.0 (基準) |
実験 2-2: 128bit / 256bit整数
GMPの処理時間が現実的になるよう、行列サイズを落として比較しました。
【計測条件】
- 行列サイズ: $256 \times 256$
-
データ型:
u128/u256(多倍長整数)
| ビット幅 | CGBN (GPU) | GMP (CPU) | 高速化倍率 |
|---|---|---|---|
| 128 bit | 6.9 ms | 343.3 ms | 約 49倍 |
| 256 bit | 12.8 ms | 537.1 ms | 約 42倍 |
結果
実験1ではCUDA，CGBNがGMPと比較して実行速度がそれぞれ約2万倍，約220倍の優位性を示しました。実験2ではCGBNはGMPよりも実行速度が40倍以上高速でした。
考察
-
GPUの圧倒的優位性:
通常の64bit整数演算 ($N=2048$) において、CUDA（GPU）はGMP（CPU）に対し 約2万倍 という驚異的な速度差があることがわかります。また、多倍長（128/256bit, $N=256$）においても、CGBNを使うことでCPUに対し 約40〜50倍 の高速化が確認できました。暗号計算のオフロード先としてGPUは非常に強力であることがわかります。 -
CGBNがC++ CUDAより遅い理由:
64bitの結果を見ると、CGBNは通常のC++ CUDA実装よりもかなり低速です。この原因は主に以下の2点だと考えられます。- グローバルメモリへの依存: CGBNの実装は、高速な共有メモリではなく、低速なグローバルメモリへのアクセス頻度が高い構造になっています。
- スレッド同期コスト: 複数のスレッドで1つの巨大な数値を表現するため、計算のたびにスレッド間の同期（ワープシャッフル等）が発生し、オーバーヘッドとなっています。
-
多倍長演算においてCGBNがGMPより高速な理由
128ビットおよび256ビット整数を用いた実験結果から、CGBNはGMPよりも高速であることがわかります。この性能差の主な理由は、GPUによる並列処理の有無です。CGBN実装では数万から数十万のスレッドが同時並列で行列要素を計算するのに対し、GMP実装ではCPU1コアで全要素を逐次処理しています。例えば、n=512の行列では262,144個の出力要素をCGBNは並列計算しますが、GMPは1つずつ順次計算するため、並列度に数万倍以上の差が生じています。
結論
今回のRust×GPU×多倍長整数の検証から得られた結論は以下の通りです。
-
RustでGPU開発は現実的
実験1より、Rust-CUDAは行列積（f32, 2048×2048）で、Rust Naiveが225.7 GFLOPS、Rust Tiledが498.5 GFLOPSを達成し、約2.21倍高速化できました。
一方でC++ CUDA（891.9 GFLOPS）には届かず、C++比55.9%でした。
ただし、この差は実験1の考察でも述べた通り、Rust-CUDAがC++ CUDAほど最適化されていないことが原因だと思われるため、これから先Rust-CUDAの共有メモリを使った最適化が進むにつれて、性能差は埋まると考えられます。CUDAに必須となるメモリ階層の制御をRustで実装できていることからRustでのGPU開発は現実的だと思われます。 -
多倍長計算におけるGPU使用の効果は大きい
実験2-1より、64bit行列積（2048×2048）ではCUDA（GPU）がGMP（CPU）より約2万倍高速でした。CGBN（GPU）もGMP比で約220倍高速でした。また実験2-2より、多倍長（256×256）でもCGBNはGMPより高速で、128bitで約49倍、256bitで約42倍でした。以上の実験結果から64bit、多倍長にかかわらずCPUよりもGPUでの実装のほうが高速であることがわかります。そのため、少なくとも今回実験したような演算においてGPUは効果的だと言えます。
既存のRust-CUDAの実装は汎用的である反面、メモリアクセス最適化の余地（共有メモリの活用など）が残されています。Rustを用いてC++での実装のような、より高速なパフォーマンスを出せる可能性があります。また、CGBNについてもGMPでの実装より高速であることから、多倍長演算をGPUのCGBNでおこない、そのラッパーをRustで書くようなことも実装の際の選択肢として挙がってくるのではないかと考えられます。
総括としては、Rustでの巨大な行列演算と多倍長整数演算は現状だとC++と比べて若干遅いですが、今後のライブラリの発展次第ではより高速なパフォーマンスを出せる可能性があり、高速な暗号ライブラリ開発の言語選択において候補に挙がってくるのではないかと思います。