我目前正在通过计算两个向量之间的点积来学习CUDA流。这些成分是一个核函数,它接受向量x和y,并返回一个大小等于块数的向量,结果,其中每个块贡献自己的减少的总和。CUDA流的性能
我也有一台主机功能dot_gpu调用内核,并降低了矢量结果到最终点积值。
同步版本不只是这一点:
// copy to device
copy_to_device<double>(x_h, x_d, n);
copy_to_device<double>(y_h, y_d, n);
// kernel
double result = dot_gpu(x_d, y_d, n, blockNum, blockSize);
而异步一个是这样:
double result[numChunks];
for (int i = 0; i < numChunks; i++) {
int offset = i * chunkSize;
// copy to device
copy_to_device_async<double>(x_h+offset, x_d+offset, chunkSize, stream[i]);
copy_to_device_async<double>(y_h+offset, y_d+offset, chunkSize, stream[i]);
// kernel
result[i] = dot_gpu(x_d+offset, y_d+offset, chunkSize, blockNum, blockSize, stream[i]);
}
for (int i = 0; i < numChunks; i++) {
finalResult += result[i];
cudaStreamDestroy(stream[i]);
}
我使用的数据流,并试图探讨的原因时变得更差的性能。我试图管理下载,内核调用和上传,但没有结果。
// accumulate the result of each block into a single value
double dot_gpu(const double *x, const double* y, int n, int blockNum, int blockSize, cudaStream_t stream=NULL)
{
double* result = malloc_device<double>(blockNum);
dot_gpu_kernel<<<blockNum, blockSize, blockSize * sizeof(double), stream>>>(x, y, result, n);
#if ASYNC
double* r = malloc_host_pinned<double>(blockNum);
copy_to_host_async<double>(result, r, blockNum, stream);
CudaEvent copyResult;
copyResult.record(stream);
copyResult.wait();
#else
double* r = malloc_host<double>(blockNum);
copy_to_host<double>(result, r, blockNum);
#endif
double dotProduct = 0.0;
for (int i = 0; i < blockNum; i ++) {
dotProduct += r[i];
}
cudaFree(result);
#if ASYNC
cudaFreeHost(r);
#else
free(r);
#endif
return dotProduct;
}
我的猜测是,这个问题是dot_gpu()功能,它不仅调用内核里面。告诉我,如果我理解正确以下流的执行
foreach stream {
cudaMemcpyAsync(device[stream], host[stream], ... stream);
LaunchKernel<<<...stream>>>(...);
cudaMemcpyAsync(host[stream], device[stream], ... stream);
}
主机执行,因为马上(cudaMemcpyAsync和内核回报所有三个指令不会被阻挡,但他们会按顺序执行的GPU,因为他们被分配到相同的流)。因此,主机继续下一个流(即使stream1谁知道它在哪个阶段,但是谁在乎......它在GPU上完成他的工作,对吧?)并且再次执行三条指令而不被阻塞......等等等等。但是,我的代码会阻止主机,然后才能处理下一个流,位于dot_gpu()函数内。是否因为我分配了空闲内容,并将内核返回的数组减少为单个值?