AWS inf1 インスタンスで推論を行うその3

前回、neuronコアを複数使用して推論を行うことを試した。
今回は、推論速度の測定を行った。

参考にしたいのは、dlshogiでの推論速度であるため、C++のdlshogiの推論部分に組み込んで、どれくらいのNPSがでるのかを確認した。

C++での推論の実装方法

C++でneuronコアを使用するには、LibTorchを使用する方法が提供されている。
LibTorch C++ Tutorial — AWS Neuron documentation

公式のLibTorchのチュートリアルのsetup.shを確認すると、
CPU版のLibTorchを使用して、Pythonのtorch-neuronに含まれるlibneuron_op.soをリンクすることで、LibTorchでneuronコアを使用できるようにしている。

setup.shから必要な部分を抜き出して実行してもよいが、今回は一旦チュートリアルのsetup.shを実行して、そこから必要なファイルをMakefileから使用するようにした。

チュートリアルのビルド

公式のページの手順の通り、実行する。

$ wget https://awsdocs-neuron.readthedocs-hosted.com/en/latest/_downloads/172f30d74345f2994f0bb5536f7dc650/libtorch_demo.tar.gz
$ tar xvf libtorch_demo.tar.gz
$ cd libtorch_demo
$ chmod +x setup.sh && ./setup.sh

dlshogiのビルドには必要ないが、チュートリアルのビルドにはRustが必要なため、事前にインストールしておく必要がある。

$ sudo apt install -y cargo

実行が成功すると、
/home/ubuntu/libtorch_demo/libtorch/lib
に必要なファイルがコピーされる。

ソース修正

dlshogiのソースをLibTorchで推論できるように修正した。
修正したソースは、feature/libtorchブランチにプッシュしている。
GitHub - TadaoYamaoka/DeepLearningShogi at feature/libtorch

ビルド

チュートリアルはCMakeでビルドするようになっているが、dlshogiではmakeを使用しているため、チュートリアルのcmakeで作成されたmakeファイル(build.makeとlink.txt)を参考にMakefileを作成した。
完全に理解していないが、リンカのオプションで、LibTorchの.soをlibneuron_op.soに置き換えることを行っていそうである。

CC = g++
MATE_SEARCH_DEPTH = 5
CFLAGS = -std=c++17 -Wextra -Ofast -MMD -MP -fopenmp -DLIBTORCH -DPV_MATE_SEARCH -DWIN_TYPE_DOUBLE -DMATE_SEARCH_DEPTH=$(MATE_SEARCH_DEPTH) -DTHREAD_POOL -DNDEBUG -DHAVE_SSE4 -DHAVE_SSE42 -DHAVE_BMI2 -msse4.2 -mbmi2 -DHAVE_AVX2 -mavx2 -D_GLIBCXX_USE_CXX11_ABI=0
LDFLAGS = -flto -Wl,-rpath,/home/ubuntu/libtorch_demo/libtorch/lib /home/ubuntu/libtorch_demo/libtorch/lib/libtorch.so /home/ubuntu/libtorch_demo/libtorch/lib/libc10.so /home/ubuntu/libtorch_demo/libtorch/lib/libneuron_op.so -Wl,--no-as-needed,"/home/ubuntu/libtorch_demo/libtorch/lib/libtorch_cpu.so" -Wl,--as-needed /home/ubuntu/libtorch_demo/libtorch/lib/libc10.so -lpthread -Wl,--no-as-needed,"/home/ubuntu/libtorch_demo/libtorch/lib/libtorch.so" -Wl,--as-needed
INCLUDE = -I../usi_libtorch -I../usi -I../cppshogi -I../cppshogi -I/home/ubuntu/libtorch_demo/libtorch/include -I/home/ubuntu/libtorch_demo/libtorch/include/torch/csrc/api/include
LIB = -L/home/ubuntu/libtorch_demo/libtorch/lib

