TadaoYamaokaの日記

山岡忠夫Homeで公開しているプログラムの開発ネタを中心に書いていきます。

水匠4とdlshogiのNPSの比較

ディープラーニング系の将棋AIは、従来の将棋AIと比べてNPSが低くても強いという特徴がある。
NPSがどれくらい違うのか質問を受けることがあるので、測定を行った。

NPSのカウント方法の違い

やねうら王(元はStockfishのソース)のNPSは、探索中にdo_moveを行った回数をノード数としてカウントしている
一方、dlshogiは、ノードをニューラルネットワークで評価した回数をノード数としてカウントしている。
そのため、NPSをそのまま比較するのは公平ではない。

そこで、やねうら王のノード数のカウント方法を、qsearchでevalを呼び出したタイミングでカウントするように変更したバージョンでも測定を行った。

測定条件

ハードウェア
CPU AMD EPYC 7742 64コア×2
GPU A100×8
ソフトウェア
ソフトウェア、バージョン 条件
水匠4 やねうら王V6.04、128スレッド、ハッシュ16GB
水匠4 NPS変更 やねうら王V6.04、128スレッド、ハッシュ16GB
dlshogi with GCT(WCSC31) ResNet10ブロック、8GPU、5スレッド/GPU
dlshogi dr2_exhi ResNet15ブロック、8GPU、4スレッド/GPU

※水匠4 NPS変更が、上記のNPSのカウント方法を変更したバージョン
※やねうら王はZEN2に最適化してビルド

思考時間

1秒/局面

測定結果

NPS
水匠4 水匠4 NPS変更 dlshogi with GCT dlshogi dr2_exhi
平均 99,890,518 23,917,691 320,066 295,643
中央値 100,784,164 23,112,924 320,377 300,384
最小値 66,594,384 15,496,863 218,794 219,697
最大値 121,220,480 40,340,437 389,231 344,105
比率
水匠4 NPS変更/水匠4 dlshogi with GCT/水匠4 dlshogi dr2_exhi/水匠4 dlshogi with GCT/水匠4 NPS変更 dlshogi dr2_exhi/水匠4 NPS変更
平均 23.94%(1/4.2) 0.32%(1/312) 0.30%(1/338) 1.34%(1/75) 1.24%(1/81)
中央値 22.93%(1/4.4) 0.32%(1/315) 0.30%(1/336) 1.39%(1/72) 1.30%(1/77)
最小値 23.27%(1/4.3) 0.33%(1/304) 0.33%(1/303) 1.41%(1/71) 1.42%(1/71)
最大値 33.28%(1/3.0) 0.32%(1/311) 0.28%(1/352) 0.96%(1/104) 0.85%(1/117)

考察

やねうら王のNPSのカウント方法をevalを呼び出した回数にすると、NPSは平均で約1/4.2になる。

dlshogiとNPSカウント方法変更前の水匠4と比較すると、
dshogi with GCT(ResNet10ブロック)は、NPSは平均で約1/312である。
dshogi dr2_exhi(ResNet15ブロック)は、NPSは平均で約1/338である。
なお、A100の場合、10ブロックと15ブロックでそれほどNPSが違わないが、V100だと大きな差がでる。

NPSカウント方法変更後の水匠4と比較すると、
dshogi with GCT(ResNet10ブロック)は、NPSは平均で約1/75である。
dshogi dr2_exhi(ResNet15ブロック)は、NPSは平均で約1/81である。

まとめ

従来の将棋AIとディープラーニング系で、NPSにどれだけ違いがあるかを調査した。
NPSのカウント方法を公平にして比較すると、平均で10ブロックで1/75、15ブロックで1/81であることがわかった。
やねうら王を使用して通常表示されるNPSでは、平均で10ブロックで1/312、15ブロックで1/338であることがわかった。

ということで、NPSはどれくらい違うのかという質問を受けたら、「公平な条件で比較すると約1/80(もしくは約1.3%)です」と答えることにする。

なお、測定環境のCPUとGPUにも依存するため、今回の測定条件での参考値である。

将棋AIモデルのテストデータの作り方

dlshogiで以前から使用しているテストデータと、GCTのノートブックのテストデータで、テスト損失と正解率の違いがあるため、テストデータの作成方法によってどう違いがでるのかを検証した。

テストデータの比較

まず、dlshogiのテストデータとGCTのテストデータの作成条件と、それぞれで同一モデルで測定したテスト損失と正解率について整理すると以下の通りになる。

dlshogiのテストデータ

