こんにちは。アプリケーション共同開発部新卒エンジニアの高坂です。
本記事では、最近特に目にする機会の増えた人工ニューラルネットワークの演算処理を支える GPGPU 演算を簡単に実装できる NVIDIA 社の Thrust ライブラリについて、その使い方を簡単に説明します。
はじめに
コンピュータにプログラムを実行させる場合、CPU に演算を行わせるのが一般的かと思います。近年の情報技術の向上により、CPU コアあたりの演算能力の向上、またコア数・スレッド数の増加による性能向上が著しく、コンピュータが現実的な時間で行えることの幅が確実に広がっています。
しかしながら、現在のフラッグシップサーバ CPU として Intel Xeon Platinum 8180 プロセッサを例にとると、一個あたりのコア数は 28 コア 56 スレッド、8ソケット構成をサポートしているので理論上 224 コア 448 スレッドが限界となりますが、非常に高価であり現実的な構成とはいえません。より現実的な構成として、ネットワークを介し、複数台のコンピュータを協調動作させて演算させる方法もありますが、この場合はネットワークの通信速度による影響があり、それほど単純に演算時間を高速化できる訳ではありません。
冒頭にも述べましたが、最近人工ニューラルネットワークのブームが来ています。このブームの火付け役として、畳み込みニューラルネットワーク（Convolutional Neural Network, CNN）等のディープラーニング（Deep Learning、深層学習）が開拓されたことが要因の一つとして挙げられるかと思います。ディープラーニングの領域が開拓された要因として、従来の4層以上の複雑なニューラルネットワークにおいて課題となっていた局所最適解や勾配消失といったものによる学習精度向上の限界が、ヒントンらによる研究成果によって改善されつつあることが挙げられます [1] 。また、それと共に、GPGPU の活用による演算能力・速度の飛躍的な向上も挙げられます。
人工ニューラルネットワークに関するプログラムを触ったことがある、作ったことがある方であれば分かるかと思いますが、人工ニューラルネットワークにおいては行列同士を何度も演算する必要があります。これを高速化するためには演算を並列化し同時に処理させる必要がありますが、一般的な CPU では並列度に限度があり、ディープラーニングのような特に複雑なネットワーク構造を持つものでは現実的な時間での利用ができません。この問題を解決するために、GPGPU が注目されています。
GPGPU とは
GPGPU とは、ざっくり説明すると従来までは CPU で行っていた演算処理を、3D グラフィックスの演算のためにより多くのコアが載った GPU 上で行わせるものです。例えば NVIDIA の現在最新の GPGPU 用 GPU である NVIDIA Tesla V100 は、NVIDIA CUDA Core を 5,120 個搭載しており、それぞれの CUDA Core ごとに1クロックで積和算を行うことができます。
また、GPU は 3D グラフィックスデータをより高速に演算することを目的に開発されていたため、同価格帯の CPU と比較してメモリバンド幅が広く、より高速なメモリアクセスが可能です。NVIDIA GPU はメモリやキャッシュ、複数個のシェーダプロセッサ（NVIDIA GPU では CUDA Core）を持つストリーミングマルチプロセッサが複数個実装され形成されており、演算する際には各ストリーミングマルチプロセッサ毎にデータを渡して同一命令を一気に行うアーキテクチャのため、例えば複雑な条件分岐を含むようなメモリへのランダムアクセスを伴う場合やスレッド間同期を何度も行う必要のある演算処理については GPU の持つ能力を発揮できませんが、人工ニューラルネットワークにおいて必要となるような比較的単純な演算処理においては GPU の持つ演算能力を活かして、処理にかかる時間をグッと短縮できます。
NVIDIA Thrust
NVIDIA の GPU を用いて GPGPU を行う際は、CUDA を用いることになります。 素の CUDA プログラミングを行う場合、cudaMalloc と cudaFree を用いて GPU デバイス上のメモリ管理をキッチリ行う必要があることをはじめとして非常に扱いが面倒な印象で、私自身、実は触ったことがありません…。
NVIDIA Thrust ライブラリを用いることで、C++ STL の vector に対応するものがメモリ管理をさほど意識せずに使えるほか、ソートや総和などの基本的な行列演算であれば、専用メソッドが用意されており一行で実装できます。
利用方法
まずは、CUDA に対応した NVIDIA GPU を搭載したコンピュータに、NVIDIA 公式のデバイスドライバをインストールします。
次に、CUDA Toolkit をインストールします。これにより、CUDA プログラミングを行うために必要な nvcc コンパイラをはじめとしたツール一式のほか、Thrust ライブラリもインストールされます。
CUDA コードを含むソースコードには .cu を、ヘッダファイルには .cuh を拡張子につけるのが一般的です。
行列間の乗算
二つの行列について、それぞれの i 番目同士を掛けた結果を得たい場合は、以下のように実装します。
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/transform.h>
#include <iostream>
int main() {
// ホスト（CPU）側メモリに領域を確保
thrust::host_vector<int> host_input_first{2, -4, 8};
thrust::host_vector<int> host_input_second{1, 2, 3};
thrust::host_vector<int> host_output(3);
// デバイス（GPU）側メモリに領域を確保
thrust::device_vector<int> device_input_first(3);
thrust::device_vector<int> device_input_second(3);
thrust::device_vector<int> device_output(3);
// ホスト側メモリで持っているデータを、デバイス側メモリにコピーする
// GPU で演算する前に、必ずこのコピー処理を行うこと。でないと速度が出ない
thrust::copy(host_input_first.begin(), host_input_first.end(), device_input_first.begin());
thrust::copy(host_input_second.begin(), host_input_second.end(), device_input_second.begin());
// device_output[i] = device_input_first[i] * device_input_second[i];
thrust::transform(device_input_first.begin(), device_input_first.end(), device_input_second.begin(), device_output.begin(), thrust::multiplies<int>());
// デバイス側メモリにある演算結果を、ホスト側メモリにコピーする
thrust::copy(device_output.begin(), device_output.end(), host_output.begin());
std::cout << host_output[0] << ", " << host_output[1] << ", " << host_output[2] << std::endl;
// 2, -8, 24
return 0;
}
行列のソート
一つの行列について、要素のソートを行いたい場合は以下のように実装します。
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/sort.h>
#include <thrust/functional.h>
#include <iostream>
#include <iterator>
int main() {
thrust::host_vector<int> host_input{5, 1, 9, 3, 7};
thrust::device_vector<int> device_vec(5);
thrust::copy(host_input.begin(), host_input.end(), device_vec.begin());
// 昇順でソートする
thrust::sort(device_vec.begin(), device_vec.end());
thrust::copy(device_vec.begin(), device_vec.end(), std::ostream_iterator<int>(std::cout, ", "));
// 1, 3, 5, 7, 9
// 降順でソートする
thrust::sort(device_vec.begin(), device_vec.end(), thrust::greater<int>());
thrust::copy(device_vec.begin(), device_vec.end(), std::ostream_iterator<int>(std::cout, ", "));
// 9, 7, 5, 3, 1
return 0;
}
任意の行列変換
Thrust ライブラリで用意されていない、任意の行列変換を行いたい場合は、以下のように実装します。
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/transform.h>
#include <iostream>
struct sample_functor {
// __device__ とつけることで、このメソッドがデバイス（GPU）側で実行される
// ホスト側で実行したいなら __host__ とつける。両方をつけることも可能
__device__ double operator () (const double& x) const {
return 2.0 * x + 1.0;
}
};
int main() {
thrust::host_vector<double> host_input{1.1, 3.3, 2.2};
thrust::host_vector<double> host_output(3);
thrust::device_vector<double> device_input(3);
thrust::device_vector<double> device_output(3);
thrust::copy(host_input.begin(), host_input.end(), device_input.begin());
// 自ら定義した行列変換メソッドを用いて行列変換を行う
thrust::transform(device_input.begin(), device_input.end(), device_output.begin(), sample_functor());
thrust::copy(device_output.begin(), device_output.end(), host_output.begin());
std::cout << host_output[0] << ", " << host_output[1] << ", " << host_output[2] << std::endl;
// 3.2, 7.6, 5.4
return 0;
}
二つの行列を用いた行列変換を行いたい場合は、以下のように実装します。
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/transform.h>
#include <iostream>
struct sample_functor {
double alpha;
double beta;
sample_functor(double _alpha, double _beta) {
alpha = _alpha;
beta = _beta;
}
__device__ double operator() (const double& x, const double& y) const {
return alpha * x + (1.0 - alpha) * (beta * y);
}
};
int main() {
thrust::host_vector<double> host_x{1.1, 3.3, 2.2};
thrust::host_vector<double> host_y{6.6, 7.7, 8.8};
thrust::host_vector<double> host_output(3);
thrust::device_vector<double> device_x(3);
thrust::device_vector<double> device_y(3);
thrust::device_vector<double> device_output(3);
double alpha = 0.005;
double beta = 0.1;
thrust::copy(host_x.begin(), host_x.end(), device_x.begin());
thrust::copy(host_y.begin(), host_y.end(), device_y.begin());
thrust::transform(device_x.begin(), device_x.end(), device_y.begin(), device_output.begin(), sample_functor(alpha, beta));
thrust::copy(device_output.begin(), device_output.end(), host_output.begin());
std::cout << host_output[0] << ", " << host_output[1] << ", " << host_output[2] << std::endl;
// 0.6622, 0.78265, 0.8866000000000002
return 0;
}
コンパイル・実行方法
CUDA コードを含むソースコードは、nvcc コンパイラを用いてコンパイルします。以下は、C++ 11 を有効にし、最適化オプションを設定し、CUDA 付属のヘッダファイル（CUDA Toolkit インストール先ディレクトリ内にあります）とプロジェクトのヘッダファイルをインクルードしてコンパイルする例になります。
$ nvcc -std=c++11 -I/PATH/TO/CUDA/INCLUDES -I/PATH/TO/PROJECT/INCLUDES -O3 *.cu -o sample
コンパイルして生成された実行ファイルは、通常の C、C++ と同じように実行可能です。
$ ./sample
CUDA プログラム開発時に意識すべき点
CUDA プログラムを開発する際は、演算速度向上の恩恵を受けるために以下のことを意識する必要があります。
- デバイスからホストメモリを参照することは出来ないので、GPU で演算を行う前に必ず対象となるデータをホストメモリからデバイスメモリにコピーしてください。また演算が終わったらデバイスメモリからホストメモリーにコピーする必要があります。
- ホストメモリ – デバイスメモリ間のデータ転送幅は、一般的に GPGPU の演算能力よりもはるかに低いです。そのため行列のうちで演算に必要となる分だけを都度ホストメモリ – デバイスメモリ間でコピーするのは良くありません。あらかじめ演算に用いる全てのデータをデバイスメモリにコピーするようにし、ホストメモリ – デバイスメモリ間のやりとりが最小限になるよう注意を払うべきです。
- ホストメモリ – デバイスメモリ間のデータ転送は非常に遅いため、演算に用いる行列のサイズが小さい場合は GPGPU を用いることによる演算速度の高速化の恩恵分をデータ転送に必要となる時間で食いつぶしてしまう可能性があります。この点のバランスを見て、GPGPU を採用するか否かを考える必要があります。
- GPU はもともと画像演算、リアルタイム 3D グラフィックスをより高速に演算するために開発されたものです。その対象となるのは整数や単精度浮動小数点の演算がメインであり、倍精度浮動小数点の演算では速度が落ちる可能性があります。GPGPU の採用を検討する際には演算対象が単精度で十分なのか、それとも倍精度が必要なのかを考慮する必要があります。
より詳しく学びたい方は、NVIDIA 公式の Performance Guildeline をご一読ください。
さいごに
CUDA による GPGPU プログラムを開発する際に、Toolkit に付属する Thrust ライブラリを用いることで極めて簡単に実装を進めることができます。上記に挙げたサンプルプログラムを見ていただければ如何に簡単であるかがわかるかと思います。
開発に際して意識すべきなのは、ホストメモリ – デバイスメモリ間のデータコピーがボトルネックになるため、少しずつデータをコピーするのではなく一気にコピーするようプログラムを実装するということです。また、そうして実装したプログラムにより、データコピーのボトルネックを上回るだけの演算速度向上が見込めるかどうかを見極める必要があるということです。また、倍精度浮動小数点を含む演算を行う必要がある場合は、演算速度が落ちる可能性があることも考慮する必要があります。
以上を踏まえて上手く GPU を扱えば、CPU での演算速度と比較して数十〜数百倍以上の速度向上が見込めます。NVIDIA 製 GPU をお持ちの方は、Thrust ライブラリを用いて CUDA プログラムをサクッと実装し、GPGPU の世界を垣間見られてはいかかでしょうか。
参考リンク
脚注
-
[1] G. E. Hinton, and R. R. Salakhutdinov. Reducing the Dimensionality of Data with Neural Networks, Science 28 Jul 2006: Vol.313, Issue 5786, pp.504-507.
フェンリルのオフィシャル Twitter アカウントでは、フェンリルプロダクトの最新情報などをつぶやいています。よろしければフォローしてください！
フェンリルの開発者アカウントでは、フェンリルプロダクトの最新開発情報などをつぶやいています。よろしければフォローしてください！
フェンリル採用チームの Twitter アカウントです。応募前のお問い合わせや、ちょっとした相談ごとなどお気軽にどうぞ！
フェンリルの Facebook ページでは、最新トピックをお知らせしています。よろしければいいね！してください！
Sleipnir の Facebook ページでは、ユーザーの方たちとのコミュニケーションや最新情報の投稿などを行なっています。よろしければいいね！してください！
関連記事
