TadaoYamaokaの開発日記

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

OpenCLのサンプルコード

ほぼ個人用メモ。

「改訂新版 OpenCL入門 1.2対応 マルチコアCPU・GPUのための並列プログラミング」を買って、OpenCLをお勉強中。


OpenCL 1.2とバージョンが古いが、KataGoもOpenCL 1.2で実装されているので問題ない。
そもそもNVIDIAOpenCLに対してやる気がなく、1.2しかサポートしていないようだ。
OpenCL 2.x support plans? - CUDA Programming and Performance - NVIDIA Developer Forums


以下、サンプルコードのWindowsでのビルド手順。

OpenCLSDK

CUDA Toolkitに含まれるため、追加でインストールはいらない。
※書籍では、GPU Computing SDKが必要とあるが、現在はCUDAに含まれる。
※ただし、OpenCLのサンプルプログラムは提供されなくなっている。
Where is gpu computing sdk ? - CUDA Setup and Installation - NVIDIA Developer Forums

プロジェクトの設定

OpenCLのインクルードファイルとライブラリは、CUDAをインストールすると以下の場所にある。
インクルード:C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\include
ライブラリ:C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\lib\x64

それぞれ、プロジェクトのインクルードとライブラリのディレクトリ設定に追加する。
リンクライブラリに、「OpenCL.lib」を追加する。

サンプルコード

書籍のサンプルコードでは、.clを動的にファイルから読むようになっていたが、ソースに埋め込むようにした。

KataGoの書き方を参考にして、

	const char source_str[] =
#include "hello.cl"
;

のようにして、文字列にした。

.clの先頭と末尾には、

R"%%(

// OpenCLのコード

)%%"

が必要になる。

修正したサンプルコード

hello.cpp
#include <iostream>

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

constexpr size_t MEM_SIZE = 128;

int main()
{
	cl_device_id device_id = nullptr;
	cl_context context = nullptr;
	cl_command_queue command_queue = nullptr;
	cl_mem memobj = nullptr;
	cl_program program = nullptr;
	cl_kernel kernel = nullptr;
	cl_platform_id platform_id = nullptr;
	cl_uint ret_num_devices;
	cl_uint ret_num_platforms;
	cl_int ret;

	char string[MEM_SIZE];

	const char source_str[] =
#include "hello.cl"
;
	size_t source_size = sizeof(source_str) - 1;
	const char* source_list[] = { source_str };

	/* プラットフォーム・デバイスの情報の取得 */
	ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
	ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);

	/* OpenCLコンテキストの作成 */
	context = clCreateContext(nullptr, 1, &device_id, nullptr, nullptr, &ret);

	/* コマンドキューの作成 */
	command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

	/* メモリバッファの作成 */
	memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(char), nullptr, &ret);

	/* 読み込んだソースからカーネルプログラムを作成 */
	program = clCreateProgramWithSource(context, 1, source_list,
		&source_size, &ret);

	/* カーネルプログラムをビルド */
	ret = clBuildProgram(program, 1, &device_id, nullptr, nullptr, nullptr);

	/* OpenCLカーネルの作成 */
	kernel = clCreateKernel(program, "hello", &ret);

	/* OpenCLカーネル引数の設定 */
	ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&memobj);

	/* OpenCLカーネルを実行 */
	ret = clEnqueueTask(command_queue, kernel, 0, nullptr, nullptr);

	/* メモリバッファから結果を取得 */
	ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0,
		MEM_SIZE * sizeof(char), string, 0, nullptr, nullptr);

	/* 結果の表示 */
	std::cout << string << std::endl;

	/* 終了処理 */
	ret = clFlush(command_queue);
	ret = clFinish(command_queue);
	ret = clReleaseKernel(kernel);
	ret = clReleaseProgram(program);
	ret = clReleaseMemObject(memobj);
	ret = clReleaseCommandQueue(command_queue);
	ret = clReleaseContext(context);

	return 0;
}
hello.cl
R"%%(
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable

__kernel void hello(__global char* string)
{
   string[0] = 'H';
   string[1] = 'e';
   string[2] = 'l';
   string[3] = 'l';
   string[4] = 'o';
   string[5] = ',';
   string[6] = ' ';
   string[7] = 'W';
   string[8] = 'o';
   string[9] = 'r';
   string[10] = 'l';
   string[11] = 'd';
   string[12] = '!';
   string[13] = '\0';
}
)%%"

まとめ

サンプルコードをWindowsでビルドして実行した。
OpenCLは、OpenCL C言語ソースコードを文字列としてランタイムに渡して動的にコンパイルされる仕組みになっていることを理解した。

ONNXファイルのパース処理

dlshogiは現在TensorRTを使用しているが、TensorRTは再配布できないなどライセンスが厳しいのと、環境構築が大変なため、OpenCLに対応させたいと思っている。

OpenCLでも、NVIDIAGPUのTensorCoreをPTXインラインアセンブリという方法で使用することができるため、上手に実装すればパフォーマンスは同等にできる見込みである。
KataGoのv1.5.0で、OpenCLでのTensorCoreサポートが実装されているため、ソースを流用すればなんとか実装できると思っている。
ただし、活性化関数にswishを使っていたりといった違いがあるため、そのままでは流用できず、KataGoのOpenCLのソースを理解した上で取り込む必要があり、なかなかハードルが高い。