2018/6/17に作成したテストデータを使い続けている。
これはfloodgateの2017年から2018年6月までの棋譜を元に、csa_to_hcpe.pyで、レーティング双方3500以上、評価値3000を超えたら打ち切りという条件で作成している。
千日手、最大手数による引き分けは出力していない。
また、hcpe_uniqで、局面と指し手が重複するデータを除外している。
局面数は、856,923局面ある。

モデルごとの精度は以下の通り。

モデル 方策損失 価値損失 方策正解率 価値正解率 方策エントロピー
GCT電竜 1.78864127 0.53541758 0.46435664 0.72764303 1.26309002
dlshogi with GCT(model-0000225kai) 1.77482759 0.49535870 0.49073551 0.74466278 1.02956448
dlshogi dr2_exhi 1.41893253 0.46856942 0.52342348 0.76075737 1.35510777
GCTのノートブックのテストデータ

使用しているツール(csa_to_hcpe.py、hcpe_uniq)は同じものを使用している。
使用している棋譜は、floodgateの2008年から2019年の棋譜で、レーティングは双方3500以上、評価値は5000を超えたら打ち切りという条件で作成している。
局面数は、2,414,974局面ある。

上記dlshogiのテストデータとの差分は、floodgateの使用している期間と、評価値の閾値である。

モデルごとの精度は以下の通り。

モデル 方策損失 価値損失 方策正解率 価値正解率 方策エントロピー
GCT電竜 1.83690576 0.51293612 0.46163161 0.73495564 1.24398525
dlshogi with GCT(model-0000225kai) 1.82079209 0.47346048 0.48964214 0.75278598 1.00554464
dlshogi dr2_exhi 1.42115222 0.44965372 0.52322546 0.76564239 1.36627845
比較

上記の2つのテストデータの精度を比較すると、方策についてはdlshogiのテストデータが良い値になり、価値についてはGCTのテストデータが良い値になっている。
floodgateの期間の違いはあるが、主な差分は評価値の閾値の差である。
終盤の局面(評価値3000以上5000以下)まで含む場合、方策の精度は低くなり、価値の精度は高くなるという傾向があることが分かった。
終盤ほど読みの比重が高くなるため方策の正解率が下がって、評価値の差が広がるほど勝敗が明確なため価値の正解率が上がるということだと理解できる。

重複局面の条件変更

np.unique

上記のテストデータは、重複局面を除外する条件として、局面と指し手が一致していることを条件としている。
これを、評価値と勝敗まで含めて一致という条件にしてみた(hcpeを単純にnp.uniqueでユニークにした)。
元データは、floodgateの2017年から2019年の棋譜で、レーティングは双方3500以上、評価値は5000を超えたら打ち切りという条件で作成した。

局面数は、重複削除前:2798640、重複削除後:2551507となった。

モデルごとの精度は以下の通り。

モデル 方策損失 価値損失 方策正解率 価値正解率 方策エントロピー
GCT電竜 1.81846763 0.51962999 0.46602104 0.72777115 1.25068274
dlshogi with GCT(model-0000225kai) 1.80058509 0.48026565 0.49412886 0.74460519 0.99275402
dlshogi dr2_exhi 1.40826132 0.45623083 0.52678000 0.75692493 1.36271331

方策が高くなり、価値は低い値になった。
評価値まで一致しないと別のデータになるため、序盤の局面の数が多くなることが原因と考えられる。
同一局面で評価値がわずかに違うだけのデータが増えると、データが偏るため、局面と指し手の一致を重複の条件とする方が適切と考える。

結果の最頻値化と評価値平均化

局面と指し手の一致で除外すると、最初に読み込んだデータの勝敗結果が使用されて、残りのデータの結果が無視されてしまう。
これは適切ではないため、結果の最頻値を採用し、評価値については平均化を行うようにした。

実装は、pandasのgroupbyを使用して行った。

    df = pd.concat([pd.DataFrame(hcpes['hcp']), pd.DataFrame(hcpes[['eval', 'bestMove16', 'gameResult', 'dummy']])], axis=1)

    # hcpとbestMove16でグループ化して平均を算出
    df2 = df.groupby(list(range(32)) + ['bestMove16'], as_index=False).mean()

    # gameResultは最頻値に変換
    df2.loc[df2['gameResult'] >= 1.5, 'gameResult'] = 2
    df2.loc[(df2['gameResult'] > 0) & (df2['gameResult'] < 2), 'gameResult'] = 1

    hcpes2 = np.zeros(len(df2), HuffmanCodedPosAndEval)
    hcpes2['hcp'] = df2[list(range(32))]
    hcpes2['eval'] = df2['eval']
    hcpes2['bestMove16'] = df2['bestMove16']
    hcpes2['gameResult'] = df2['gameResult']

