并行编程实战——CUDA编程的异步拷贝

一、异步拷贝拷贝是一个非常容易理解的动作但是如果加上异步可能就需要看具体的场景下的实现了。对于CUDA来说异步拷贝Asynchronous Copy指的是在不阻塞某一方主机或Grid内的其它线程的情况下实现数据双方的数据传输机制。它与流水线中的重叠作业有些类似可以有效的隐藏数据传输的延迟进而提高硬件的利用率特别是GPU。对于CUDA编程来说主机和GPU间异步拷贝的基础主要有两个技术流流内的操作串行但不同流操作可以并行固定内存主机端分配的内存被锁定而无法交换到硬盘GPU可以通过DMA技术直接访问它是异步拷贝的前提。要想高效物实现异步拷贝就必须根据实际情况来选择相应的数据拷贝方式否则很可能是一种伪异步。而对于网格内的异步拷贝基础技术也有两个Pipeline和异步屏障。前者用于全局内存与共享内存间的异步拷贝允许线程在数据加载时继续进行其它计算任务而后者用于协调数据就绪与消费的顺序。二、CUDA中的API在CUDA的编程框架中提供了以下相关的API主机与设备端的异步拷贝cudaMemcpyAsync()核心API。它既可以处理主机到设备也可以支持设备间的异步拷贝。前提是目标内存必须是固定内存cudaMemcpy2DAsync()、cudaMemcpy3DAsync()和cudaMemcpyPeerAsync()分别用于处理矩阵、3D数据和多GPU间的异步传输cudaMemcpyFrom/ToArrayAsync()用于CUDA数组cudaArray的异步拷贝接口cudaMemcpyFrom/ToSymbolAsync()读写设备端全局符号的异步拷贝接口设备端网格内的异步拷贝cuda::memcpy_async与cuda::barrier或cuda::pipeline共同使用组织协调异步数据流cooperative_groups::memcpy_async配合协同组Cooperative Groups使用除了上面的异步拷贝API外还有一些其它的API用于辅助支持异步拷贝如刚刚提到的cuda::barrier或cuda::pipeline等接口以及cudaEvent等相关的异步接口。对于CUDA来说底层的指令也提供了一些异步拷贝的支持如LDGSTS、TMA (Tensor Memory Accelerator)和STAS等。大家都明白的硬件支持往往意味着比软件的效率更高。三、应用场景在并行编程中异步拷贝的主要目的是为了提高效率。而异步拷贝的应用场景就可以从它的目的出发一般可见于以下常用场景主机到设备端的大数据量拷贝且核函数的执行时间少于数据传输所需时间。比如深度学习和科学计算的大数据的预处理或结果的异步回传GPU全局内存与共享内存间的异步拷贝。它比较适合于访问密集型的情况特别是数据可按块划分的场景。比如图像的卷积操作或矩阵的分块计算等多GPU的点对点传输。这个属于比较大的应用了多个GPU间的异步拷贝。常见的就是分布式训练或多卡模型训练TMATensor Memory Accelerator高维异步拷贝。在计算力9以上的架构上TMA支持多维张量的异步拷贝。四、应用下面看一个简单的例程#includecuda_runtime.h#includedevice_launch_parameters.h#includestdio.h#includeiostream#includevector#includecmath// Kernel: add two vectors__global__voidvecAdd(constfloat*a,constfloat*b,float*c,intn){intidxblockIdx.x*blockDim.xthreadIdx.x;if(idxn){c[idx]a[idx]b[idx];}}intmain(){constintn120;// total 1M elementsconstintblockSize256;constintnumStreams4;// use 4 CUDA streamssize_tbytesn*sizeof(float);size_tchunkBytesbytes/numStreams;// chunk size per stream// Allocate pinned host memory for asynchronous copiesfloat*hA,*hB,*hC;cudaHostAlloc(hA,bytes,cudaHostAllocDefault);cudaHostAlloc(hB,bytes,cudaHostAllocDefault);cudaHostAlloc(hC,bytes,cudaHostAllocDefault);// Initialize host datafor(inti0;in;i){hA[i]1.0f;hB[i]2.0f;}// Allocate device memoryfloat*dA,*dB,*dC;cudaMalloc(dA,bytes);cudaMalloc(dB,bytes);cudaMalloc(dC,bytes);// Create streamscudaStream_tstreams[numStreams];for(inti0;inumStreams;i){cudaStreamCreate(streams[i]);}// Pipeline execution: each stream processes one chunkfor(inti0;inumStreams;i){size_toffseti*chunkBytes/sizeof(float);// element offsetsize_tchunkSizechunkBytes;// Async copy: host - device (H2D)cudaMemcpyAsync(dA[offset],hA[offset],chunkSize,cudaMemcpyHostToDevice,streams[i]);cudaMemcpyAsync(dB[offset],hB[offset],chunkSize,cudaMemcpyHostToDevice,streams[i]);// Kernel compute (can overlap with later D2H)intthreadsblockSize;intblocks(chunkSize/sizeof(float)threads-1)/threads;vecAddblocks,threads,0,streams[i](dA[offset],dB[offset],dC[offset],chunkSize/sizeof(float));// Async copy: device - host (D2H)cudaMemcpyAsync(hC[offset],dC[offset],chunkSize,cudaMemcpyDeviceToHost,streams[i]);}// Wait for all streams to completefor(inti0;inumStreams;i){cudaStreamSynchronize(streams[i]);}// Verify resultsbool correcttrue;for(inti0;in;i){if(std::fabs(hC[i]-3.0f)1e-5){correctfalse;break;}}std::coutrun result: (correct?OK:NO)std::endl;// Clean up resourcesfor(inti0;inumStreams;i){cudaStreamDestroy(streams[i]);}cudaFree(dA);cudaFree(dB);cudaFree(dC);cudaFreeHost(hA);cudaFreeHost(hB);cudaFreeHost(hC);return0;}代码很简单上机运行即可知道整体的流程。五、总结CUDA中的异步与CPU主机编程中的异步其实本质是相似的。所以是不是可以理解二者对内存的优化方向很有可能也是相似的。因此触类旁通的将二者进行对比分析学习很可能让学习的效率提高不少。