また、dlshogiはモデルファイルのフォーマットにONNXを採用しているため、モデルファイルをパースして、グラフ定義とパラメータを取り出す処理も実装が必要である。

ということで初めの一歩として、ONNXファイルのパース処理を試してみた。

ONNXファイルのフォーマット

ONNXファイルのフォーマットは、公式サイトで公開されている。
onnx/IR.md at main · onnx/onnx · GitHub

フォーマットは、Protocol Buffersで定義されており、onnx.proto3から、各言語向けのパース用コードが生成できる。

Protocol Buffersのインストール

WindowsC++向けProtocol Buffersのインストールを行う。

インストール方法は、公式のページに記載がある。
protobuf/src at master · protocolbuffers/protobuf · GitHub

C++ Installation - Windows」の記載内容にしたがって、vcpkgを使用してインストールする。
vcpkgのインストールは、vcpkgのページに従って行う。
GitHub - microsoft/vcpkg: C++ Library Manager for Windows, Linux, and MacOS

スタティックビルドを行いため、vcpkgからインストールする際に
x64-windows-static
を指定する。
インストールコマンドは以下の通りになる。

>vcpkg install protobuf protobuf:x64-windows-static

.proto3からC++に変換

.proto3からC++のコードへの変換は、protocというツールで行う。
protocは以下の場所にインストールされている。

C:\vcpkg\packages\protobuf_x64-windows\tools\protobuf\protoc.exe

onnxのソースをgit cloneして、以下のコマンドを実行するとC++のコードが生成される。

>git clone https://github.com/onnx/onnx.git
>cd onnx
>C:\vcpkg\packages\protobuf_x64-windows\tools\protobuf\protoc.exe --cpp_out=. onnx.prot3

生成されるソースコード

  • onnx.proto3.pb.cc
  • onnx.proto3.pb.h

コンパイル

生成されたソースをビルドするには、Protocol Buffersのインクルードディレクトリとライブラリディレクトリの設定が必要である。
Vissual Studioのプロジェクトの設定に、
インクルードディレクトリ:「C:\vcpkg\installed\x64-windows-static\include」
ライブラリディレクトリ:「C:\vcpkg\installed\x64-windows-static\lib」
を追加する。
また、リンクするライブラリに、「libprotobuf.lib」を追加する。

ビルド構成がDEBUGの場合は、
ライブラリディレクトリを「C:\\vcpkg\\installed\\x64-windows-static\\debug\\lib」に変更して、リンクするライブラリを「libprotobufd.lib」にする。

また、スタティックリンクするようにランタイムライブラリをスタティック版に設定する。

コンパイルエラーが起きる問題

生成されたソースコードをVisual C++コンパイルすると以下の箇所でコンパイルエラーが発生する。

  ::PROTOBUF_NAMESPACE_ID::internal::memswap<
      PROTOBUF_FIELD_OFFSET(AttributeProto, type_)
      + sizeof(AttributeProto::type_)
      - PROTOBUF_FIELD_OFFSET(AttributeProto, t_)>(
          reinterpret_cast<char*>(&t_),
          reinterpret_cast<char*>(&other->t_));

※同様の処理が何か所かある。

memswapのテンプレート引数が定数でないために発生している。
PROTOBUF_FIELD_OFFSETというマクロ定義の中で、reinterpret_castを使用しており、定数扱いにならないことが原因であった。

Protocol Buffersのバグのようだが、生成されたソースコードの方を修正することにした。
上記の箇所を、テンプレート引数が定数になるように以下のように修正した。

  switch (PROTOBUF_FIELD_OFFSET(AttributeProto, type_)
      + sizeof(AttributeProto::type_)
      - PROTOBUF_FIELD_OFFSET(AttributeProto, t_)) {
  case 0:
      ::PROTOBUF_NAMESPACE_ID::internal::memswap<0>(
              reinterpret_cast<char*>(&t_),
              reinterpret_cast<char*>(&other->t_));
      break;
  case 2:
      ::PROTOBUF_NAMESPACE_ID::internal::memswap<2>(
          reinterpret_cast<char*>(&t_),
          reinterpret_cast<char*>(&other->t_));
      break;
  case 8:
      ::PROTOBUF_NAMESPACE_ID::internal::memswap<8>(
          reinterpret_cast<char*>(&t_),
          reinterpret_cast<char*>(&other->t_));
      break;
  case 16:
      ::PROTOBUF_NAMESPACE_ID::internal::memswap<16>(
          reinterpret_cast<char*>(&t_),
          reinterpret_cast<char*>(&other->t_));
      break;
  }

※同様の修正を複数個所に行う。

パースのサンプルプログラム

ONNXファイルをパースして、内容を表示するサンプルプログラムを作成した。

#include <iostream>
#include <fstream>
#include "onnx.proto3.pb.h"