元データは先ほどと同じく、floodgateの2017年から2019年の棋譜で、レーティングは双方3500以上、評価値は5000を超えたら打ち切りという条件で作成した。
局面数は、重複削除前:2798640、重複削除後:2383564となった。

モデルごとの精度は以下の通り。

モデル 方策損失 価値損失 方策正解率 価値正解率 方策エントロピー
GCT電竜 1.84447462 0.50861321 0.46116580 0.73828041 1.24337653
dlshogi with GCT(model-0000225kai) 1.82565162 0.46669739 0.48931699 0.75633276 1.00591798
dlshogi dr2_exhi 1.42368041 0.44125965 0.52300734 0.76926015 1.37016454

GCTのノートブックのテストデータに比較的近い値になっている。
重複のない局面の方が多いため傾向としては同じで、重複のある局面での勝敗がより統計的に正しく測れるようになったと考える。

さらに統計的に測りたければ、勝敗結果を分布するのがよいが、その場合、正解率という指標は使えなくなるので、正解率を使うなら最頻値を使うのが良いと考える。

同一局面で指し手ごとにデータを分けるのも、数が多い指し手と少ない指し手を同様に扱っているので良くないが、指し手の方は最頻値を採用すると、序盤は一手のみが正解とも言えないので、指し手ごとにデータを分けるようにしている。

まとめ

テストデータの作成条件による精度の傾向を調べた。
また、重複削除の方法として、局面の指し手が一致するデータの勝敗結果の最頻値を使用するという方法を試した。

棋譜から条件を指定してhcpeを作成する方法と、最頻値を使う重複削除はツール化しているので活用して欲しい(今回の調査中にcsa_to_hcpe.pyにバグが見つかったので修正している)。

dlshogiのOpenCL対応 その1

dlshogiをOpenCLに対応させたいと思っている。
現在のdlhsogiはTensorRTを使用しているため、再配布の制限があり、環境構築が大変になっている。
OpenCLは、必要なdllを同梱すれば環境構築が不要になるので利用者にメリットがある。

OpenCLのコードを一から作成するのは難易度が高すぎるので、KataGoで、OpenCLでTensorCoreを使う実装がされているので、ソースを流用して対応しようと思っている。

OpenCLのチューニング処理

今回は、KataGoのソースの内、OpenCLのチューニングを行う部分だけ抜き出して試してみた。
OpenCL 1.2は、OpenCL C言語のソースを文字列でライブラリに渡す仕組み(オンラインコンパイル)で、KataGoではOpenCLのソースで使用するパラメータを初回にチューニングして、ソースにパラメータを埋め込んで編集するようにしている。
行列演算(gemm)と、畳み込み演算で使うパラメータがチューニング対象になっている。

ソース修正

KataGoのチューニングの処理は、KataGoのモデルに特化したテストを行っているため、そのままではdlshogiのモデルでは使用できない箇所があった。
具体的には、盤のサイズが19x19でなくて、9x9になる。
また、KataGoは入力層が5x5の畳み込みになっているので、conv5x5のテストがあるが、dlshogiでは不要である。
また、グローバルプーリングに特化した処理もあったがdlshogiでは不要である。
畳み込みの入力チャネルサイズと出力チャネルサイズのパターンもdlshogiに合わせた変更が必要であった。

依存しているソースが多かったので諸々スリムにして、最小限のソースで試せるようにした。
opencltune/tune at master · TadaoYamaoka/opencltune · GitHub

実行結果

チューニングを実行すると以下のように表示される。