target = bin/usi
target_make_book = bin/make_book
sources = nn_libtorch.cpp
usi_sources = main.cpp dfpn.cpp Message.cpp UctSearch.cpp Node.cpp PvMateSearch.cpp
cppshogi_sources = bitboard.cpp book.cpp common.cpp cppshogi.cpp generateMoves.cpp hand.cpp init.cpp move.cpp mt64bit.cpp position.cpp search.cpp square.cpp usi.cpp
objects = $(addprefix obj/, $(sources:.cpp=.o))
usi_objects = $(addprefix obj/, $(usi_sources:.cpp=.o))
cppshogi_objects = $(addprefix obj/, $(cppshogi_sources:.cpp=.o))

$(target): $(objects) $(usi_objects) $(cppshogi_objects)
	@[ -d bin ] || mkdir -p bin
	$(CC) -o $@ $^ $(LIB) $(LDFLAGS) $(CFLAGS)

obj/%.o: %.cpp
	@[ -d obj ] || mkdir -p obj
	$(CC) $(CFLAGS) $(INCLUDE) -o $@ -c $<

obj/%.o: ../usi/%.cpp
	@[ -d obj ] || mkdir -p obj
	$(CC) $(CFLAGS) $(INCLUDE) -o $@ -c $<

obj/%.o: ../cppshogi/%.cpp
	@[ -d obj ] || mkdir -p obj
	$(CC) $(CFLAGS) $(INCLUDE) -o $@ -c $<

all: $(target)

clean:
	rm -f $(objects) $(cppshogi_objects) $(target)

$(target_make_book): $(objects) $(cppshogi_objects)
	@[ -d bin ] || mkdir -p bin
	$(CC) -o $@ $^ $(LIB) $(LDFLAGS) $(CFLAGS)

make_book: CFLAGS += -DMAKE_BOOK -DWIN_TYPE_DOUBLE
make_book: $(target_make_book)

測定

10ブロックのモデル

10ブロック192フィルタのモデルを使って、初期局面でのNPSを測定した。

USIオプションのDNN_Modelには、以前記事にした方法で変換したaws neuron用のモデルを指定する。
inf1.xlargeインスタンスのvCPUは4しかないため、GPUあたりのスレッド数(UCT_Threads)は2で測定した。

neuronコア1つでの測定結果は、以下の通り。

setoption name DNN_Model value /home/ubuntu/model-pre5_resnet15_swish_b4096lr004-008.neuron.pt
setoption name UCT_Threads2 value 2
isready
position startpos
go byoyomi 1000
info nps 1285 time 39064 nodes 50204 hashfull 5 score cp 189 depth 23 pv 2g2f 3c3d 7g7f 8c8d 2f2e 8d8e 6i7h 8e8f 8g8f 8b8f 2e2d 2c2d 2h2d 4a3b 2d3d 2b3c 5i5h 5a5b 3g3f 8f7f 8h7g 7c7d P*2b
bestmove 2g2f

NPSは、1285であった。

参考として、NVIDIAのGPU RTX3090でのOnnxRuntimeとTensorRTを使用した際のNPSはそれぞれ以下の通りである。

バージョン	NPS
OnnxRuntime版	5715
TensorRT版	41280

OnnxRuntime版に対して約22.5%、TensorRT版に対して約3.1%の推論速度である。
カタログスペックから期待されるような速度はでていない。

浮動小数点の精度は自動で変換される仕組みのため、FP16で計算できているはずである(変換したモデルサイズも元の半分になっている)。
それにしても遅い。

neuronコアが使用されていることは、neuron-topコマンドで確認できている。

neuron-top - 08:40:18
Models: 1 loaded, 1 running. NeuronCores: 1 used.
0000:00:1f.0 Utilizations: NC0 48.64%, NC1 0.00%, NC2 0.00%, NC3 0.00%,
Model ID   Device    NeuronCore%   Device Mem   Host Mem   Model Name
10008      nd0:nc0   48.64           14 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmp4ki80ors

マルチコアでの測定

vCPUが4しかないため、おそらくCPUがボトルネックになるが、2コアと4コアでも試してみた。
モデルを複数ロードするとneuronコアが割り当てられるため、特にソースの修正はいらない(USIオプションのUCT_Threads2、UCT_Threads3、・・・を設定すればよい)。