int main(int argc, char* argv[]) {
	GOOGLE_PROTOBUF_VERIFY_VERSION;

	if (argc != 2) {
		std::cerr << "Usage:  " << argv[0] << " onnx" << std::endl;
		return -1;
	}

	onnx::ModelProto model;

	std::fstream in(argv[1], std::ios::in | std::ios::binary);
	model.ParseFromIstream(&in);

	std::cout
		<< model.ir_version() << "\n"
		<< model.producer_name() << std::endl;

	const auto& graph = model.graph();

	std::cout << "---- inputs ----" << std::endl;
	for (int i = 0; i < graph.input_size(); i++) {
		const auto& input = graph.input(i);
		std::cout << input.name() << "\t";
		const auto& tensor_type = input.type().tensor_type();
		const auto& shape = tensor_type.shape();
		std::cout << tensor_type.elem_type() << "[";
		for (int n = 0; n < shape.dim_size(); n++) {
			if (n != 0) std::cout << ",";
			std::cout << shape.dim(n).dim_value();
		}
		std::cout << "]" << std::endl;
	}

	std::cout << "---- outputs ----" << std::endl;
	for (int i = 0; i < graph.output_size(); i++) {
		const auto& output = graph.output(i);
		std::cout << output.name() << "\t";
		const auto& tensor_type = output.type().tensor_type();
		const auto& shape = tensor_type.shape();
		std::cout << tensor_type.elem_type() << "[";
		for (int n = 0; n < shape.dim_size(); n++) {
			if (n != 0) std::cout << ",";
			std::cout << shape.dim(n).dim_value();
		}
		std::cout << "]" << std::endl;
	}

	std::cout << "---- nodes ----" << std::endl;
	for (int i = 0; i < graph.node_size(); i++) {
		const auto& node = graph.node(i);
		std::cout << node.name() << "\tinputs[";
		for (int n = 0; n < node.input_size(); n++) {
			if (n != 0) std::cout << ",";
			std::cout << node.input(n);
		}
		std::cout << "]\toutputs[";
		for (int n = 0; n < node.output_size(); n++) {
			if (n != 0) std::cout << ",";
			std::cout << node.output(n);
		}
		std::cout << "]\n";
	}

	std::cout << "---- initializers ----" << std::endl;
	for (int i = 0; i < graph.initializer_size(); i++) {
		const auto& initializer = graph.initializer(i);
		std::cout << initializer.name() << "\t";
		std::cout << initializer.data_type() << ":" << initializer.dims_size() << "[";
		for (int n = 0; n < initializer.dims_size(); n++) {
			if (n != 0) std::cout << ",";
			std::cout << initializer.dims(n);
		}
		std::cout << "]" << std::endl;
	}

	return 0;
}

dlshogiのモデルファイルを入力にして実行すると、以下のように表示される。