Found OpenCL Platform 0: NVIDIA CUDA (NVIDIA Corporation) (OpenCL 1.2 CUDA 11.2.152)
Found 1 device(s) on platform 0 with type CPU or GPU or Accelerator
Found OpenCL Device 0: GeForce RTX 3090 (NVIDIA Corporation) (score 11000102)
Creating context for OpenCL Platform: NVIDIA CUDA (NVIDIA Corporation) (OpenCL 1.2 CUDA 11.2.152)
Using OpenCL Device 0: GeForce RTX 3090 (NVIDIA Corporation) OpenCL 1.2 CUDA (Extensions: cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_fp64 cl_khr_byte_addressable_store cl_khr_icd cl_khr_gl_sharing cl_nv_compiler_options cl_nv_device_attribute_query cl_nv_pragma_unroll cl_nv_d3d10_sharing cl_khr_d3d10_sharing cl_nv_d3d11_sharing cl_nv_copy_opts cl_nv_create_buffer cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_device_uuid)
Beginning GPU tuning for GeForce RTX 3090 channels 224
Setting winograd3x3TileSize = 4
------------------------------------------------------
Tuning xGemmDirect for 1x1 convolutions and matrix mult
Testing 56 different configs
Tuning 0/56 (reference) Calls/sec 16094.4 L2Error 0 WGD=8 MDIMCD=1 NDIMCD=1 MDIMAD=1 NDIMBD=1 KWID=1 VWMD=1 VWND=1 PADA=1 PADB=1
Tuning 1/56 Calls/sec 16281.7 L2Error 0 WGD=8 MDIMCD=1 NDIMCD=1 MDIMAD=1 NDIMBD=1 KWID=1 VWMD=1 VWND=1 PADA=1 PADB=1
Tuning 2/56 Calls/sec 86246.6 L2Error 0 WGD=8 MDIMCD=8 NDIMCD=8 MDIMAD=8 NDIMBD=8 KWID=1 VWMD=1 VWND=1 PADA=1 PADB=1
Tuning 3/56 Calls/sec 112680 L2Error 0 WGD=16 MDIMCD=8 NDIMCD=8 MDIMAD=8 NDIMBD=8 KWID=1 VWMD=1 VWND=1 PADA=1 PADB=1
Tuning 20/56 ...
Tuning 40/56 ...
------------------------------------------------------
Tuning xGemm for convolutions
Testing 70 different configs
Tuning 0/70 (reference) Calls/sec 2668.78 L2Error 0 MWG=8 NWG=8 KWG=8 MDIMC=1 NDIMC=1 MDIMA=1 NDIMB=1 KWI=1 VWM=1 VWN=1 STRM=0 STRN=0 SA=0 SB=0
Tuning 1/70 Calls/sec 2675.99 L2Error 0 MWG=8 NWG=8 KWG=8 MDIMC=1 NDIMC=1 MDIMA=1 NDIMB=1 KWI=1 VWM=1 VWN=1 STRM=0 STRN=0 SA=0 SB=0
Tuning 2/70 Calls/sec 13413.4 L2Error 0 MWG=8 NWG=8 KWG=8 MDIMC=8 NDIMC=8 MDIMA=8 NDIMB=8 KWI=1 VWM=1 VWN=1 STRM=0 STRN=0 SA=0 SB=0
Tuning 3/70 Calls/sec 20150.1 L2Error 0 MWG=16 NWG=16 KWG=16 MDIMC=8 NDIMC=8 MDIMA=8 NDIMB=8 KWI=1 VWM=1 VWN=1 STRM=0 STRN=0 SA=0 SB=0
Tuning 4/70 Calls/sec 26924.6 L2Error 0 MWG=16 NWG=16 KWG=16 MDIMC=8 NDIMC=8 MDIMA=8 NDIMB=8 KWI=2 VWM=2 VWN=2 STRM=0 STRN=0 SA=0 SB=0
Tuning 9/70 Calls/sec 27470.9 L2Error 0 MWG=16 NWG=16 KWG=32 MDIMC=8 NDIMC=8 MDIMA=8 NDIMB=8 KWI=2 VWM=2 VWN=2 STRM=0 STRN=0 SA=0 SB=0
Tuning 17/70 Calls/sec 29074.5 L2Error 0 MWG=32 NWG=32 KWG=32 MDIMC=16 NDIMC=8 MDIMA=16 NDIMB=8 KWI=2 VWM=2 VWN=2 STRM=0 STRN=0 SA=1 SB=1
Tuning 22/70 Calls/sec 36589.1 L2Error 0 MWG=64 NWG=64 KWG=32 MDIMC=16 NDIMC=8 MDIMA=16 NDIMB=8 KWI=2 VWM=4 VWN=4 STRM=0 STRN=0 SA=1 SB=1
Tuning 24/70 Calls/sec 39082.6 L2Error 0 MWG=64 NWG=64 KWG=32 MDIMC=16 NDIMC=16 MDIMA=16 NDIMB=16 KWI=2 VWM=4 VWN=4 STRM=0 STRN=0 SA=1 SB=1
Tuning 31/70 Calls/sec 39529.5 L2Error 0 MWG=32 NWG=32 KWG=32 MDIMC=8 NDIMC=8 MDIMA=8 NDIMB=8 KWI=2 VWM=4 VWN=4 STRM=0 STRN=0 SA=1 SB=1
Tuning 44/70 Calls/sec 41162 L2Error 0 MWG=32 NWG=32 KWG=16 MDIMC=8 NDIMC=8 MDIMA=8 NDIMB=8 KWI=2 VWM=4 VWN=4 STRM=0 STRN=0 SA=1 SB=1
Tuning 60/70 ...
------------------------------------------------------
Tuning hGemmWmma for convolutions
Testing 146 different configs
Tuning 0/146 (reference) Calls/sec 71175 L2Error 0 MWG=16 NWG=16 KWG=16 MWAVE=16 NWAVE=16 MWARP=16 NWARP=16 VWM=2 VWN=2 SA=0 SB=0
Tuning 1/146 Calls/sec 77126.9 L2Error 0 MWG=16 NWG=16 KWG=16 MWAVE=16 NWAVE=16 MWARP=16 NWARP=16 VWM=2 VWN=2 SA=0 SB=0
Tuning 20/146 ...
Tuning 40/146 ...
Tuning 60/146 ...
Tuning 64/146 Calls/sec 87848.6 L2Error 0 MWG=32 NWG=32 KWG=16 MWAVE=32 NWAVE=8 MWARP=32 NWARP=8 VWM=2 VWN=2 SA=0 SB=0
Tuning 80/146 ...
Tuning 96/146 Calls/sec 88581 L2Error 0 MWG=32 NWG=32 KWG=32 MWAVE=16 NWAVE=16 MWARP=16 NWARP=16 VWM=2 VWN=2 SA=0 SB=0
Tuning 120/146 ...
Tuning 140/146 ...
Enabling FP16 tensor cores due to better performance
------------------------------------------------------
Using FP16 storage!
Using FP32 compute!
Using FP16 tensor cores!
------------------------------------------------------
Tuning winograd transform for convolutions
Testing 47 different configs
Tuning 0/47 (reference) Calls/sec 34136.7 L2Error 0  transLocalSize0=1 transLocalSize1=1
Tuning 1/47 Calls/sec 34885.4 L2Error 0  transLocalSize0=1 transLocalSize1=1
Tuning 2/47 Calls/sec 184904 L2Error 0  transLocalSize0=128 transLocalSize1=2
Tuning 8/47 Calls/sec 217745 L2Error 0  transLocalSize0=16 transLocalSize1=2
Tuning 9/47 Calls/sec 230045 L2Error 0  transLocalSize0=64 transLocalSize1=1
Tuning 20/47 ...
Tuning 40/47 ...
------------------------------------------------------
Tuning winograd untransform for convolutions
Testing 111 different configs
Tuning 0/111 (reference) Calls/sec 58050.2 L2Error 0  untransLocalSize0=1 untransLocalSize1=1 untransLocalSize2=1
Tuning 1/111 Calls/sec 58849.8 L2Error 0  untransLocalSize0=1 untransLocalSize1=1 untransLocalSize2=1
Tuning 2/111 Calls/sec 173990 L2Error 0  untransLocalSize0=8 untransLocalSize1=1 untransLocalSize2=8
Tuning 4/111 Calls/sec 210017 L2Error 0  untransLocalSize0=8 untransLocalSize1=2 untransLocalSize2=4
Tuning 20/111 ...
Tuning 32/111 Calls/sec 211785 L2Error 0  untransLocalSize0=8 untransLocalSize1=4 untransLocalSize2=4
Tuning 40/111 Calls/sec 236954 L2Error 0  untransLocalSize0=8 untransLocalSize1=4 untransLocalSize2=2
Tuning 60/111 ...
Tuning 80/111 ...
Tuning 100/111 ...
Done tuning
------------------------------------------------------

