本記事はNVIDIAの協力に基づいて記事作成しております。
データ転送を速くする
前回のサンプル：1000×1000行列のかけ算の実行結果を再掲します。
簡単なサンプルとして行列積を取り上げましたけど、実アプリではこんなのわざわざ実装する必要はありません。基本的な線型代数ルーチンを集めたBLAS（Basic Linear Algebra Subprograms）のCUDA実装:cuBLASがCUDA Toolkitに含まれています。デバイス側コードはcuBLASライブラリの中にあるので、.cuを書くこともなく通常のC/C++ライブラリとしてリンクし、利用できます（cuBLASによる行列積をサンプルファイルに納めました）。
デバイス内の処理は20GFLOPS超のナイスなスピードを叩き出しています。が、ホスト-デバイス間のデータ転送を含めると……ちょっと残念。データ転送時間が足を引っ張っています。PCI-busを越えて大量のデータをやりとりするのはそれなりの時間がかかるものです。
ホスト-デバイス間のデータ転送はDMA（Direct Memory Access）を用い、CPUの動作とは独立してホスト側のメモリとビデオ・メモリとの間でコピーを行っています。ホスト側プログラム（プロセス）は仮想アドレス空間で動いており、プログラム内のポインタが指す値（アドレス）と実メモリ上の物理アドレスは一致しません。さらにOSはメモリの使用状況を監視していて、必要に応じてメモリの内容をディスクに保存し（ページアウト）、他の目的で使ったりします。
DMAにしてみればページアウト領域がコピー元／先では意図しない値がコピーされることとなり、困ってしまいます。そこで物理メモリ上に特定の固定領域を用意し、コピー元からいったんそちらに転送したのちにDMAがデバイスにコピーを行います（逆方向も同様）。CUDAではページアウトしない"ピン留め"された領域（page-lockedメモリ／pinnedメモリ）を確保するAPIが用意されています。ピン留めされた領域なら、そこから直接DMA転送できるってスンポーです。
page-lockedメモリの確保は簡単、cudaMemcpyによるホスト-デバイス間コピーの対象となるホスト側領域をcudaMallocHostで確保するだけ。
型* p = new 型[要素数];
や
型* p = (型*)malloc(要素数*sizeof(型));
を
型* p = nullptr; cudaMallocHost(&p, 要素数*sizeof(型));
に置き換え、delete[] p;やfree(p)はcudaFreeHost(p)に改めます。
早速実験。10メガbyteのデータをホスト-デバイス間で100往復させ、所要時間を測ります。
#include "cuda_runtime.h"
#include <iostream>
#include <chrono>
#include <cassert>
using namespace std;
chrono::milliseconds memory_test(char* host, char* device, size_t size, unsigned int repeat) {
chrono::high_resolution_clock::time_point start = chrono::high_resolution_clock::now();
while ( --repeat ) {
cudaMemcpy(device, host, size, cudaMemcpyHostToDevice);
cudaMemcpy(host, device, size, cudaMemcpyDeviceToHost);
}
chrono::high_resolution_clock::time_point stop = chrono::high_resolution_clock::now();
return chrono::duration_cast<chrono::milliseconds>(stop - start);
}
int main() {
const size_t size = 1024*1024*10; // 10M-bytes
const unsigned int repeat = 100; // 100 times
char* device;
cudaMalloc(&device, size);
chrono::milliseconds duration;
char* host = nullptr;
for ( int i = 0; i < 5; ++i ) {
cout << "--- normal -------- ";
host = new char[size];
assert( host != nullptr );
duration = memory_test(host, device, size, repeat);
long long nms = duration.count();
cout << duration.count() << " [ms]" << endl;
delete[] host;
host = nullptr;
cout << "--- page-locked --- ";
cudaError_t status = cudaMallocHost(&host, size);
assert( status == cudaSuccess );
duration = memory_test(host, device, size, repeat);
long long pms = duration.count();
cout << duration.count() << " [ms] (" << pms * 1.0f / nms << ")" << endl;
cudaFreeHost(host);
host = nullptr;
cout << endl;
}
return 0;
}
40％ほど時間短縮できているみたいです。この差は決して小さくないですよ。
ホスト-デバイス間のデータ転送はその量／回数ともにできるだけ少なくなるような実装設計がパフォーマンスのキモとなりそうです。いったんホストからデバイスにデータを投げたら、極力デバイス内で処理を完結し、最終結果のみをデバイスから引き取るのがベストでしょうね。