コア数	NPS
1コア	1285
2コア	1868
4コア	2727

コアを増やすのNPSが上がることが確認できた。
NPSが線形に伸びていないのは、おそらくCPUボトルネックになっているためである。

neuron-topコマンドで、コアの使用率を確認すると、1コアの時より低くなっている。

2コアの場合：

neuron-top - 08:41:18
Models: 2 loaded, 2 running. NeuronCores: 2 used.
0000:00:1f.0 Utilizations: NC0 42.08%, NC1 42.11%, NC2 0.00%, NC3 0.00%,
Model ID   Device    NeuronCore%   Device Mem   Host Mem   Model Name
10010      nd0:nc1   42.11           14 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmp4ki80ors
10009      nd0:nc0   42.08           14 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmp4ki80ors

4コアの場合：

neuron-top - 08:43:15
Models: 4 loaded, 4 running. NeuronCores: 4 used.
0000:00:1f.0 Utilizations: NC0 34.38%, NC1 34.79%, NC2 34.71%, NC3 34.59%,
Model ID   Device    NeuronCore%   Device Mem   Host Mem   Model Name
10014      nd0:nc3   34.59           14 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmp4ki80ors
10013      nd0:nc2   34.71           14 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmp4ki80ors
10012      nd0:nc1   34.79           14 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmp4ki80ors
10011      nd0:nc0   34.38           14 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmp4ki80ors

15ブロックでの測定

15ブロック224フィルタのモデルでも測定した。

コア数	NPS
1コア	713
2コア	1285
4コア	1897

10ブロックの比較して、1コアの場合でNPSが、約55.5%になっている。
RTX 3090でTensorRTを使用した場合、10ブロック 49106、15ブロック 31243で、約63.6%になる。

neuronコアの使用率は、

2コアの場合：

neuron-top - 08:27:36
Models: 2 loaded, 2 running. NeuronCores: 2 used.
0000:00:1f.0 Utilizations: NC0 64.75%, NC1 64.69%, NC2 0.00%, NC3 0.00%,
Model ID   Device    NeuronCore%   Device Mem   Host Mem   Model Name
10003      nd0:nc1   64.69           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q
10002      nd0:nc0   64.75           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q

4コアの場合：

neuron-top - 08:29:36
Models: 4 loaded, 4 running. NeuronCores: 4 used.
0000:00:1f.0 Utilizations: NC0 64.95%, NC1 64.75%, NC2 64.95%, NC3 64.89%,
Model ID   Device    NeuronCore%   Device Mem   Host Mem   Model Name
10007      nd0:nc3   64.89           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q
10006      nd0:nc2   64.95           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q
10004      nd0:nc0   64.95           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q
10005      nd0:nc1   64.75           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q

10ブロックよりも使用率が高く、コア数を増やしても使用率が下がらないことから、GPUもボトルネックになっていそうである(線形に伸びていないのでCPUもボトルネック)。

まとめ

dlshogiでneuronコアを使用して、初期局面でNPSがどれくらいでるか測定した。
結果、NVIDIAの最新のGPUでTensorRTを使用した場合と比較して、推論速度は1コアで約3.1%程度であった。

1つのInferentia チップで、4コア使用できるため、NPSが線形に伸びると仮定すると、約12%程度である。
ただし、今回測定したインスタンスはvCPUが少ないため、線形には伸びなかった。
なお、OnnxRuntime版と比較すると、1コアで約22.5%の推論速度のため、4コアでNPSが線形に伸びると仮定すると、約90%の推論速度になる。

NVIDIAのT4が使用できるインスタンスとの価格差が半分程度であることを考慮すると、T4のインスタンスでTensorRTを使用した方がお得かもしれない。
ただし、T4はRTX 3090より性能は低いため測定して比較してみる必要がある。

dlshogiでのNPSはあまりでないことが分かったが、TensorRTはライセンス的にインスタンを他の人と共有することができないため、inf1インスタンスでneuronコアを使用するのも用途はありそうである。

TadaoYamaokaの開発日記

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