パラメータの組み合わせのテストを実行してスループットが計測されている。
TensorCoreも使用できている。

チューニングを実行した結果、以下のようにパラメータが保存された。

VERSION=8
#shouldUseFP16Storage
1
#shouldUseFP16Compute
0
#shouldUseFP16TensorCores
1
#xGemmDirect
WGD=16 MDIMCD=8 NDIMCD=8 MDIMAD=8 NDIMBD=8 KWID=1 VWMD=1 VWND=1 PADA=1 PADB=1
#xGemm
MWG=32 NWG=32 KWG=16 MDIMC=8 NDIMC=8 MDIMA=8 NDIMB=8 KWI=2 VWM=4 VWN=4 STRM=0 STRN=0 SA=1 SB=1
#xGemm16
MWG=32 NWG=32 KWG=16 MDIMC=8 NDIMC=8 MDIMA=8 NDIMB=8 KWI=2 VWM=4 VWN=4 STRM=0 STRN=0 SA=1 SB=1
#hGemmWmma
MWG=32 NWG=32 KWG=32 MWAVE=16 NWAVE=16 MWARP=16 NWARP=16 VWM=2 VWN=2 SA=0 SB=0
#conv3x3
INTILE_XSIZE=6 INTILE_YSIZE=6 OUTTILE_XSIZE=4 OUTTILE_YSIZE=4 transLocalSize0=64 transLocalSize1=1 untransLocalSize0=8 untransLocalSize1=4 untransLocalSize2=2

