以前にCUDAマルチストリームに対応した際、8GPUだと、CPU-GPU間の帯域がボトルネックとなり、NPSが上がらない課題があることがわかった。
対策として、FP16で転送することを検討し、NPSが平均で18.4%向上することを確認した。
GPU数とNPS
GPU数とNPSの関係を調べたところ以下の通りであった。
FP32ではGPUを増やしてもNPSが線形に伸びていないことが確認でき、CPU-GPU間の転送がボトルネックになっていることがわかる。
転送量の削減
年末にDiscordでやりとりしている中で、CPU-GPU間の転送がボトルネックを解消するアイディアとして、入力特徴量をbitで送る案をMizarさんが提案し、さっそく実装して、NPSが大幅に向上すると報告があった。
昨年末、やねさん山岡さん48さんと議論してたホスト-GPU間の転送帯域を減らす実装をTensorRT版ふかうら王で試してみた。マルチストリームはまだ未実装。
— Mizar/みざー (@mizarjp) 2022年1月6日
resnet10x192で NPS 60832 → 66642
resnet5x192で NPS 86047 → 103426
8GPUマルチストリーム化でどうなるかは山岡さんの実験に期待。
現状のdlshogiでは、入力特徴量は、9×9の2値画像×119チャンネルを、浮動小数(FP32)で転送している。
各画素の値は、0か1のため、情報量は1bitで表せる。
また、持ち駒の数など特徴量は、9×9の画像のすべての画素を0か1にしているため、チャンネルを1bitで表すことができる。
そこで、画素またはチャンネルの値を1bitで転送し、GPU側で浮動小数にすることで、転送量を削減できる。
dlshogiでも、この転送量の削減を実装し、NPSがどれくらい上がるか測定してみた。
GPU側で展開する処理
転送したデータの各bitを、浮動小数の画像に展開する処理は、CUDAのプログラムで実装する。
画素の浮動小数は、FP16とする。
Mizarさんの実装では、NVRTCを使用して、実行時にCUDAのプログラムをコンパイルしていたが、初期化の処理が増えるため、事前にnvccでコンパイルしてリンクするようにした。
また、Mizaerさんの実装では、全てのバッチのデータをbit単位で詰めて転送量を減らしているが、バッチ単位でデータクリアする処理が煩雑になるため、バッチ単位では、バイト境界をまたがないように実装した（転送量は少し増えるがCPU側の処理は軽くなり、既存のコードと処理を共通化しやすくなる）。
CUDAプログラム(unpack.cu)
#include "unpack.h"
constexpr int features1_size = sizeof(features1_t) / sizeof(DType) / SquareNum;
constexpr int features2_size = sizeof(features2_t) / sizeof(DType) / SquareNum;
__global__ void unpack_features1_kernel(char* p1, short* x1) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int p1_offset = sizeof(packed_features1_t) * 8 * blockIdx.x + threadIdx.x * 81;
int x1_offset = tid * 81;
#pragma unroll
for (int i = 0; i < 81; ++i) {
int j = p1_offset + i;
// p1[j / 8] >> (j % 8)で下位1bitに設定する値を持ってくる
// 下位1bitのマスクを行い、符号を負にすることで1の場合1byteの全bitを1にする
// 0x3c00と論理積を取ることでfloat16の1.0にする
x1[x1_offset + i] = (-(short)((p1[j >> 3] >> (j & 7)) & 1)) & 0x3c00;
}
}
__global__ void unpack_features2_kernel(char* p2, short* x2) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int j = sizeof(packed_features2_t) * 8 * blockIdx.x + threadIdx.x;
short v = (-(short)((p2[j >> 3] >> (j & 7)) & 1)) & 0x3c00;
int x2_offset = tid * 81;
#pragma unroll
for (int i = 0; i < 81; ++i) {
x2[x2_offset + i] = v;
}
}
void unpack_features1(const int batch_size, packed_features1_t* p1, features1_t* x1, cudaStream_t stream)
{
unpack_features1_kernel<<<batch_size, features1_size, 0, stream>>>((char*)p1, (short*)x1);
}
void unpack_features2(const int batch_size, packed_features2_t* p2, features2_t* x2, cudaStream_t stream)
{
unpack_features2_kernel<<<batch_size, features2_size, 0, stream>>> ((char*)p2, (short*)x2);
}
nvccを使用するVisual Studioのプロジェクト設定
.cuをnvccでコンパイルするように、Visual Studioのプロジェクト設定を行う。
NVIDIA CUDA Visual Studio Integrationを使うと、特に設定しなくても.cuをコンパイルできるが、ビルド環境構築の前提条件が増えてしまうため、カスタムビルドツールの設定で対応した。
.cuのプロパティで、カスタムビルドツールを選択し、
全般のコマンドラインと出力ファイルに以下のように設定する。
コマンドライン
Releaseビルド
"$(CUDA_PATH_V11_1)\bin\nvcc.exe" --use-local-env -ccbin "$(VC_ExecutablePath_x86_x64)" -x cu --keep-dir $(Platform)\$(TargetedSDKConfiguration) -maxrregcount=0 --machine 64 --compile -cudart static -DFP16 -DNDEBUG -D_CONSOLE -D_UNICODE -DUNICODE -Xcompiler "/EHsc /W3 /nologo /O2 /Fdx64\Release\vc142.pdb /FS /MT" -o "$(Platform)\$(TargetedSDKConfiguration)\%(Filename)%(Extension).obj" "%(FullPath)"
Debugビルド
"$(CUDA_PATH_V11_1)\bin\nvcc.exe" --use-local-env -ccbin "$(VC_ExecutablePath_x86_x64)" -x cu -G --keep-dir $(Platform)\$(TargetedSDKConfiguration) -maxrregcount=0 --machine 64 --compile -cudart static -g -DFP16 -D_DEBUG -D_CONSOLE -D_UNICODE -DUNICODE -Xcompiler "/EHsc /W3 /nologo /Od /Fdx64\Debug\vc142.pdb /FS /Zi /RTC1 /MTd" -o "$(Platform)\$(TargetedSDKConfiguration)\%(Filename)%(Extension).obj" "%(FullPath)"
出力ファイル
$(Platform)\$(TargetedSDKConfiguration)\%(Filename)%(Extension).obj
Makefile
LinuxのMakefileでは、.cuをnvccコマンドでコンパイルする定義を追加する。
obj/unpack.cu.o: unpack.cu nvcc -x cu -I../cppshogi -maxrregcount=0 --machine 64 --compile -cudart static -DFP16 -DNDEBUG -Xcompiler "-Ofast" -o obj/unpack.cu.o unpack.cu
測定結果
floodgateの棋譜からサンプリングした100局面で、1秒探索した際のNPSの統計は以下の通り。
各局面を10回測定して平均値を使用した。
モデルサイズは15ブロック224フィルタ。
RTX3090、2スレッド
| master | FP16 | pack | FP16/master | pack/master | |
|---|---|---|---|---|---|
| 平均 | 29417 | 29689 | 30518 | 100.8% | 103.7% |
| 中央値 | 30238 | 30726 | 31278 | 101.6% | 103.9% |
| 最大値 | 31593 | 32261 | 33232 | 103.1% | 105.7% |
| 最小値 | 18337 | 17976 | 19077 | 96.8% | 99.6% |
※masterが現在のdlshogi
※FP16が転送時の浮動小数をFP16にしたバージョン
※packが今回実装した画素をbitで転送して、GPU側で浮動小数に展開したバージョン
※比は、同じ局面での比
RTX3090では、平均で3.9%NPSが向上した。
Mizarさんの報告では10ブロックで10%弱、5ブロックで12%NPSが向上しているが、15ブロックでは、GPUの処理の方がボトルネックとなるようである。
A100x8、GPUあたり4スレッド
| master | FP16 | pack | FP16/master | pack/master | |
|---|---|---|---|---|---|
| 平均 | 287561 | 346906 | 392338 | 120.7% | 136.6% |
| 中央値 | 289233 | 351813 | 398223 | 120.9% | 136.6% |
| 最大値 | 339627 | 394866 | 438219 | 149.8% | 179.2% |
| 最小値 | 221239 | 257580 | 291222 | 102.7% | 109.3% |
A100x8では、平均で36.6%NPSが向上した。
RTX3090 1枚に比べて、大幅に向上している。
8GPUではCPU-GPU間の帯域がボトルネックになっていたことがわかる。
まとめ
入力特徴量のCPU-GPU間の転送量を削減することで、8GPUの場合にNPSが平均で36.6%向上することがわかった。
今回の実験によっても、8GPUではCPU-GPU間の帯域がボトルネックとなっていたことが裏付けられた。
ニューラルネットワークの出力についても、現在は非合法手の指し手の確率も転送しているが、合法手をbitで表現して送って、転送量を減らす案がある。
こちらも別途検証してみたい。
また、CUDAマルチストリームを使用した場合についても別途検証したい。
ソース
feature/pack_featuresブランチにプッシュしている。
GitHub - TadaoYamaoka/DeepLearningShogi at feature/pack_features
※既存処理とのifdefによる条件分けは未実装