6
pytorch
---- inputs ----
input1  1[0,62,9,9]
input2  1[0,57,9,9]
---- outputs ----
output_policy   1[0,2187]
output_value    1[0,1]
---- nodes ----
Conv_0  inputs[input1,l1_1_1.weight]    outputs[142]
Conv_1  inputs[input1,l1_1_2.weight]    outputs[143]
Conv_2  inputs[input2,l1_2.weight]      outputs[144]
Add_3   inputs[142,143] outputs[145]
Add_4   inputs[145,144] outputs[146]
BatchNormalization_5    inputs[146,norm1.weight,norm1.bias,norm1.running_mean,norm1.running_var]        outputs[147]
Sigmoid_6       inputs[147]     outputs[148]
Mul_7   inputs[147,148] outputs[149]
Conv_8  inputs[149,256,257]     outputs[255]
Sigmoid_9       inputs[255]     outputs[152]
Mul_10  inputs[255,152] outputs[153]
Conv_11 inputs[153,259,260]     outputs[258]
Add_12  inputs[258,149] outputs[156]
Sigmoid_13      inputs[156]     outputs[157]
Mul_14  inputs[156,157] outputs[158]
Conv_15 inputs[158,262,263]     outputs[261]
Sigmoid_16      inputs[261]     outputs[161]
Mul_17  inputs[261,161] outputs[162]
Conv_18 inputs[162,265,266]     outputs[264]
Add_19  inputs[264,158] outputs[165]
Sigmoid_20      inputs[165]     outputs[166]
Mul_21  inputs[165,166] outputs[167]
Conv_22 inputs[167,268,269]     outputs[267]
Sigmoid_23      inputs[267]     outputs[170]
Mul_24  inputs[267,170] outputs[171]
Conv_25 inputs[171,271,272]     outputs[270]
Add_26  inputs[270,167] outputs[174]
Sigmoid_27      inputs[174]     outputs[175]
Mul_28  inputs[174,175] outputs[176]
Conv_29 inputs[176,274,275]     outputs[273]
Sigmoid_30      inputs[273]     outputs[179]
Mul_31  inputs[273,179] outputs[180]
Conv_32 inputs[180,277,278]     outputs[276]
Add_33  inputs[276,176] outputs[183]
Sigmoid_34      inputs[183]     outputs[184]
Mul_35  inputs[183,184] outputs[185]
Conv_36 inputs[185,280,281]     outputs[279]
Sigmoid_37      inputs[279]     outputs[188]
Mul_38  inputs[279,188] outputs[189]
Conv_39 inputs[189,283,284]     outputs[282]
Add_40  inputs[282,185] outputs[192]
Sigmoid_41      inputs[192]     outputs[193]
Mul_42  inputs[192,193] outputs[194]
Conv_43 inputs[194,286,287]     outputs[285]
Sigmoid_44      inputs[285]     outputs[197]
Mul_45  inputs[285,197] outputs[198]
Conv_46 inputs[198,289,290]     outputs[288]
Add_47  inputs[288,194] outputs[201]
Sigmoid_48      inputs[201]     outputs[202]
Mul_49  inputs[201,202] outputs[203]
Conv_50 inputs[203,292,293]     outputs[291]
Sigmoid_51      inputs[291]     outputs[206]
Mul_52  inputs[291,206] outputs[207]
Conv_53 inputs[207,295,296]     outputs[294]
Add_54  inputs[294,203] outputs[210]
Sigmoid_55      inputs[210]     outputs[211]
Mul_56  inputs[210,211] outputs[212]
Conv_57 inputs[212,298,299]     outputs[297]
Sigmoid_58      inputs[297]     outputs[215]
Mul_59  inputs[297,215] outputs[216]
Conv_60 inputs[216,301,302]     outputs[300]
Add_61  inputs[300,212] outputs[219]
Sigmoid_62      inputs[219]     outputs[220]
Mul_63  inputs[219,220] outputs[221]
Conv_64 inputs[221,304,305]     outputs[303]
Sigmoid_65      inputs[303]     outputs[224]
Mul_66  inputs[303,224] outputs[225]
Conv_67 inputs[225,307,308]     outputs[306]
Add_68  inputs[306,221] outputs[228]
Sigmoid_69      inputs[228]     outputs[229]
Mul_70  inputs[228,229] outputs[230]
Conv_71 inputs[230,310,311]     outputs[309]
Sigmoid_72      inputs[309]     outputs[233]
Mul_73  inputs[309,233] outputs[234]
Conv_74 inputs[234,313,314]     outputs[312]
Add_75  inputs[312,230] outputs[237]
Sigmoid_76      inputs[237]     outputs[238]
Mul_77  inputs[237,238] outputs[239]
Conv_78 inputs[239,l22.weight]  outputs[240]
Constant_79     inputs[]        outputs[241]
Reshape_80      inputs[240,241] outputs[242]
Add_81  inputs[242,l22_2.bias]  outputs[output_policy]
Conv_82 inputs[239,316,317]     outputs[315]
Sigmoid_83      inputs[315]     outputs[246]
Mul_84  inputs[315,246] outputs[247]
Constant_85     inputs[]        outputs[248]
Reshape_86      inputs[247,248] outputs[249]
Gemm_87 inputs[249,l23_v.weight,l23_v.bias]     outputs[250]
Sigmoid_88      inputs[250]     outputs[251]
Mul_89  inputs[250,251] outputs[252]
Gemm_90 inputs[252,l24_v.weight,l24_v.bias]     outputs[253]
Sigmoid_91      inputs[253]     outputs[output_value]
---- initializers ----
256     1:4[192,192,3,3]
257     1:1[192]
259     1:4[192,192,3,3]
260     1:1[192]
262     1:4[192,192,3,3]
263     1:1[192]
265     1:4[192,192,3,3]
266     1:1[192]
268     1:4[192,192,3,3]
269     1:1[192]
271     1:4[192,192,3,3]
272     1:1[192]
274     1:4[192,192,3,3]
275     1:1[192]
277     1:4[192,192,3,3]
278     1:1[192]
280     1:4[192,192,3,3]
281     1:1[192]
283     1:4[192,192,3,3]
284     1:1[192]
286     1:4[192,192,3,3]
287     1:1[192]
289     1:4[192,192,3,3]
290     1:1[192]
292     1:4[192,192,3,3]
293     1:1[192]
295     1:4[192,192,3,3]
296     1:1[192]
298     1:4[192,192,3,3]
299     1:1[192]
301     1:4[192,192,3,3]
302     1:1[192]
304     1:4[192,192,3,3]
305     1:1[192]
307     1:4[192,192,3,3]
308     1:1[192]
310     1:4[192,192,3,3]
311     1:1[192]
313     1:4[192,192,3,3]
314     1:1[192]
316     1:4[27,192,1,1]
317     1:1[27]
l1_1_1.weight   1:4[192,62,3,3]
l1_1_2.weight   1:4[192,62,1,1]
l1_2.weight     1:4[192,57,1,1]
l22.weight      1:4[27,192,1,1]
l22_2.bias      1:1[2187]
l23_v.bias      1:1[256]
l23_v.weight    1:2[256,2187]
l24_v.bias      1:1[1]
l24_v.weight    1:2[1,256]
norm1.bias      1:1[192]
norm1.running_mean      1:1[192]
norm1.running_var       1:1[192]
norm1.weight    1:1[192]

まとめ