KataGoを初回起動して実行されたチューニングで保存された結果は、以下の通りになる。

VERSION=8
#shouldUseFP16Storage
1
#shouldUseFP16Compute
0
#shouldUseFP16TensorCores
1
#xGemmDirect
WGD=32 MDIMCD=8 NDIMCD=8 MDIMAD=16 NDIMBD=16 KWID=8 VWMD=2 VWND=2 PADA=1 PADB=1
#xGemm
MWG=64 NWG=64 KWG=32 MDIMC=16 NDIMC=8 MDIMA=16 NDIMB=8 KWI=2 VWM=4 VWN=4 STRM=0 STRN=0 SA=1 SB=1
#xGemm16
MWG=64 NWG=64 KWG=32 MDIMC=16 NDIMC=8 MDIMA=16 NDIMB=8 KWI=2 VWM=4 VWN=4 STRM=0 STRN=0 SA=1 SB=1
#hGemmWmma
MWG=64 NWG=64 KWG=64 MWAVE=16 NWAVE=16 MWARP=16 NWARP=16 VWM=2 VWN=2 SA=0 SB=0
#conv3x3
INTILE_XSIZE=6 INTILE_YSIZE=6 OUTTILE_XSIZE=4 OUTTILE_YSIZE=4 transLocalSize0=128 transLocalSize1=1 untransLocalSize0=8 untransLocalSize1=4 untransLocalSize2=1
#conv5x5
INTILE_XSIZE=6 INTILE_YSIZE=6 OUTTILE_XSIZE=2 OUTTILE_YSIZE=2 transLocalSize0=128 transLocalSize1=1 untransLocalSize0=8 untransLocalSize1=4 untransLocalSize2=1
#gPool
XYSTRIDE=16 CHANNELSTRIDE=2 BATCHSTRIDE=1

ベストなパラメータがdlshogi向けと比べて変わっているのが確認できる。

まとめ

OpenCL対応のとっかかりとして、KataGoのチューニング処理を抜き出して試してみた。
これで、行列演算と畳み込み演算という基本部品は動かせるようになった。
次はdlshogiモデルの推論処理を試す予定だが、KataGoとdlshogiのモデルの違いから、OpenCL C言語のソースを修正する必要がある。
OpenCLのコードの理解して修正する必要があり、まだ先は長そうである。

やねうら王バグのdlshogiへの影響

先日の水匠とdlshogiの長時間マッチの第1局で、水匠側にバグが発生するというハプニングがあった。
原因などの詳細は、やねうら王ブログを参照して欲しい。
先日の電竜戦、長時間マッチで現れたやねうら王のバグについて | やねうら王 公式サイト

ここで、このやねうら王のバグは、dlshogiがバグが起きやすい局面に誘導したのではないかという疑問が浮かんでくる。
やねうら王ブログでも可能性について言及している。

はじめ私もその可能性はあると考えた。
しかし、再現性が低いことがわかったため可能性は低そうだが、検証してみないと否定はできない状況である。

dlshogiは、強化学習をリーグ戦で実施しており、割合は高くないが水匠も混ぜている。
バグの影響があったとするとモデルの精度にも悪影響していることになり、はっきりさせておきたい。


やねうら王に修正のプルリクエストが上がったので、修正前後で勝率を比較し、バグの影響がどれくらいあったか検証した。

勝率比較

平手初期局面から

平手初期局面から持ち時間5分2秒加算の結果は以下の通り。

_bugfixがPRを適用したもの。
先後は1局ずつ入れ替えて対局。
dlshogiは、1GPU、3スレッド。
やねうら王は、12スレッド、ハッシュは4096。

