streamで速くする(2)
streamはその名の通り"一連の(処理の)流れ"です。 streamに乗せられたjob(仕事)は並んだ順に実行されます。 複数のstreamにjobを乗せておけば、GPUはその中から実行可能なものを一つ選んで処理します。 ホスト・メモリ間のデータ転送にGPUは関与しないので、データ転送の間でも他のstreamに乗ったjobを処理できます。
ここまでにいくつかサンプル・コードを紹介してきましたが、それらの中には stream を明示的に使ったものはありません。 けども実は暗黙の default stream が一本だけあって、その default stream に cudaMemcpy
やら kernel呼び出しやらのjobを乗っけていたんです。
3本のstream(ssin, scos, stan)を用意し、sin,cos,tanの表作成をそれぞれの stream に乗せて処理しましょう。
#include <cuda_runtime.h> #include <device_launch_parameters.h> #include <math_functions.h> // 時間稼ぎ __device__ float time_waste() { float result = 0.0f; const int TIMES = 10U; for ( int i = 0; i < TIMES; ++i ) { result += 1.0f / TIMES; } return result; } // 度→ラジアン __device__ __forceinline__ float deg2rad(float angle) { return angle / 180.0f * 3.01416f * time_waste(); } // sine __global__ void kernel_sin(float* out, const float* angle, unsigned int size) { unsigned int i = blockDim.x * blockIdx.x + threadIdx.x; if ( i < size ) { out[i] = sinf(deg2rad(angle[i])); } } // cosine __global__ void kernel_cos(float* out, const float* angle, unsigned int size) { unsigned int i = blockDim.x * blockIdx.x + threadIdx.x; if ( i < size ) { out[i] = cosf(deg2rad(angle[i])); } } // tangent __global__ void kernel_tan(float* out, const float* angle, unsigned int size) { unsigned int i = blockDim.x * blockIdx.x + threadIdx.x; if ( i < size ) { out[i] = tanf(deg2rad(angle[i])); } } // helper funcs template<typename T> inline T* allocate_host(size_t size) { T* p = nullptr; cudaMallocHost(&p, size*sizeof(T)); return p; } template<typename T> inline void free_host(T* host) { cudaFreeHost(host); } template<typename T> inline T* allocate_device(size_t size) { T* p = nullptr; cudaMalloc(&p, size*sizeof(T)); return p; } template<typename T> inline void free_device(T* device) { cudaFree(device); } inline cudaEvent_t create_event() { cudaEvent_t e = nullptr; cudaEventCreate(&e); return e; }; inline void destroy_event(cudaEvent_t event) { cudaEventDestroy(event); } inline cudaStream_t create_stream() { cudaStream_t s = nullptr; cudaStreamCreate(&s); return s; }; inline void destroy_stream(cudaStream_t stream) { cudaStreamDestroy(stream); } #include <iostream> using namespace std; const unsigned int N = 1000U * 1000U; const unsigned int BN = N*sizeof(float); void multiple_stream() { cout << "--- multiple stream ---" << endl; float* angle = allocate_host<float>(N); // in float* sin = allocate_host<float>(N); // out float* cos = allocate_host<float>(N); // out float* tan = allocate_host<float>(N); // out for ( unsigned int i = 0; i < N; ++i ) { angle[i] = 360.0f * i / N; } float* dangle = allocate_device<float>(N); float* dsin = allocate_device<float>(N); float* dcos = allocate_device<float>(N); float* dtan = allocate_device<float>(N); unsigned int size = N; unsigned int block = 128U; unsigned int grid = (size + block -1U) / block; cudaEvent_t start = create_event(); cudaEvent_t mid = create_event(); cudaEvent_t stop = create_event(); cudaStream_t ssin = create_stream(); cudaStream_t scos = create_stream(); cudaStream_t stan = create_stream(); cudaEventRecord(start); cudaMemcpy(dangle, angle, BN, cudaMemcpyHostToDevice); cudaEventRecord(mid); kernel_sin<<<grid,block,0,ssin>>>(dsin, dangle, size); cudaMemcpyAsync(sin, dsin, BN, cudaMemcpyDeviceToHost, ssin); kernel_cos<<<grid,block,0,scos>>>(dcos, dangle, size); cudaMemcpyAsync(cos, dcos, BN, cudaMemcpyDeviceToHost, scos); kernel_tan<<<grid,block,0,stan>>>(dtan, dangle, size); cudaMemcpyAsync(tan, dtan, BN, cudaMemcpyDeviceToHost, stan); cudaEventRecord(stop); cudaDeviceSynchronize(); float elapsed; cudaEventElapsedTime(&elapsed, start, mid); cout << elapsed << " + "; cudaEventElapsedTime(&elapsed, mid, stop); cout << elapsed << " = "; cudaEventElapsedTime(&elapsed, start, stop); cout << elapsed << " [ms]" << endl; for ( unsigned int i = 0; i < 5; ++i ) { cout << angle[i] << " :" << " sin=" << sin[i] << " cos=" << cos[i] << " tan=" << tan[i] << endl; } cout << "..." << endl; for ( auto event : { start, mid, stop }) destroy_event(event); for ( auto stream : { ssin, scos, stan }) destroy_stream(stream); for ( auto host : { angle, sin, cos, tan }) free_host(host); for ( auto device : { dangle, dsin, dcos, dtan }) free_device(device); } int main() { multiple_stream(); }
default stream 1本で処理した時のタイムラインがコレ。
対して3本(+default)のstreamでやったのがコチラ。
ね、データ転送とkernel実行が同時に行われてるでしょ。重なった分だけ速くなるんです。
実装上のキモはふたつ:
cudaMemcpy
じゃなくcudaMemcpyAsync
を使うべし。 後者はコピー完了を待たずに(job投入後ただちに)返って来るので、すぐさま次のjobを投入できます。ホスト側メモリを
malloc
/free
やnew
/delete
じゃなくcudaMallocHost
/cudaFreeHost
で確保/解放すべし。 そうしないとデータ転送中に kernelが動いてくれないのですよ。
... streamのおハナシはまだまだ続きます。
(original: 2015-04-16)