OpenCL対応の準備として、ONNXをパースする処理を試した。
パースすることができるようになったので、ONNXファイルからパラメータを取り出して、OpenCLの計算に使用することができる。

ニューラルネットワークの各層の処理は、ONNXファイルのnode一覧の上から順に計算していけばよい。
nodeのOPコードに対応するOpenCLカーネルを呼び出していけばよいが、KataGoのOpenCLカーネルは畳み込みとBNとReLUをまとめていたりするので、汎用的に対応するには何パターンもカーネルを実装しておく必要がある。
現実的には現在のdlshogiで使用するものだけ対応するようになりそうである。

次は、KataGoのOpenCLのソースを流用して簡単なテストをしてみるつもりである。

dlshogi(第2回世界将棋AI電竜戦エキシビジョンバージョン)のWindows版ビルド済みファイル公開

dlshogi(第2回世界将棋AI電竜戦エキシビジョンバージョン)のWindows版ビルド済みファイルを公開します。

モデルファイル

モデルファイルは別のzipファイル(model-dr2_exhi.zip)になっています。
ダウンロード前に、下記のライセンスを参照してください。

ダウンロードしたモデルファイルを使用するにはエンジン設定で、モデルファイルのパスの設定が必要です。
DNN_Modelに解凍したモデルファイル(model-dr2_exhi.onnx)のパスを設定してください。

モデルファイルのニューラルネットワークは15ブロックのResNetになっているため、以前のバージョン(10ブロックResNet)よりNPSが低下します。
GPUの性能によっては、同一の時間での対局で弱くなる場合があります。

倍精度バージョンの追加

1千万ノード以上探索すると、floatの桁落ちによる誤差により評価値の精度が落ちるため、今までのバージョンは長時間の検討には不向きでした。
そのため、検討用に倍精度浮動小数点数でビルドしたバイナリを追加しました。

長時間検討する場合は、ファイル名に「_double」が付く.exeファイルを使用してください。

倍精度バージョンは、メモリ使用量が増えます。
メモリ使用量の目安は、以下の通りです。

バージョン ノード数 メモリ使用量
通常バージョン 1000万ノード 18.2GB
倍精度バージョン 1000万ノード 23.3GB

GPU(RTX 3090)1枚で1千万ノード以上探索するには、5分以上かかります。
長時間検討しない場合や、メモリが少ないマシンでは通常バージョンを使用してください。

なお、1千万ノード以上探索する場合は、UCT_NodeLimitの値を増やす必要があります。

USIオプション追加

読み筋(PV)上に詰みがないかチェックするオプションを追加しました。
頓死をある程度防ぐことができます。
有効にするには、PV_Mate_Search_Threadsを1以上に設定してください。

ただし、CPUコアが余っていないとかえって弱くなる場合があります。

モデル別の探索パラメータの設定

モデルによって最適な探索パラメータが異なるため、探索パラメータのデフォルト値をモデルファイルのパスに拡張子「.ini」を付けたファイルで設定できるようにしました。

USIオプションがデフォルト値から変更されている場合は、変更した値の方が使用されます。
将棋所だと現在の値がデフォルト値かが確認できないため、心配な場合はエンジンの設定を一旦削除して、登録からやり直すとよいです。

.iniファイルには、

オプション名=値

を行区切りで記述します。
※=の前後にスペースは記述しません。

設定例:

C_init=100
C_base=20000
C_fpu_reduction=10
C_init_root=200
C_base_root=21000
Softmax_Temperature=100

モデルファイルのzipには、.iniを含んでいますので、モデルファイル(.onnx)と同一フォルダに配置してください。

モデルファイルのライセンス

モデルファイルを公開してしまうと、大会で不利になってしまうため、モデルファイルだけライセンスを制限します。
使用するには、以下のライセンスに同意が必要です。

1. 本モデルファイルの著作権者は山岡忠夫に帰属します。

2. 私的使用の範囲内で、本モデルファイルを無料で使用できます。

3. コンピュータ将棋の大会向けに本モデルを使用する場合、
    ・モデル学習用の棋譜生成
    ・勝率測定
    を目的とする場合のみ使用を許可します。

    以下は、許可しません。
    ・定跡生成に使用すること
    ・本モデルファイルを直接対局に使用すること
    ・本モデルファイルに追加学習すること
    ・本モデルファイルのパラメータを転用すること

4. リバースエンジニアリング、改変、再配布は禁止します。

ライセンスを確認しているかを担保するため、モデルファイルのzipにはパスワードをかけています。
「dlshogi-dr2_exhi」で解凍できます(※このパスワードは、他には転記しないでください)。

※2022/1/23追記
営利目的(本モデルを使って評価値や候補手を放送して広告収入を得るなど)は許可していませんので、注意喚起します。

最新のモデルファイルは、棋神アナリティクスでご利用いただけます。

電竜戦長時間マッチ「水匠 vs dlshogi」

明日8/15(日) 17:00から、私の開発した「dlshogi」と電竜戦TSEC優勝ソフト「水匠」との長時間マッチを行います。
対局の様子は、YouTubeニコニコ動画で生配信されます。
先手、後手入れ替えて計2局行います。
第1局は、阿部健治郎七段と佐々木勇気七段のダブル解説、そしてゲスト解説に渡辺明名人をお呼びしています。

