关于异步数据传输的使用方法与性能分析
提供异步执行的函数主要有以下几类:
Ø Kernel函数,通过配置流参数实现
Ø 以Async结尾的函数
Ø Device to device内存拷贝函数
Ø 存储器初始化函数,如cudaMemset等
比较适合用于数据量很大的情况,让处理时间与数据拷贝时间重叠。
异步数据传输使用方法:第一步创建流,
cudaStream_t *stream=(cudaStream_t*)calloc(streamNo,sizeof(cudaStream_t));
for (int i = 0; i < streamNo; i++)
cudaStreamCreate(&stream[i]);
然后使用异步函数和设置流参数:
for (int i = 0; i < streamNo; i++)
{ cudaMemcpy2DAsync(d_Adata+i*size,fpitch,h_fddata+i*size,sample*depth*sizeof(float),sample*depth*sizeof(float),line,cudaMemcpyHostToDevice,stream[i]); cudaMemcpy2DAsync(d_Bdata+i*size,fpitch,h_fddata+i*size,sample*depth*sizeof(float),sample*depth*sizeof(float),line,cudaMemcpyHostToDevice,stream[i]);
Mul<<<blocks,threads,0,stream[i]>>>(d_Cdata,d_Adata,d_Bdata,sample*depth,line,fpitch);
}
性能分析:
Ø 在使用分页内存时,异步处理方式明显优于同步方式,大略快4倍以上且性能波动很小(测试平台9800GT)。GPU上的计算资源是有限的,让数据处理与数据传输的时间产生重叠是一种不错的选择。
Ø 页锁定内存操作性能反而是同步优于异步,我的一种解释是页锁定内存传输的带宽三倍于分页内存,数据分块传输的代价大于处理与传输时间重叠换来的收益,只是个人测试结果,未必正确!