リーグ戦のレーティング
# PLAYER              :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 dlshogi-1gpu_3th    :    87.3   28.2   167.5     248    68     100  151   33   64    13
2 suisho4-12th        :   -12.7   27.1   115.5     245    47      99   87   57  101    23
3 suisho4-12th_bugfix :   -74.6   28.5    84.0     241    35     ---   60   48  133    20

White advantage = 37.14 +/- 17.36
Draw rate (equal opponents) = 20.20 % +/- 2.15

1対1のレーティング

# PLAYER              :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 suisho4-12th        :    33.1   27.0    69.5     118    59      99   52   35   31    30
2 suisho4-12th_bugfix :   -33.1   27.0    48.5     118    41     ---   31   35   52    30

White advantage = 78.07 +/- 28.27
Draw rate (equal opponents) = 31.47 % +/- 4.53

# PLAYER              :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 dlshogi-1gpu_3th    :    51.3   29.6    79.5     124    64     100   69   21   34    17
2 suisho4-12th        :   -51.3   29.6    44.5     124    36     ---   34   21   69    17

White advantage = -23.29 +/- 30.44
Draw rate (equal opponents) = 17.86 % +/- 3.51

# PLAYER              :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 dlshogi-1gpu_3th    :    79.7   34.5    85.0     120    71     100   79   12   29    10
2 suisho4-12th_bugfix :   -79.7   34.5    35.0     120    29     ---   29   12   79    10

White advantage = 57.38 +/- 34.96
Draw rate (equal opponents) = 11.31 % +/- 3.23

※White advantageは先手のレーティング

バグ修正前に対して特に勝率が高くなっているということはなかった。
バグ修正後に対しての方がむしろ勝率が高くなっているが、平手初期局面からだと結果が偏るため、バグ修正後が弱くなっているとは言えない。

たややん互角局面24手目から

たややん互角局面24手目から持ち時間5分2秒加算の結果は以下の通り。

リーグ戦のレーティング
# PLAYER              :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 dlshogi-1gpu_3th    :    79.6   24.6   196.0     296    66     100  170   52   74    18
2 suisho4-12th        :   -32.9   25.0   127.0     294    43      73   93   68  133    23
3 suisho4-12th_bugfix :   -46.7   24.7   118.0     292    40     ---   84   68  140    23

White advantage = 43.31 +/- 15.26
Draw rate (equal opponents) = 22.59 % +/- 2.00

1対1のレーティング

# PLAYER              :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 suisho4-12th        :     3.7   23.5    73.0     143    51      62   52   42   49    29
2 suisho4-12th_bugfix :    -3.7   23.5    70.0     143    49     ---   49   42   52    29

White advantage = 2.40 +/- 25.59
Draw rate (equal opponents) = 29.38 % +/- 3.84

# PLAYER              :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 dlshogi-1gpu_3th    :    52.2   26.5    95.0     148    64     100   82   26   40    18
2 suisho4-12th        :   -52.2   26.5    53.0     148    36     ---   40   26   82    18

White advantage = 51.98 +/- 27.50
Draw rate (equal opponents) = 18.75 % +/- 3.27

# PLAYER              :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 dlshogi-1gpu_3th    :    72.0   28.4   100.0     146    68     100   87   26   33    18
2 suisho4-12th_bugfix :   -72.0   28.4    46.0     146    32     ---   33   26   87    18

White advantage = 85.79 +/- 30.30
Draw rate (equal opponents) = 20.28 % +/- 3.60

互角局面でも、バグ修正前に対して特に勝率が高くなっているということはなかった。
やねうら王のバグ修正前後の強さは誤差の範囲である。

強化学習への影響

強化学習でのリーグ戦での、修正前後のやねうら王(水匠4)に対する勝率も計測した。

修正前
[2021-08-19 13:21:29.823] [info] Made 25000061 teacher nodes in 213932 seconds. games:259910, draws:18227, ply/game:96.1874, usi_games:18224, usi_win:9216, usi_draw:656, usi_winrate:52.46%
修正後
[2021-08-22 21:03:38.594] [info] Made 25000138 teacher nodes in 218892 seconds. games:259670, draws:18612, ply/game:96.2766, usi_games:18216, usi_win:9212, usi_draw:720, usi_winrate:52.65%

usi_winrateがやねうら王(水匠4)に対する勝率を示す。

修正前後で、勝率は約18200対局で52.46%と52.65%であり、誤差の範囲である。