予告動画


電竜戦長時間マッチ「水匠 vs dlshogi」第1局

youtu.be

電竜戦長時間マッチ「水匠 vs dlshogi」第2局

youtu.be

ソフトの特徴

従来型将棋AI

「水匠」は、第2回電竜戦TSECで優勝した現在最強の将棋AIです。

  • 探索部

やねうら王ライブラリを使用しており、探索はαβ法を基本として、チェスのStockfishを参考に開発されています。
αβ法ベースの探索は、非常に長い歴史があり1940年代(正確にはミニマックス法、αβ法は1950年代)より改良が続けられています。

  • 評価関数

評価関数には、CPUのみで処理できるニューラルネットワーク(NNUE)を使用しています。
入力特徴量に2駒関係を使った4層のニューラルネットワークで、局面単位に低レイテンシで処理ができることが特徴です。

ディープラーニング

一方「dlshogi」は、従来の将棋AIとは異なるディープラーニングを使用した将棋AIになります。
近年急速に強くなっており、第1回電竜戦でGCTが優勝し、dlshogiが第2回電竜戦TSECのB級で優勝しています。

従来のAIと大きく異なる特徴としては、盤面を画像として入力して指し手を予測する点です(正確には精度を上げるため利きなどの情報も入力しています)。
特に将棋のルールを実装することなく、floodgateの上位ソフトの指し手を、いっさい探索を行わず50%以上の精度で予測できます。
ニューラルネットワークには、畳み込みニューラルネットワークという画像認識で使用されるモデルを使用します。
モデルサイズは、15ブロック(30層)のResNetという深いニューラルネットワークを使用します(最新dlshogiの場合。囲碁では40ブロックとかもっと大きい)。
計算量が多いため、GPUを使用することが前提になっています。

ただし、これだけでは終盤の細い読みが必要になる局面に弱いため、モンテカルロ木探索という探索を組み合わせて動作します。
モンテカルロ木探索は、コンピュータ囲碁で開発された方法で比較的新しい方法(2000年代後半)になります。
ディープラーニングモンテカルロ木探索の組み合わせは、AlphaGoがコンピュータ囲碁で用いて、棋力が飛躍的に向上しました。

  • 探索部

Stockfish系の将棋AIでは、静止探索やキラーヒューリスティックといった、将棋において経験的に効果のあるヒューリスティックを用いて探索効率を上げていますが、モンテカルロ木探索は統計的な手法が用いられておりヒューリスティックがありません。
ヒューリスティックに相当する部分は、ディープラーニングによる指し手の予測(ポリシーネットワーク)が担っています。
ヒューリスティックを用いることで探索の効率を上げることができますが、十分に学習されたポリシーネットワークの方が良いパフォーマンスになります。
ヒューリスティックがない分、モンテカルロ木探索の探索部の実装はシンプルです。

dlshogiの強さ

floodgateのレーティングは4500を超えており、最上位です。
第2回電竜戦TSECのB級で優勝時点からは、追加学習を行ってR+40くらいになっています。

現在の精度は、GCTのノートブックで使用しているテスト局面(2008年~2019年のfloodgateのR3500以上の棋譜)を使用した場合、

バージョン 方策(指し手)一致率 価値(勝敗)一致率
GCT電竜 0.46163575 0.73494528
dlshogi with GCT 0.48964297 0.75278556
dlshogi 第2回TSCE時点 0.51964970 0.76499052
dlshogi 最新 0.52322504 0.76564239

と、着実に精度が向上しています。

まとめ

水匠の方も追加学習を行っていると思うので、明日の対局はどっちが勝つかは全く分からないです。
従来型とディープラーニングが、ちょうど同じくらい強さになった絶妙なタイミングなので、どっちが勝つか対局を楽しみにしてください。

内心では勝敗よりも長時間でちゃんと動くかが一番心配です。。。

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であった。

参考として、NVIDIAGPU 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コアを使用するのも用途はありそうである。

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

前回AWS inf1 インスタンスでdlshogiのモデルの推論ができることを確認したが、今回はマルチGPUで推論できるかを確かめてみた。

マルチGPUにするには、Inferentia チップが複数あるAWS inf1 インスタンス(inf1.6xlargeとか)が必要と思っていたが、よく調べるとInferentia チップには、4つのNeuronコアが搭載されていて、4並列で推論ができることがわかった。

inf1.xlargeで、/opt/aws/neuron/bin/neuron-lsを実行すると4コアあることが確かめられる。

$ /opt/aws/neuron/bin/neuron-ls
+--------+--------+--------+-----------+--------------+---------+---------+---------+
| NEURON | NEURON | NEURON | CONNECTED |     PCI      | RUNTIME | RUNTIME | RUNTIME |
| DEVICE | CORES  | MEMORY |  DEVICES  |     BDF      | ADDRESS |   PID   | VERSION |
+--------+--------+--------+-----------+--------------+---------+---------+---------+
| 0      | 4      | 8 GB   |           | 0000:00:1f.0 | NA      | 1062    | NA      |
+--------+--------+--------+-----------+--------------+---------+---------+---------+

