企业官网建设流程全解析

一、异步拷贝

拷贝是一个非常容易理解的动作，但是如果加上异步，可能就需要看具体的场景下的实现了。对于CUDA来说，异步拷贝，Asynchronous Copy，指的是在不阻塞某一方（主机或Grid内的其它线程）的情况下，实现数据双方的数据传输机制。它与流水线中的重叠作业有些类似，可以有效的隐藏数据传输的延迟进而提高硬件的利用率（特别是GPU）。
对于CUDA编程来说，主机和GPU间异步拷贝的基础主要有两个技术：流，流内的操作串行，但不同流操作可以并行；固定内存，主机端分配的内存被锁定而无法交换到硬盘，GPU可以通过DMA技术直接访问，它是异步拷贝的前提。要想高效物实现异步拷贝，就必须根据实际情况来选择相应的数据拷贝方式，否则很可能是一种伪异步。
而对于网格内的异步拷贝基础技术也有两个：Pipeline和异步屏障。前者用于全局内存与共享内存间的异步拷贝（允许线程在数据加载时继续进行其它计算任务）而后者用于协调数据就绪与消费的顺序。

二、CUDA中的API

在CUDA的编程框架中，提供了以下相关的API：

1. 主机与设备端的异步拷贝
   cudaMemcpyAsync()：核心API。它既可以处理主机到设备也可以支持设备间的异步拷贝。前提是目标内存必须是固定内存
   cudaMemcpy2DAsync()、cudaMemcpy3DAsync()和cudaMemcpyPeerAsync()：分别用于处理矩阵、3D数据和多GPU间的异步传输
   cudaMemcpyFrom/ToArrayAsync()：用于CUDA数组（cudaArray）的异步拷贝接口
   cudaMemcpyFrom/ToSymbolAsync()：读写设备端全局符号的异步拷贝接口
2. 设备端网格内的异步拷贝
   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等。大家都明白的，硬件支持往往意味着比软件的效率更高。

三、应用场景

在并行编程中，异步拷贝的主要目的是为了提高效率。而异步拷贝的应用场景就可以从它的目的出发，一般可见于以下常用场景：

1. 主机到设备端的大数据量拷贝且核函数的执行时间少于数据传输所需时间。比如深度学习和科学计算的大数据的预处理或结果的异步回传
2. GPU全局内存与共享内存间的异步拷贝。它比较适合于访问密集型的情况，特别是数据可按块划分的场景。比如图像的卷积操作或矩阵的分块计算等
3. 多GPU的点对点传输。这个属于比较大的应用了，多个GPU间的异步拷贝。常见的就是分布式训练或多卡模型训练
4. TMA（Tensor Memory Accelerator）高维异步拷贝。在计算力9以上的架构上，TMA支持多维张量的异步拷贝。

四、应用

下面看一个简单的例程：

#include"cuda_runtime.h"#include"device_launch_parameters.h"#include<stdio.h>#include<iostream>#include<vector>#include<cmath>// Kernel: add two vectors__global__voidvecAdd(constfloat*a,constfloat*b,float*c,intn){intidx=blockIdx.x*blockDim.x+threadIdx.x;if(idx<n){c[idx]=a[idx]+b[idx];}}intmain(){constintn=1<<20;// total 1M elementsconstintblockSize=256;constintnumStreams=4;// use 4 CUDA streamssize_tbytes=n*sizeof(float);size_tchunkBytes=bytes/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(inti=0;i<n;++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(inti=0;i<numStreams;++i){cudaStreamCreate(&streams[i]);}// Pipeline execution: each stream processes one chunkfor(inti=0;i<numStreams;++i){size_toffset=i*chunkBytes/sizeof(float);// element offsetsize_tchunkSize=chunkBytes;// 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)intthreads=blockSize;intblocks=(chunkSize/sizeof(float)+threads-1)/threads;vecAdd<<<blocks,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(inti=0;i<numStreams;++i){cudaStreamSynchronize(streams[i]);}// Verify resultsbool correct=true;for(inti=0;i<n;++i){if(std::fabs(hC[i]-3.0f)>1e-5){correct=false;break;}}std::cout<<"run result: "<<(correct?"OK":"NO")<<std::endl;// Clean up resourcesfor(inti=0;i<numStreams;++i){cudaStreamDestroy(streams[i]);}cudaFree(dA);cudaFree(dB);cudaFree(dC);cudaFreeHost(hA);cudaFreeHost(hB);cudaFreeHost(hC);return0;}

代码很简单，上机运行即可知道整体的流程。

五、总结

CUDA中的异步与CPU主机编程中的异步其实本质是相似的。所以是不是可以理解，二者对内存的优化方向很有可能也是相似的。因此，触类旁通的将二者进行对比分析学习，很可能让学习的效率提高不少。