TadaoYamaokaの開発日記

個人開発しているスマホアプリや将棋AIの開発ネタを中心に書いていきます。

将棋AI実験ノート:入力特徴量の転送量削減

以前にCUDAマルチストリームに対応した際、8GPUだと、CPU-GPU間の帯域がボトルネックとなり、NPSが上がらない課題があることがわかった。
対策として、FP16で転送することを検討し、NPSが平均で18.4%向上することを確認した。

GPU数とNPS

GPU数とNPSの関係を調べたところ以下の通りであった。
f:id:TadaoYamaoka:20220110202547p:plain

FP32ではGPUを増やしてもNPSが線形に伸びていないことが確認でき、CPU-GPU間の転送がボトルネックになっていることがわかる。

転送量の削減

年末にDiscordでやりとりしている中で、CPU-GPU間の転送がボトルネックを解消するアイディアとして、入力特徴量をbitで送る案をMizarさんが提案し、さっそく実装して、NPSが大幅に向上すると報告があった。

現状の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のプロパティで、カスタムビルドツールを選択し、
f:id:TadaoYamaoka:20220110205519p:plain

全般のコマンドラインと出力ファイルに以下のように設定する。
f:id:TadaoYamaoka:20220110205550p:plain

コマンドライン

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

LinuxMakefileでは、.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による条件分けは未実装