マルチコアの使い方

マルチコアの使い方は、単にモデルを複数回loadするだけである。

models = []
for i in range(4):
    model = torch.jit.load('model-pre5_resnet15_swish_b4096lr004-008.neuron.pt')
    model.eval()
    models.append(model)

のようにして、4回ロードすればよい。
それぞれのモデルが別のNeronCoreに割り当てられる。

環境変数NEURONCORE_GROUP_SIZESで、NeuronCoreをグループ化できるが、特にその必要はない。
デフォルトで1コアが1グループになっている。

並列で推論

以下のようなコードで並列で推論を試した。

from concurrent.futures import ThreadPoolExecutor

executor = ThreadPoolExecutor(max_workers=4)

def infer(model):
    for i in range(10000):
        y1, y2 = model(x1, x2)

futures = []
for i in range(4):
    futures.append(executor.submit(infer, models[i]))

for i in range(4):
    futures[i].result()


/opt/aws/neuron/bin/neuron-topで、4つのNeuronコアが使用されていることが確認できた。

$/opt/aws/neuron/bin/neuron-top
neuron-top - 12:34:51
Models: 4 loaded, 4 running. NeuronCores: 4 used.
0000:00:1f.0 Utilizations: NC0 67.13%, NC1 67.41%, NC2 67.32%, NC3 67.30%,
Model ID   Device    NeuronCore%   Device Mem   Host Mem   Model Name
10004      nd0:nc3   67.30           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q
10003      nd0:nc2   67.32           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q
10001      nd0:nc0   67.13           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q
10002      nd0:nc1   67.41           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q

まとめ

AWS inf1 インスタンスでマルチコアでの推論をためした。
一番安いインスタンスでも4コアが搭載されているので、これで1コアの推論性能が良ければNvidiaGPUインスタンよりかなりお得である。
推論性能は別途検証したい。

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

dlshogiをAWS inf1 インスタンスで動かせないか試している。
AWS inf1 インスタンスは、推論に特化したAWS Inferentia チップが搭載されている。

現在のdlshogiは、TensorRTを使用して推論を行っているため、NvidiaGPUで動かすことが前提になっている。
onnxruntime版も用意しているが、推論速度は7.2倍ほど差がある。

dlshogiをAWSで動かそうとすると、NvidiaGPUが搭載されたインスタンスを借りる必要がある。
しかし、NvidiaGPUが搭載されたAWSインスタンスの料金は高い。
AWS inf1 インスタンスは、NvidiaGPUインスタンスに比べるとリーズナブルである。

料金

推論向けのT4が搭載されたg4dnインスタンスと比較すると、バージニア北部で、以下の通りである。

インスタンス vCPU GPU メモリ(GB) スポットインスタンス料金
g4dn.2xlarge 8 1 32 $0.2256 /1 時間
g4dn.12xlarge 48 4 192 $1.1761 /1 時間
inf1.2xlarge 8 1 16 $0.1086 /1 時間
inf1.6xlarge 24 4 48 $0.354 /1 時間

CPU、メモリが同じインスタンスがないが比較的近いスペックだと、同じGPU数で価格は半分以下になる。

推論性能

推論性能は、カタログスペックで、以下の通りである。

T4 65 TFLOP(FP16)
Inferentia 128 TOPS

ただし、dlshogiで実際に性能がでるかは測ってみないとわからない。


ということで、まずはPythonで、dlshogiのモデルの推論ができるか試してみた。

概要

Inferentiaを使用するには、neuron SDKを使用する必要がある。
neuron SDKは、TensorFlowやPyTorchをカスタマイズしたパッケージとして提供される。
Deep Learming AMIを使用すると、フレームワーク別のconda環境が用意されている。
ドキュメント参照:
Compile with Framework API and Deploy on EC2 Inf1 — AWS Neuron documentation

PyTorchで学習したモデルは、直接は利用できないため、Inferentia向けに変換を行う必要がある。
残念ながらモデルの形式は、onnxには対応していない。
onnxに変換前のモデルが必要である(dlshogi with GCTのモデルはonnxでしか公開していないので各自学習したモデルを使用する必要がある)。

なお、onnxからPyTorchに変換するツールを試したが、変換後の推論結果が一致せずうまくいかなかった。

手順

変換

モデルの変換は、inf1インスタンスでなくても実行できる。
c4.xlargeで、Deep Learming AMI(Ubuntu18)を使用して行った。

  • ログイン後、
source activate aws_neuron_pytorch_p36

を実行し、PyTorchのneuron SDKの環境に切り替える。

  • dlshogiのモデルをアップロードする。
  • dlshogiのソースをgit cloneして、pip install -e .でパッケージをインストールする。
  • 以下のスクリプトでモデルを変換する(ネットワークタイプとファイル名は書き換える)。
import torch
import torch_neuron

import torch.nn as nn
from dlshogi.common import *
from dlshogi.network.policy_value_network import policy_value_network
from dlshogi import serializers
from dlshogi import cppshogi