そもそも、強化学習は、平手初期局面からではなく、初期局面集からさらにランダムムーブを行った局面から開始しているため、今回起きた局面になる可能性はかなり低い。
今回計測した勝率からも、バグの局面を学習してしまっているということはないと言えそうだ。

まとめ

dlshogiがやねうら王のバグを学習してしまっていることがないか確認した。
モデルに悪影響がでていないか心配だったのと、ちゃんと計測して事実を書いておかないと誤解が起きると考えたので、今回計測を行った。
バグ修正前後の勝率と、強化学習での勝率の計測結果から、影響はないと言えそうである。

dlshogiの先手勝率

dlshogiの先手勝率について質問されることがあるので、計測してみた。

持ち時間1分1秒加算で、dlshogi(dr2_exhi)の自己対局と、dlshogiと水匠4の対局で計測した。
対局条件が持ち時間だと、ある程度手が揺らぐため完全に同じ手順になるのは防ぐことができる。

測定条件

  • 持ち時間1分1秒加算
  • dlshogiは、1GPU(V100)、2スレッド
  • 水匠4は、10スレッド、ハッシュサイズ4096

やねうら王は読み筋の成れない飛車が成るバグフィックスのプルリクバージョンを使用した。

dlshogiを2エンジンと水匠4を1エンジンで、リーグ戦で測定した。
リーグ戦では、1局ずつ先手後手を入れ変えて対局している。

測定結果

リーグ戦のレーティング
# PLAYER             :  RATING  ERROR  POINTS  PLAYED   (%)  CFS(%)    W    D    L  D(%)
1 dlshogi-1gpu_2th-2 :    59.6   16.9   506.5     826    61      91  497   19  310     2
2 dlshogi-1gpu_2th-1 :    40.4   16.4   477.5     828    58     100  473    9  346     1
3 suisho4-10th       :  -100.0   17.5   254.0     822    31     ---  241   26  555     3
エンジンごとの先手、後手の勝率

同一手順の対局を除いた、エンジンごとの先手、後手の勝率は以下の通り。

dlshogi-1gpu_2th-1 vs dlshogi-1gpu_2th-2: 175-190-1 (48.0%)
Black vs White: 240-125-1 (65.7%)
dlshogi-1gpu_2th-1 playing Black: 117-67-0 (63.6%)
dlshogi-1gpu_2th-1 playing White: 58-123-1 (32.1%)
dlshogi-1gpu_2th-2 playing Black: 123-58-1 (67.9%)
dlshogi-1gpu_2th-2 playing White: 67-117-0 (36.4%)

dlshogi-1gpu_2th-1 vs suisho4-10th: 275-127-8 (68.0%)
Black vs White: 236-166-8 (58.5%)
dlshogi-1gpu_2th-1 playing Black: 155-46-5 (76.5%)
dlshogi-1gpu_2th-1 playing White: 120-81-3 (59.6%)
suisho4-10th playing Black: 81-120-3 (40.4%)
suisho4-10th playing White: 46-155-5 (23.5%)

dlshogi-1gpu_2th-2 vs suisho4-10th: 280-112-18 (70.5%)
Black vs White: 223-169-18 (56.6%)
dlshogi-1gpu_2th-2 playing Black: 152-41-12 (77.1%)
dlshogi-1gpu_2th-2 playing White: 128-71-6 (63.9%)
suisho4-10th playing Black: 71-128-6 (36.1%)
suisho4-10th playing White: 41-152-12 (22.9%)

dlshogi同士だと、先手勝率は、65.7%とかなり高い値になっている。

dlshogiと水匠4では、先手勝率は、57.55%※となっている。
※「dlshogi-1gpu_2th-1 vs suisho4-10th」と「dlshogi-1gpu_2th-2 vs suisho4-10th」の平均

まとめ

dlshogiの平手初期局面からの先手勝率を測定した。

dlshogi同士の対局では、65.7%と予想以上に先手の勝率が高かった。
ニューラルネットワークには、後手の場合は盤面を180度回転して入力しているため、先手後手の区別は特にしていない。
それが、偏った先手勝率の原因になっている可能性がある。
本来ある盤面が与えられたら、最善手は手番がどちらであるかによらないはず(入玉宣言を除く)なので、手番を区別する必要はないと考えていたが、後手特有の戦略を学習するように、手番を区別して学習した方がよいのかもしれない。

dlshogiと水匠4では、先手勝率は57.55%と高い値になったが、他のソフトでもこれくらいの値になるという話は聞いているので、それほど偏った値ではないと思う。

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 master · 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のソースを流用して簡単なテストをしてみるつもりである。