将棋AIをChainerを使用した実装からcuDNNを使用した実装に変更できたので、マルチGPUでの性能を測定した。
Python経由でChainerを使用しているPythonのGIL機構によってマルチスレッドの性能に制限がある。
Chainerを使用した場合の、マルチGPUによる効果は1.33倍程度であった。
| playout/sec | |
| シングルGPU(Titan V) | 5724 |
| マルチGPU(Titan V+GeForce 1080 Ti) | 7640 |
cuDNNを直接使用した場合の測定結果は、以下の通りであった。
| playout/sec | |
| シングルGPU(Titan V) | 6266 |
| マルチGPU(Titan V+GeForce 1080 Ti) | 10779 |
cuDNNを直接使用した場合、マルチGPU化により1.72倍の高速化の効果があった。
ChainerをシングルGPUで使用した場合と比べると、1.88倍になった。
モデルの学習にディープラーニングフレームワークのChainerを使用していますが、対局時にChainerで推論を行うと、Python経由で呼び出すためマルチGPUで動かす際、Python経由だとGILによってマルチスレッドの性能が出なくなる。
また、実行環境にPythonが必要になるため、実行バイナリを配布しても利用者側で環境構築が必要になってしまう。
それらの問題を解決するため、対局時にはChainerを使用しないで、cuDNNを直接しようして推論を行えるようにした。
実装方法の詳細は以下の記事を参照してほしい。
Chainerで学習したモデルを使ってcuDNNで推論する - TadaoYamaokaの開発日記
Chainerで学習したモデルを使ってcuDNNで推論する(時間計測) - TadaoYamaokaの開発日記
Chainerで学習したモデルを使ってcuDNNで推論する(BatchNormalization) - TadaoYamaokaの開発日記
Chainerで学習したモデルを使ってcuDNNで推論する(dropout) - TadaoYamaokaの開発日記
Chainerで学習したモデルを使ってcuDNNで推論する(ResNet) - TadaoYamaokaの開発日記
Chainerで学習したモデルを使ってcuDNNで推論する(マルチGPU) - TadaoYamaokaの開発日記
はじめ、キューに溜まっている要求がバッチサイズ未満でも固定のバッチサイズで推論を行うように実装したら、Python経由でChainerを呼び出すよりも遅くなってしまった。
| playout/sec | |
| Chainer | 4542 |
| cuDNN | 4095 |
期待した結果にならなかったので、バッチサイズを可変にして、キューに溜まっているサイズのバッチサイズで推論を行うようにしたところ、Python経由でChainerを使用する場合を上回った。
| playout/sec | |
| Chainer | 4542 |
| cuDNN | 4854 |
GPUの並列性を期待して、バッチサイズによらず速度は一定と考えたが、メモリの転送時間とかの影響があるため、バッチサイズは可変にした方が良いことがわかった。
しかし、MNISTの推論で比較した場合に比べて、期待したほど速度が向上していない。
ChainerはPythonで実行している分無駄があるので、もっと差が出てよさそうである。
Chainerは、速度を出すためにチューニングしているため、素朴な実装だとGPUの性能が出せていないと思われる。
速度を出すには、メモリ転送をStreamを使うなどチューニングが必要そうだ。
CUDAのチューニングノウハウをほとんど持っていないので、Chainerのソースをみるなどしてもう少し調査してみるつもりだ。
マルチGPUについては、効果があるはずなので、別途測定を行う予定。
メモリ転送をストリームを使って並列化してみたが、
| playout/sec | |
| 直列 | 4853 |
| ストリーム | 4202 |
という結果になり、かえって遅くなってしまった。
GPU計算中にメモリ転送できる箇所が入力層と出力層に一部しかなく、cudaStreamSynchronizeを実行する分、遅くなってしまったようだ。
ストリームに対応させる際に、ホストメモリの確保をC++のデフォルトヒープから、cudaHostAllocによって割り当てられるピンメモリに変更したが、そのことによって速くなることはなかった。
以上から、推論の実行時間のほとんどはGPUの計算に費やされており、Python経由でChainerから実行しても、cuDNNで直接実行しても差は限定的ということが言えそうだ。それでも、Chainerよりも7%程速くなっているので、少しは効果があった。
前回実装した、Chainerで学習したモデルを使用してcuDNNで推論するコードを、マルチGPUで並列に動作するようにした。
cuDNNをマルチスレッドで、スレッドに別々のGPUを割り当てて使用する場合、それぞれのスレッドでcudaSetDevice()を呼び出し、GPU IDを指定する。
cudaSetDevice()を呼び出し後に、cudnnCreate()を呼び出しcuDNNのハンドルを取得し、スレッドごとに異なるcuDNNのハンドルを使用する。
前回までのコードは、cudnnCreate()をNNクラスの静的クラス変数の初期化で呼び出していたが、上記の仕様に対応するため、NNクラスの初期化時に呼び出し、
スレッドごとにNNクラスのインスタンスを作成するように変更した。
以上の変更で、マルチGPUで動作できるようになった。
マルチGPUによって、高速化の効果があるか確認するために、以下の通り測定を行った。
| Chainer(シングルGPU) | 4767 ms |
| シングルGPU | 1657 ms |
| マルチGPU | 915 ms |
比較対象として、Chainerでシングルスレッドで推論した時間、cuDNNでシングルスレッドシングルGPUで推論した時間を記載している。
マルチGPUは、シングルGPUの約1.8倍高速になっている。
以前に将棋AIの推論でChainerを使ってマルチGPUに対応させたが、マルチGPU化の効果はほとんどなかった。
ChainerをPython経由で使用しているため、PythonのGIL機構によりマルチスレッドでの性能がでないためである。
mnistCUDNN\mnistCUDNN>..\x64\Release\mnistCUDNN.exe 2 There are 2 CUDA capable devices on your machine : device 0 sms 80 Capabilities 7.0, SmClock 1455 Mhz, MemSize (Mb) 12288, MemClock 850 Mhz, Ecc=0, boardGroupID=0 device 1 sms 28 Capabilities 6.1, SmClock 1670.5 Mhz, MemSize (Mb) 11264, MemClock 5505 Mhz, Ecc=0, boardGroupID=1 gpu:0 1000 iterations 867 [ms] gpu:1 1000 iterations 751 [ms] total time = 894 [ms]
gpu:0がTitan Vで、推論に867 [ms]を使っている。
gpu:1がGeForce GTX 1080 Tiで、推論に751 [ms]を使っている。
Titan Vの方が遅いため、トータルの時間は、Titan Vの実行時間+αになっている。
Titan Vの方が遅いのは、クロック数の違いによる。
| ベースクロック (MHz) | メモリクロック (MHz) | |
| Titan V | 1200 | 850 |
| Geforce GTX 1080 Ti | 1480 | 1100 |
モデルのサイズ、バッチサイズが小さい場合は、Geforce GTX 1080 Tiの方が高速に動作する。
Titan Vの性能を引き出せるのは、サイズの大きいモデルでバッチサイズを大きくして学習する場合である。
推論の速度だけであれば、モデルによってはクロック数の高いGeforce GTX 1080 Tiを複数枚用意した方がよさそうである。
測定に使用したソースコードを公開しました。
github.com
前回実装した、Chainerで学習したモデルを使用してcuDNNで推論するコードを、Residual Network(ResNet)構成にした。
推論時には、テンソルの加算を行うだけで特に難しいことはない。
ResNetは1ブロックのみで、ブロック内の畳み込み層は1層のみとした。
from chainer import Chain import chainer.functions as F import chainer.links as L # ネットワーク定義 k = 16 fcl = 256 class NN(Chain): def __init__(self): super(NN, self).__init__() with self.init_scope(): self.conv1 = L.Convolution2D(in_channels = 1, out_channels = k, ksize = 3, pad = 1) self.conv2 = L.Convolution2D(in_channels = k, out_channels = k, ksize = 3, pad = 1) self.conv3 = L.Convolution2D(in_channels = k, out_channels = k, ksize = 3, pad = 1) self.l4 = L.Linear(7*7*k, fcl) self.l5 = L.Linear(fcl, 10) self.bn1 = L.BatchNormalization(k) self.bn2 = L.BatchNormalization(k) def __call__(self, x): h = self.conv1(F.reshape(x, (len(x), 1, 28, 28))) h = self.bn1(h) h1 = F.relu(h) # resnet block h = self.conv2(h1) h = self.bn2(h) h = h + h1 h = F.max_pooling_2d(F.relu(h), 2) h = self.conv3(h) h = F.max_pooling_2d(F.relu(h), 2) h = F.relu(self.l4(h)) h = F.dropout(h, ratio=0.4) return self.l5(h)
cuDNNを使用してテンソルの加算を行うクラスを追加した。
class Add { public: void operator() (cudnnHandle_t handle, cudnnTensorDescriptor_t xDesc, float* x, float* y) { const float alpha = 1.0f; const float beta = 1.0f; checkCUDNN(cudnnAddTensor(handle, &alpha, xDesc, x, &beta, xDesc, y)); } };
ResNet構成のネットワーク定義を以下の通り実装した。
class NN { public: typedef float x_t[batch_size][1][IMAGE_H][IMAGE_W]; typedef float y_t[batch_size][10]; NN(); ~NN(); void load_model(const char* filename); void foward(x_t x, y_t y); private: static CudnnHandle cudnnHandle; static CublasHandle cublasHandle; static const int k = 16; static const int fcl = 256; ConvLayer<k, 1, 3, 1> conv1; Bias<k, 1, 1> bias1; ConvLayer<k, k, 3, 1> conv2; Bias<k, 1, 1> bias2; ConvLayer<k, k, 3, 1> conv3; Bias<k, 1, 1> bias3; Linear<7 * 7 * k, fcl> l4; Bias<fcl, 1, 1> bias4; Linear<fcl, 10> l5; Bias<10, 1, 1> bias5; BatchNormalization<k> bn1; BatchNormalization<k> bn2; ReLU relu; MaxPooling2D<2> max_pooling_2d; Add add; CudnnTensorDescriptor xDesc; CudnnTensorDescriptor h1Desc; CudnnTensorDescriptor h2Desc; CudnnTensorDescriptor h3Desc; CudnnTensorDescriptor h4Desc; CudnnTensorDescriptor h5Desc; CudnnTensorDescriptor h6Desc; CudnnTensorDescriptor yDesc; float* x_dev; float* h1_dev; float* h1_bn_dev; float* h2_dev; float* h2_bn_dev; float* h3_dev; float* h4_dev; float* h5_dev; float* h6_dev; float* y_dev; };
NN::NN() { conv1.get_xdesc(xDesc, batch_size, IMAGE_H, IMAGE_W); const int h1_h = conv1.get_yh(IMAGE_H); const int h1_w = conv1.get_yw(IMAGE_W); conv1.get_ydesc(h1Desc, batch_size, h1_h, h1_w); const int h3_h = max_pooling_2d.get_yh(h1_h); const int h3_w = max_pooling_2d.get_yw(h1_w); conv3.get_xdesc(h3Desc, batch_size, h3_h, h3_w); const int h4_h = conv3.get_yh(h3_h); const int h4_w = conv3.get_yw(h3_w); conv3.get_ydesc(h4Desc, batch_size, h4_h, h4_w); const int h5_h = max_pooling_2d.get_yh(h4_h); const int h5_w = max_pooling_2d.get_yw(h4_w); max_pooling_2d.get_desc(h5Desc, batch_size, k, h5_h, h5_w); l4.get_ydesc(h6Desc, batch_size); l5.get_ydesc(yDesc, batch_size); // init conv layers conv1.init(cudnnHandle, xDesc, h1Desc); conv2.init(cudnnHandle, h1Desc, h1Desc); conv3.init(cudnnHandle, h3Desc, h4Desc); // malloc checkCudaErrors(cudaMalloc((void**)&x_dev, conv1.get_xsize(batch_size, IMAGE_H, IMAGE_W))); checkCudaErrors(cudaMalloc((void**)&h1_dev, conv1.get_ysize(batch_size, h1_h, h1_w))); checkCudaErrors(cudaMalloc((void**)&h1_bn_dev, conv1.get_ysize(batch_size, h1_h, h1_w))); checkCudaErrors(cudaMalloc((void**)&h2_dev, conv2.get_ysize(batch_size, h1_h, h1_w))); checkCudaErrors(cudaMalloc((void**)&h2_bn_dev, conv2.get_ysize(batch_size, h1_h, h1_w))); checkCudaErrors(cudaMalloc((void**)&h3_dev, conv3.get_xsize(batch_size, h3_h, h3_w))); checkCudaErrors(cudaMalloc((void**)&h4_dev, conv3.get_ysize(batch_size, h4_h, h4_w))); checkCudaErrors(cudaMalloc((void**)&h5_dev, batch_size * k * h5_h * h5_w * sizeof(float))); checkCudaErrors(cudaMalloc((void**)&h6_dev, batch_size * fcl * sizeof(float))); checkCudaErrors(cudaMalloc((void**)&y_dev, batch_size * 10 * sizeof(float))); } NN::~NN() { checkCudaErrors(cudaFree(x_dev)); checkCudaErrors(cudaFree(h1_dev)); checkCudaErrors(cudaFree(h1_bn_dev)); checkCudaErrors(cudaFree(h2_dev)); checkCudaErrors(cudaFree(h2_bn_dev)); checkCudaErrors(cudaFree(h3_dev)); checkCudaErrors(cudaFree(h4_dev)); checkCudaErrors(cudaFree(h5_dev)); checkCudaErrors(cudaFree(h6_dev)); checkCudaErrors(cudaFree(y_dev)); }
前回までの畳み込みの読み込みと同じでResNet構成によって特別な処理はない。
void NN::load_model(const char* filepath) { // load nn params ParamMap params; load_npz(filepath, params); conv1.set_param(params["conv1/W.npy"].data); bias1.set_bias(params["conv1/b.npy"].data); conv2.set_param(params["conv2/W.npy"].data); bias2.set_bias(params["conv2/b.npy"].data); conv3.set_param(params["conv3/W.npy"].data); bias3.set_bias(params["conv3/b.npy"].data); l4.set_param(params["l4/W.npy"].data); bias4.set_bias(params["l4/b.npy"].data); l5.set_param(params["l5/W.npy"].data); bias5.set_bias(params["l5/b.npy"].data); bn1.set_param(params["bn1/gamma.npy"].data, params["bn1/beta.npy"].data, params["bn1/avg_mean.npy"].data, params["bn1/avg_var.npy"].data); bn2.set_param(params["bn2/gamma.npy"].data, params["bn2/beta.npy"].data, params["bn2/avg_mean.npy"].data, params["bn2/avg_var.npy"].data); }
ResNetブロックの処理とテンソルの加算処理を追加した。
void NN::foward(x_t x, y_t y) { // input checkCudaErrors(cudaMemcpy(x_dev, x, sizeof(x_t), cudaMemcpyHostToDevice)); // conv1 conv1(cudnnHandle, xDesc, x_dev, h1Desc, h1_dev); bias1(cudnnHandle, h1Desc, h1_dev); bn1(cudnnHandle, h1Desc, h1_dev, h1_bn_dev); relu(cudnnHandle, h1Desc, h1_bn_dev); // resnet block conv2(cudnnHandle, h1Desc, h1_bn_dev, h1Desc, h2_dev); bias2(cudnnHandle, h1Desc, h2_dev); bn2(cudnnHandle, h1Desc, h2_dev, h2_bn_dev); add(cudnnHandle, h1Desc, h1_bn_dev, h2_bn_dev); relu(cudnnHandle, h1Desc, h2_bn_dev); max_pooling_2d(cudnnHandle, h1Desc, h2_bn_dev, h3Desc, h3_dev); // conv3 conv3(cudnnHandle, h3Desc, h3_dev, h4Desc, h4_dev); bias3(cudnnHandle, h4Desc, h4_dev); relu(cudnnHandle, h4Desc, h4_dev); max_pooling_2d(cudnnHandle, h4Desc, h4_dev, h5Desc, h5_dev); // fcl l4(cublasHandle, batch_size, h5_dev, h6_dev); bias4(cudnnHandle, h6Desc, h6_dev); relu(cudnnHandle, h6Desc, h6_dev); l5(cublasHandle, batch_size, h6_dev, y_dev); bias5(cudnnHandle, yDesc, y_dev); // output checkCudaErrors(cudaMemcpy(y, y_dev, sizeof(y_t), cudaMemcpyDeviceToHost)); }
テンソル加算の実際の処理は、layers.hのAddクラスのoperator()に実装している。
1層目の畳み込み層の活性化関数の出力h1_bn_devをResNetブロックの入力にして、ResNetブロックの出力h2_bn_devをh1_bn_devに加算している。
ResNetブロックの入力と出力のテンソルのサイズは変わらないため、テンソルディスクリプタは1層目の畳み込み層と共通のh1Descを使用している。
前回実装した、Chainerで学習したモデルを使用してcuDNNで推論するコードに、dropoutを追加した。
dropoutは学習時のみ処理を行うため、推論時には何もしなくてよい。
つまり、推論のネットワーク定義にはdropoutは必要ない。
学習用のネットワーク定義にdropoutを追加した。
from chainer import Chain import chainer.functions as F import chainer.links as L # ネットワーク定義 k = 16 fcl = 256 class NN(Chain): def __init__(self): super(NN, self).__init__() with self.init_scope(): self.conv1 = L.Convolution2D(in_channels = 1, out_channels = k, ksize = 3, pad = 1) self.conv2 = L.Convolution2D(in_channels = k, out_channels = k, ksize = 3, pad = 1) self.l3 = L.Linear(7*7*k, fcl) self.l4 = L.Linear(fcl, 10) self.bn1 = L.BatchNormalization(k) self.bn2 = L.BatchNormalization(k) def __call__(self, x): h = self.conv1(F.reshape(x, (len(x), 1, 28, 28))) h = self.bn1(h) h = F.max_pooling_2d(F.relu(h), 2) h = self.conv2(h) h = self.bn2(h) h = F.max_pooling_2d(F.relu(h), 2) h = F.relu(self.l3(h)) h = F.dropout(h, ratio=0.4) return self.l4(h)
推論のネットワーク定義は、前回のままでよい。
前回実装した、Chainerで学習したモデルを使用してcuDNNで推論するコードに、BatchNormalizationを追加した。
BatchNormalizationも、cuDNNにAPIが用意されているため、簡単に使用できる。
まず、Chainerで学習するモデルにBatchNormalizationを追加した。
from chainer import Chain import chainer.functions as F import chainer.links as L # ネットワーク定義 k = 16 fcl = 256 class NN(Chain): def __init__(self): super(NN, self).__init__() with self.init_scope(): self.conv1 = L.Convolution2D(in_channels = 1, out_channels = k, ksize = 3, pad = 1) self.conv2 = L.Convolution2D(in_channels = k, out_channels = k, ksize = 3, pad = 1) self.l3 = L.Linear(7*7*k, fcl) self.l4 = L.Linear(fcl, 10) self.bn1 = L.BatchNormalization(k) self.bn2 = L.BatchNormalization(k) def __call__(self, x): h = self.conv1(F.reshape(x, (len(x), 1, 28, 28))) h = self.bn1(h) h = F.max_pooling_2d(F.relu(h), 2) h = self.conv2(h) h = self.bn2(h) h = F.max_pooling_2d(F.relu(h), 2) h = F.relu(self.l3(h)) return self.l4(h)
cuDNNを使用して推論を行うC++のコードにBatchNormalizationのレイヤーを追加した。
template<const int k> class BatchNormalization { public: BatchNormalization() : bnScale(nullptr), bnBias(nullptr), estimatedMean(nullptr), estimatedVariance(nullptr) { const size_t size = k; checkCudaErrors(cudaMalloc((void**)&bnScale, size * sizeof(float))); checkCudaErrors(cudaMalloc((void**)&bnBias, size * sizeof(float))); checkCudaErrors(cudaMalloc((void**)&estimatedMean, size * sizeof(float))); checkCudaErrors(cudaMalloc((void**)&estimatedVariance, size * sizeof(float))); } ~BatchNormalization() { checkCudaErrors(cudaFree(bnScale)); checkCudaErrors(cudaFree(bnBias)); checkCudaErrors(cudaFree(estimatedMean)); checkCudaErrors(cudaFree(estimatedVariance)); } void operator() (cudnnHandle_t handle, cudnnTensorDescriptor_t xDesc, float* x, float* y) { const float alpha = 1.0f; const float beta = 0.0f; const double eps = 2e-5; checkCUDNN(cudnnDeriveBNTensorDescriptor(bnScaleBiasMeanVarDesc, xDesc, CUDNN_BATCHNORM_SPATIAL)); checkCUDNN(cudnnBatchNormalizationForwardInference(handle, CUDNN_BATCHNORM_SPATIAL, &alpha, &beta, xDesc, x, xDesc, y, bnScaleBiasMeanVarDesc, bnScale, bnBias, estimatedMean, estimatedVariance, eps)); } void set_param(float* bnScale, float *bnBias, float *estimatedMean, float *estimatedVariance) { const size_t size = k; checkCudaErrors(cudaMemcpy(this->bnScale, bnScale, size * sizeof(float), cudaMemcpyHostToDevice)); checkCudaErrors(cudaMemcpy(this->bnBias, bnBias, size * sizeof(float), cudaMemcpyHostToDevice)); checkCudaErrors(cudaMemcpy(this->estimatedMean, estimatedMean, size * sizeof(float), cudaMemcpyHostToDevice)); checkCudaErrors(cudaMemcpy(this->estimatedVariance, estimatedVariance, size * sizeof(float), cudaMemcpyHostToDevice)); } private: CudnnTensorDescriptor bnScaleBiasMeanVarDesc; float *bnScale; float *bnBias; float *estimatedMean; float *estimatedVariance; };
ネットワークの定義に、上記で実装したBatchNormalizationの層を追加した。
#pragma once #include "layers.h" const int IMAGE_H = 28; const int IMAGE_W = 28; const int batch_size = 2; class NN { public: typedef float x_t[batch_size][1][IMAGE_H][IMAGE_W]; typedef float y_t[batch_size][10]; NN(); ~NN(); void load_model(const char* filename); void foward(x_t x, y_t y); private: static CudnnHandle cudnnHandle; static CublasHandle cublasHandle; static const int k = 16; static const int fcl = 256; ConvLayer<k, 1, 3, 1> conv1; Bias<k, 1, 1> bias1; ConvLayer<k, k, 3, 1> conv2; Bias<k, 1, 1> bias2; Linear<7 * 7 * k, fcl> l3; Bias<fcl, 1, 1> bias3; Linear<fcl, 10> l4; Bias<10, 1, 1> bias4; BatchNormalization<k> bn1; BatchNormalization<k> bn2; ReLU relu; MaxPooling2D<2> max_pooling_2d; CudnnTensorDescriptor xDesc; CudnnTensorDescriptor h1Desc; CudnnTensorDescriptor h2Desc; CudnnTensorDescriptor h3Desc; CudnnTensorDescriptor h4Desc; CudnnTensorDescriptor h5Desc; CudnnTensorDescriptor yDesc; float* x_dev; float* h1_dev; float* h1_bn_dev; float* h2_dev; float* h3_dev; float* h3_bn_dev; float* h4_dev; float* h5_dev; float* y_dev; };
Chainerで学習したモデルの読み込みにBatchNormalizationのパラメータの読み込みを追加した。
void NN::load_model(const char* filepath) { // load nn params ParamMap params; load_npz(filepath, params); conv1.set_param(params["conv1/W.npy"].data); bias1.set_bias(params["conv1/b.npy"].data); conv2.set_param(params["conv2/W.npy"].data); bias2.set_bias(params["conv2/b.npy"].data); l3.set_param(params["l3/W.npy"].data); bias3.set_bias(params["l3/b.npy"].data); l4.set_param(params["l4/W.npy"].data); bias4.set_bias(params["l4/b.npy"].data); bn1.set_param(params["bn1/gamma.npy"].data, params["bn1/beta.npy"].data, params["bn1/avg_mean.npy"].data, params["bn1/avg_var.npy"].data); bn2.set_param(params["bn2/gamma.npy"].data, params["bn2/beta.npy"].data, params["bn2/avg_mean.npy"].data, params["bn2/avg_var.npy"].data); }
Chainerで学習したモデルのBatchNormalization層は、gamma.npy、beta.npy、avg_mean.npy、avg_var.npyというファイル名で格納されている。
推論の処理にBatchNormalization層を追加した。
void NN::foward(x_t x, y_t y) { // input checkCudaErrors(cudaMemcpy(x_dev, x, sizeof(x_t), cudaMemcpyHostToDevice)); // conv1 conv1(cudnnHandle, xDesc, x_dev, h1Desc, h1_dev); bias1(cudnnHandle, h1Desc, h1_dev); bn1(cudnnHandle, h1Desc, h1_dev, h1_bn_dev); relu(cudnnHandle, h1Desc, h1_bn_dev); max_pooling_2d(cudnnHandle, h1Desc, h1_bn_dev, h2Desc, h2_dev); // conv2 conv2(cudnnHandle, h2Desc, h2_dev, h3Desc, h3_dev); bias2(cudnnHandle, h3Desc, h3_dev); bn2(cudnnHandle, h3Desc, h3_dev, h3_bn_dev); relu(cudnnHandle, h3Desc, h3_bn_dev); max_pooling_2d(cudnnHandle, h3Desc, h3_bn_dev, h4Desc, h4_dev); // fcl l3(cublasHandle, batch_size, h4_dev, h5_dev); bias3(cudnnHandle, h5Desc, h5_dev); relu(cudnnHandle, h5Desc, h5_dev); l4(cublasHandle, batch_size, h5_dev, y_dev); bias4(cudnnHandle, yDesc, y_dev); // output checkCudaErrors(cudaMemcpy(y, y_dev, sizeof(y_t), cudaMemcpyDeviceToHost)); }
cuDNNでBatchNormalizationの順伝播は、cudnnBatchNormalizationForwardInferenceで行う。
実際の処理は、layers.hのBatchNormalizationクラスのoperator()に実装している。
Chainerで学習したモデルから読み込んだ、gamma.npy、beta.npy、avg_mean.npy、avg_var.npyをcudnnBatchNormalizationForwardInferenceのパラメータに渡してやればよい。
cudnnBatchNormalizationForwardInferenceの2番目のパラメータのmodeには、CUDNN_BATCHNORM_SPATIALを指定する。
畳み込み層の直後でBatchNormalizationを使用する場合は、CUDNN_BATCHNORM_SPATIALを使用する。
ChainerのソースでもCUDNN_BATCHNORM_SPATIALが使用されていることを確認した。
cudnnBatchNormalizationForwardInferenceの最後の引数のepsの値は、Chainerのデフォルト値(2e-5)を固定で入力している。
昨日の日記の続きです。
C++のコードでcuDNNを直接使用して推論を行った場合と、Chainerを使用して推論を行った場合の実行時間の比較を行った。
| Chainer | 989.7 ms |
| cuDNN | 55 ms |
cuDNNを直接使用して推論を行った場合、Chainerを使用して推論を行った場合より、約18倍高速になった。
シングルスレッドでの計測のため、マルチスレッドでマルチGPUを使う場合、PythonにはGILの制約があるため、さらに高速化の効果があると思われる。
推論に速度を求める場合は、cuDNNを直接使用することも検討した方が良さそうだ。
計測には以下のコードを使用した。
github.com
※layers.hのConvLayerのset_paramの「const size_t size = k * k * fsize * fsize;」は、「const size_t size = c * k * fsize * fsize;」の間違いです。