model = policy_value_network('resnet15_swish', add_sigmoid=False)
serializers.load_npz('model-pre5_resnet15_swish_b4096lr004-008', model, False)

class PolicyValueNetworkAddSigmoid(nn.Module):
    def __init__(self, model):
        super(PolicyValueNetworkAddSigmoid, self).__init__()
        self.model = model
        
    def forward(self, x1, x2):
        y1, y2 = self.model(x1, x2)
        return y1, torch.sigmoid(y2)

model = PolicyValueNetworkAddSigmoid(model)
model.eval()

def mini_batch(hcpevec):
    features1 = np.empty((len(hcpevec), FEATURES1_NUM, 9, 9), dtype=np.float32)
    features2 = np.empty((len(hcpevec), FEATURES2_NUM, 9, 9), dtype=np.float32)
    move = np.empty((len(hcpevec)), dtype=np.int64)
    result = np.empty((len(hcpevec)), dtype=np.float32)
    value = np.empty((len(hcpevec)), dtype=np.float32)

    cppshogi.hcpe_decode_with_value(hcpevec, features1, features2, move, result, value)

    z = result.astype(np.float32) - value + 0.5

    return (torch.tensor(features1),
            torch.tensor(features2),
            torch.tensor(move.astype(np.int64)),
            torch.tensor(result.reshape((len(hcpevec), 1))),
            torch.tensor(z),
            torch.tensor(value.reshape((len(value), 1)))
            )

batchsize = 1
hcpevec = np.array([([ 88, 164,  73,  33,  12, 215,  87,  33, 126, 142,  77,  33,  44, 175,  66, 120,  20, 194, 171,  16, 158,  77,  33,  44, 215,  95,  33,  62, 142,  73,  33,  12], 0, 7739, 1, 0)] * batchsize, HuffmanCodedPosAndEval)
x1, x2, t1, t2, z, value = mini_batch(hcpevec)

model_neuron = torch.neuron.trace(model, example_inputs=[x1, x2], dynamic_batch_size=True)

model_neuron.save('model-pre5_resnet15_swish_b4096lr004-008.neuron.pt')
推論

inf1.xlargeをスポットインスタンスで借りて試した。
※スポットインスタンスのvCPUの上限がデフォルト0になっていたので、制限解除の申請が必要だったがすぐに許可された。

  • ログイン後、
source activate aws_neuron_pytorch_p36

を実行し、PyTorchのneuron SDKの環境に切り替える。

  • 変換したモデルをアップロードする。
import torch
import torch_neuron

import torch.nn as nn
from dlshogi.common import *
from dlshogi import cppshogi

model = torch.jit.load('model-pre5_resnet15_swish_b4096lr004-008.neuron.pt')
model.eval()

def mini_batch(hcpevec):
    features1 = np.empty((len(hcpevec), FEATURES1_NUM, 9, 9), dtype=np.float32)
    features2 = np.empty((len(hcpevec), FEATURES2_NUM, 9, 9), dtype=np.float32)
    move = np.empty((len(hcpevec)), dtype=np.int64)
    result = np.empty((len(hcpevec)), dtype=np.float32)
    value = np.empty((len(hcpevec)), dtype=np.float32)

    cppshogi.hcpe_decode_with_value(hcpevec, features1, features2, move, result, value)

    z = result.astype(np.float32) - value + 0.5

    return (torch.tensor(features1),
            torch.tensor(features2),
            torch.tensor(move.astype(np.int64)),
            torch.tensor(result.reshape((len(hcpevec), 1))),
            torch.tensor(z),
            torch.tensor(value.reshape((len(value), 1)))
            )


batchsize = 128
hcpevec = np.array([([ 88, 164,  73,  33,  12, 215,  87,  33, 126, 142,  77,  33,  44, 175,  66, 120,  20, 194, 171,  16, 158,  77,  33,  44, 215,  95,  33,  62, 142,  73,  33,  12], 0, 7739, 1, 0)] * batchsize, HuffmanCodedPosAndEval)
x1, x2, t1, t2, z, value = mini_batch(hcpevec)

print('start')
for i in range(10000):
    y1, y2 = model(x1, x2)
  • 別のコンソールから、/opt/aws/neuron/bin/neuron-topを実行することで、Inferentiaが使用されていることを確認できる。
neuron-top - 12:53:58
Models: 1 loaded, 1 running. NeuronCores: 1 used.
0000:00:1f.0 Utilizations: NC0 66.96%, NC1 0.00%, NC2 0.00%, NC3 0.00%,
Model ID   Device    NeuronCore%   Device Mem   Host Mem   Model Name
10016      nd0:nc0   66.96           29 MB       151 KB    1.5.5.0+3cc38c60b-/tmp/tmpmwpges2q

まとめ

AWS inf1 インスタンスで、dlshogiのモデルの推論ができることが確認できた。
C++から推論するには、LibTorchを使用して推論する必要があるため、dlshogiの推論部分のソース修正が必要になるため、別途試したい。
推論速度もC++で試せるようになったら測定する予定である。