データ転送とカーネル実行をオーバーラップさせようとすると、使用するストリームに関係なく、カードがすべてのメモリ転送を順番に実行しているように見えます。
したがって、次を発行すると:
- ストリーム 1: MemcpyA_HtoD_1; カーネル_1; MemcpyA_DtoH_1
- ストリーム 2: MemcpyA_HtoD_2; カーネル_2; MemcpyA_DtoH_2
MemcpyA_HtoD_2 は MemcpyA_DtoH_1 が完了するまで待機します。したがって、オーバーラップは達成されません。使用するストリームの構成に関係なく、Memcpy 操作は常に順番に発行されます。したがって、オーバーラップを実現する唯一の方法は、出力をバッファリングするか、次の反復まで出力転送を遅らせることです。
CUDA 5.5、Windows 7 x64、および GTX Titan を使用しています。すべての CPU メモリが固定され、data_transfers は非同期バージョンを使用して行われます。
次の画面で動作を確認してください。
発行、host_to_device -> kernel -> device_to_host (通常の動作) となり、オーバーラップを取得できません。
host_to_device -> kernel (kernel の後に device_to_host を回避) を発行するとオーバーラップが発生します...どのようなストリーム構成を試しても、すべてのメモリ コピーが順番に実行されるためです。
アップデート
誰かがこの問題を再現することに興味がある場合は、この望ましくない動作を示す合成プログラムをコーディングしました. CUDA 5.5 を使用した完全な VS2010 ソリューション
オーバーラップをテストするために誰かがこれを Linux で実行できますか?
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#define N 1024*1024
__global__ void someKernel(int *d_in, int *d_out) {
for (int i = threadIdx.x; i < threadIdx.x + 1024; i++) {
d_out[i] = d_in[i];
}
}
int main () {
int *h_bufferIn[100];
int *h_bufferOut[100];
int *d_bufferIn[100];
int *d_bufferOut[100];
//allocate some memory
for (int i = 0; i < 100; i++) {
cudaMallocHost(&h_bufferIn[i],N*sizeof(int));
cudaMallocHost(&h_bufferOut[i],N*sizeof(int));
cudaMalloc(&d_bufferIn[i], N*sizeof(int));
cudaMalloc(&d_bufferOut[i], N*sizeof(int));
}
//create cuda streams
cudaStream_t st[2];
cudaStreamCreate(&st[0]);
cudaStreamCreate(&st[1]);
//trying to overlap computation and memcpys
for (int i = 0; i < 100; i+=2) {
cudaMemcpyAsync(d_bufferIn[i], h_bufferIn[i], N*sizeof(int), cudaMemcpyHostToDevice, st[i%2]);
someKernel<<<1,256, 0, st[i%2]>>>(d_bufferIn[i], d_bufferOut[i]);
cudaMemcpyAsync(h_bufferOut[i], d_bufferOut[i], N*sizeof(int), cudaMemcpyDeviceToHost, st[i%2]);
cudaStreamQuery(0);
cudaMemcpyAsync(d_bufferIn[i+1], h_bufferIn[i+1], N*sizeof(int), cudaMemcpyHostToDevice, st[(i+1)%2]);
someKernel<<<1,256, 0, st[(i+1)%2]>>>(d_bufferIn[i+1], d_bufferOut[i+1]);
cudaMemcpyAsync(h_bufferOut[i+1], d_bufferOut[i+1], N*sizeof(int), cudaMemcpyDeviceToHost, st[(i+1)%2]);
cudaStreamQuery(0);
}
cudaDeviceSynchronize();
}