前回、単純な二値化処理を4通りの並列化手法で比較しました。
今回はCUDA関連でもう少し検証してみます。
CUDAのプログラミングで最初に気にするところと言えば、やはりcudaMalloc, cudaFreeのメモリ管理かと思います。出来ることならこの辺りは何も考えなくても勝手にやってくれるようにして欲しいものです。
というわけで今回はNVIDIAが開発している並列アルゴリズムライブラリであるThrustを使って前回の.cuコードを書き直してみました。
参考記事：
An Introduction to the Thrust Parallel Algorithms Library
いくつかとても気になる部分がありましたので、以下、そのメモになります。
書き直したコードは以下の通りです。
|
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 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 |
// Copyright SCIEMENT, Inc. // by Hirofumi Seo, M.D., CEO & President #include "CUDA_test.h" #include "thrust/device_vector.h" #include "thrust/copy.h" #include <chrono> #include <iostream> #include <stdio.h> __global__ void Kernel_make_bit_vertices(const int threshold, const int* voxels, int* bit_vertices, const int voxels_size) { int index = blockIdx.x * blockDim.x + threadIdx.x; const int stride = blockDim.x * gridDim.x; for (int i = index; i < voxels_size; i += stride) { bit_vertices[i] = (voxels[i] < threshold) ? 1 : 0; } } cudaError_t Cuda_make_bit_vertices(const int threshold, const std::vector<int>& voxels, std::vector<int>* bit_vertices) { cudaError_t cuda_status; auto start = std::chrono::system_clock::now(); // Choose which GPU to run on, change this on a multi-GPU system. cuda_status = cudaSetDevice(0); if (cuda_status != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); return cuda_status; } auto end = std::chrono::system_clock::now(); auto execution_time = end - start; std::cout << "Initialize: " << std::chrono::duration_cast<std::chrono::milliseconds>(execution_time).count() << " msec." << std::endl; const int voxels_size = voxels.size(); bit_vertices->resize(voxels_size); start = std::chrono::system_clock::now(); //thrust::device_vector<int> dev_voxels_vector(voxels); thrust::device_vector<int> dev_voxels_vector(voxels_size); thrust::device_vector<int> dev_bit_vertices_vector(voxels_size); end = std::chrono::system_clock::now(); execution_time = end - start; std::cout << "GPU Malloc: " << std::chrono::duration_cast<std::chrono::milliseconds>(execution_time).count() << " msec." << std::endl; //std::cout << "GPU Malloc and CPU -> GPU: " << std::chrono::duration_cast<std::chrono::milliseconds>(execution_time).count() << " msec." << std::endl; start = std::chrono::system_clock::now(); cudaMemcpy(thrust::raw_pointer_cast(dev_voxels_vector.data()), voxels.data(), voxels_size * sizeof(int), cudaMemcpyHostToDevice); end = std::chrono::system_clock::now(); execution_time = end - start; std::cout << "CPU -> GPU: " << std::chrono::duration_cast<std::chrono::milliseconds>(execution_time).count() << " msec." << std::endl; start = std::chrono::system_clock::now(); // Launch a kernel on the GPU with one thread for each element. const int block_size = 256; // MAX: 1024 const int num_blocks = (voxels_size + block_size - 1) / block_size; Kernel_make_bit_vertices <<<num_blocks, block_size >>>(threshold, thrust::raw_pointer_cast(dev_voxels_vector.data()), thrust::raw_pointer_cast(dev_bit_vertices_vector.data()), voxels_size); // Check for any errors launching the kernel cuda_status = cudaGetLastError(); if (cuda_status != cudaSuccess) { fprintf(stderr, "Kernel_make_bit_vertices launch failed: %s\n", cudaGetErrorString(cuda_status)); return cuda_status; } // cudaDeviceSynchronize waits for the kernel to finish, and returns // any errors encountered during the launch. cuda_status = cudaDeviceSynchronize(); if (cuda_status != cudaSuccess) { fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching Kernel_make_bit_vertices!\n", cuda_status); return cuda_status; } end = std::chrono::system_clock::now(); execution_time = end - start; std::cout << "GPU: " << std::chrono::duration_cast<std::chrono::microseconds>(execution_time).count() / 1000.0 << " msec." << std::endl; start = std::chrono::system_clock::now(); // Copy output vector from GPU buffer to host memory. //thrust::copy(dev_bit_vertices_vector.begin(), dev_bit_vertices_vector.end(), bit_vertices->begin()); cudaMemcpy(bit_vertices->data(), thrust::raw_pointer_cast(dev_bit_vertices_vector.data()), voxels_size * sizeof(int), cudaMemcpyDeviceToHost); end = std::chrono::system_clock::now(); execution_time = end - start; std::cout << "GPU -> CPU: " << std::chrono::duration_cast<std::chrono::milliseconds>(execution_time).count() << " msec." << std::endl; return cuda_status; } |
cudaMalloc, cudaFreeが不要になります。
最初は元々のコードの最後の部分にあったcudaFreeのところをコメントアウトして、
|
1 2 3 4 5 |
Error: //cudaFree(dev_voxels); //cudaFree(dev_bit_vertices); return cudaStatus; |
として
|
1 2 3 |
if (cuda_status != cudaSuccess) { goto Error; } |
のようにしていたのですが、何故かコンパイル時に
initialization of ‘dev_bit_vertices_vector’ is skipped by ‘goto Error’
というエラーが出てしまいました（この原因は全くわかりませんでした…）。
cudaFreeが必要なくなりましたので、Error:を消してreturnで返すようにしたらエラーは消えたのですが、実行時ではなくコンパイル時にこのエラーが出てしまった理由はわかりませんでした…。
せっかくThrustを使っているので、
・thrust::device_vector = std::vectorでのdevice_vectorの初期化
・thrust::copyでdevice_vectorからstd::vectorへのデータの書き込み
を行ってみたところ、信じられないくらい速度が遅くなりました…。
Thurstを使っているのに毎回raw pointerを取得してcudaMemcpyを使うほうが圧倒的にCPU, GPU間のデータ処理は速いようです。
以下、実行時間の比較です。出力結果に少し後付けで説明を加えています。
もちろん、数字はいじっていません。
|
CUDA（前回のもの）: Initialize: 170 msec. GPU Malloc: 99 msec. CPU -> GPU: 9 msec. GPU: 0.622 msec. GPU -> CPU: 9 msec. |
Thrust（非cudaMemcpy版）： Initialize: 140 msec. Thrust Malloc: 167 msec. + CPU -> GPU GPU: 0.605 msec. GPU -> CPU: 52 msec. |
Thrust（cudaMemcpy版）： Initialize: 139 msec. Thrust Malloc: 108 msec. CPU -> GPU: 9 msec. GPU: 0.583 msec. GPU -> CPU: 8 msec. |
データ転送時間が6倍近くも変わってしまっていますね…。なんとも残念な結果です。
とは言え、cudaMemcpyを使いさえすればThrust版でも元々のCUDA版と同じパフォーマンスが得られ、且つメモリ解放し忘れを気にしなくて良くなるのはとても大きいですね。
raw pointerを取得すれば既存のカーネル関数も使えるわけですし、さらにThrustの各種アルゴリズムを手軽に使えるのも嬉しい限りです。
※本記事内容は、国立研究開発法人 日本医療研究開発機構（AMED）の平成29年度 「未来医療を実現する医療機器・システム研究開発事業『術中の迅速な判断・決定を支援するための診断支援機器・システム開発』」採択課題である「術前と術中をつなぐスマート手術ガイドソフトウェアの開発」（代表機関名：東京大学、研究開発代表者名：齊藤延人）に、東京大学大学院情報理工学系研究科の学術支援専門職員として参画している瀬尾拡史が、研究開発として行っているものやその成果を一部含んでいます。