はじめに
先のブログでPyTorch C++を紹介した（その1,その2,その3）。CUDAプログラミングを知らなくても、GPUを利用する深層学習のプログラムをC++で実装できる。ところで、NVIDIAはCUDAを利用するためのC++インターフェースThrustを公開している。これはテンプレートライブラリである。今回はこれを紹介する。
コンテナ
Thrustは2種類のコンテナを提供する。
前者はCPU側のメモリ上に、後者はGPU側のメモリ上にバッファを確保する。その使い方は以下の通り。
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 |
void test_container() { //_/_/_/ CPUメモリ上に確保する。 // thrust::host_vector<int> hv {1, 2, 3}; <-- エラーになります。 // 以下を真似することはできません。 std::vector<int> sv {1, 2, 3}; // ひとつずつ代入します。 thrust::host_vector<int> hv(3); hv[0] = 1; hv[1] = 2; hv[2] = 3; //_/_/_/ GPUメモリ上に確保する。 // 簡単です。 thrust::device_vector<int> dv(3); dv[0] = 1; dv[1] = 2; dv[2] = 3; //_/_/_/ CPUメモリ上のバッファをGPUメモリ上へ転送する。 // 簡単です。 dv = hv; //_/_/_/ GPUメモリ上のバッファをCPUメモリ上へ転送する。 // copy関数を使います。 thrust::copy(dv.begin(), dv.end(), hv.begin()); } |
thrust::host_vectorはstd::vectorとほぼ同じ機能を持つ。前者を使うメリットはデバイス側とのデータの遣り取りがスムーズになることである。
アルゴリズム
C++の標準テンプレートライブラリ(Standard Template Library:STL）には、「アルゴリズム」と総称される非常にたくさんの関数が用意されている。その１つにstd::transformがある。使い方は以下の通り。
|
1 2 3 4 5 6 7 8 9 |
void test_stl_algorithm() { std::vector<int> sv {1, 2, 3}; std::cout << sv[0] << " " << sv[1] << " " << sv[2] << std::endl; std::vector<int> dv(3); std::transform(sv.begin(), sv.end(), dv.begin(), std::negate<int>()); std::cout << dv[0] << " " << dv[1] << " " << dv[2] << std::endl; } |
上のstd::transformは、svの各値にマイナスを付けdvに代入している。これと同じ機能を持つGPU版は以下の通りである。
|
1 2 3 4 5 6 7 8 9 10 |
void test_algorithm() { thrust::device_vector<int> sv(3); thrust::sequence(sv.begin(), sv.end(), 1); std::cout << sv[0] << " " << sv[1] << " " << sv[2] << std::endl; thrust::device_vector<int> dv(3); thrust::transform(sv.begin(), sv.end(), dv.begin(), thrust::negate<int>()); std::cout << dv[0] << " " << dv[1] << " " << dv[2] << std::endl; } |
std::transformと同じ名前の関数thrust::transformが用意されているので、これを用いる。ところで、thrust::sequenceはSTLにない関数である。上の引数の場合、1始まりの連続整数値を生成する。
ここまでに紹介したインターフェースを用いて簡単な画像処理のプログラムを作成し、速度比較を行う。
画像処理プログラム
今回実装する画像処理の内容は以下の2つである。
- RGB画像をグレイ画像に変換する。
- グレイ画像に閾値を与えて2値化する。
まず最初に、グレイ化に使用する関数オブジェクト（Functor）を示す。
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 |
struct gray_converter { const float r_factor_ {0.299}; const float g_factor_ {0.587}; const float b_factor_ {0.114}; __host__ __device__ inline uchar operator()(const uchar3& rgb) const { float g = r_factor_ * rgb.z + g_factor_ * rgb.y + b_factor_ * rgb.x; if (g < 0) { g = 0; } if (g > 255) { g = 255; } return (uchar)(g); } }; |
10行目にグレイ画像への変換式が定義されている。それ以降の行で0から255までの値に収める処理をする。同様に2値化に使用するFunctorは以下のようになる。
|
1 2 3 4 5 6 7 8 9 10 11 12 13 |
struct binary_generator { const uint8_t threshold_; binary_generator(uint8_t threshold) : threshold_{threshold} {} __host__ __device__ inline uchar operator()(const uchar v) const { return v > threshold_ ? 255 : 0; } }; |
11行目で2値化処理をしている。これらのFunctorが持つメソッドoperator()には識別子「__host__ __device__」が付いている。これをつけることにより、CPU上でもGPU上でも動作する関数になる。この辺りはCUDAプログラミングの流儀である（詳細は割愛する）。これらのFunctorを用いた画像処理プログラムは以下の通り。
CPU版：
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 |
void test_cpu() { // load RGB image auto src_image = cv::imread(PATH); auto rows = src_image.rows; auto cols = src_image.cols; // copy image to source buffer // uchar3 means a sequence of (uchar,uchar,uchar) uchar3* ptr = reinterpret_cast<uchar3*>(src_image.data); std::vector<uchar3> src_buffer(ptr, ptr + rows * cols); // make destination buffer std::vector<uchar> dst_buffer(src_buffer.size()); cv::Mat gray_image {}; auto gc = gray_converter(); auto bg = binary_generator(128); auto start = std::chrono::system_clock::now(); for (auto i = 0; i < ITERATIONS; ++i) { std::transform(src_buffer.begin(), src_buffer.end(), dst_buffer.begin(), gc); std::transform(dst_buffer.begin(), dst_buffer.end(), dst_buffer.begin(), bg); gray_image = cv::Mat(rows, cols, CV_8UC1, dst_buffer.data()); } auto end = std::chrono::system_clock::now(); cv::imwrite("cpu_hoge.jpg", gray_image); // display time auto t = std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count(); std::cout << "cpu " << t << " [ms]" << std::endl; } |
GPU版：
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 |
void test_gpu() { // load RGB image auto src_image = cv::imread(PATH); auto rows = src_image.rows; auto cols = src_image.cols; // make ouput buffer on GPU device thrust::device_vector<uchar> dst_buffer(rows * cols); thrust::host_vector<uchar> host_image(rows * cols); cv::Mat gray_image {}; auto gc = gray_converter(); auto bg = binary_generator(128); auto start = std::chrono::system_clock::now(); // copy image to container on GPU device // uchar3 means a sequence of (uchar,uchar,uchar) uchar3* ptr = reinterpret_cast<uchar3*>(src_image.data); thrust::device_vector<uchar3> src_buffer(ptr, ptr + rows * cols); for (auto i = 0; i < ITERATIONS; ++i) { thrust::transform(src_buffer.begin(), src_buffer.end(), dst_buffer.begin(), gc); thrust::transform(dst_buffer.begin(), dst_buffer.end(), dst_buffer.begin(), bg); thrust::copy(dst_buffer.begin(), dst_buffer.end(), host_image.begin()); gray_image = cv::Mat(rows, cols, CV_8UC1, host_image.data()); } auto end = std::chrono::system_clock::now(); cv::imwrite("gpu_hoge.jpg", gray_image); auto t = std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count(); std::cout << "gpu " << t << " [ms]" << std::endl; } |
それぞれのプログラム内で、forループをITERATIONS（=1000）回まわす。その処理時間を計測すると以下のようになる。
| CPU版 | 8803 [ms] |
| GPU版 | 103 [ms] |
明らかにGPU版の方が高速である。今回検証に使用したマシーン（EC2インスタンス p2.xlarge）のスペックは以下の通りである。
| CPU | Intel Xeon 2.3GHz(x4), CPUメモリ 60GB |
| GPU | Tesla K80, GPUメモリ 11GB |
まとめ
今回は、C++のテンプレートライブラリであるThrustを紹介し、高速に処理できることを示した。インターフェースもSTLを意識したものになっており使いやすい。CUDA環境で高速化を実現する際の手っ取り早いツールである。ただし、行列演算ができないため、昨今流行りの深層学習には取り入れ難い。
