CUDA程序优化之数据传输
一、设备端和主机端的数据相互拷贝
设备端指GPU端，数据存放在显存中；主机端指CPU，数据存放在内存中。一般情况下，一般情况下设备端是不能直接访问主机端内存的，而我们的数据通常情况下都是存放在主机端内存中，要在GPU中执行算法运算就必须先把数据拷贝至设备端，运算完成再把结果拷回至主机端。这个传输过程，显然是会耗时的。
传输需要多少耗时？ 这和PCIe总线带宽正相关。PCIe是CPU和GPU之间数据传输的接口，发展至今有多代技术，从之前的PCIe 1.0到现在的PCIe 3.0、PCIe 4.0，带宽越来越大，传输也是越来越快。一般PCIe会有多条Lane并行传输，理论传输速度成倍增加，我这里列一下多路PCIe 3.0、PCIe 4.0各自的带宽数值：
我用GPU-Z查了查我的MX150，显示是PCIe x4 3.0，对应上表中的400MB/s的带宽。
我们可以通过总线带宽来计算数据传输耗时，以一张1280x960的灰度图像为例，1个像素占1个字节，则传输数据量为 1280x960x1 B = 1228800 B = 1200 KB = 1.172 MB。若用我的MX150，则传输耗时 t = 1.172/4000 s ≈ 0.29 ms。看起来很少对不对，但我们算的可是理论峰值带宽，你见过有几个产品能到理论峰值的呢？最后的时间基本是要打较大折扣的，时间估计在0.35ms左右，你可能还是觉得很少，但是如果你传的是彩色图（一个像素3个字节）呢？要是一次需要传两张图呢？t = 0.35 x 3 x 2 = 2.1 ms，对于GPU算法来说，这个时间就不该被忽视了。
二、不同的内存分配/传输方式，传输效率有何不同？
（1）常规方式传输：cudaMemcpy
在CUDA中常规的传输接口是cudaMemcpy，我想这也是被使用最多的接口，他可以将数据从主机端拷贝至设备端，也可以从设备端拷贝至主机端，函数声明如下：
__host__ cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )
cudaMemcpyKind决定拷贝的方向，有以下取值：
cudaMemcpyHostToHost = 0
Host -> Host
cudaMemcpyHostToDevice = 1
Host -> Device
cudaMemcpyDeviceToHost = 2
Device -> Host
cudaMemcpyDeviceToDevice = 3
Device -> Device
cudaMemcpyDefault = 4
Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing
该方式使用非常简单，很多情况下效率也足以满足性能需求。
（2）高维矩阵传输：cudaMemcpy2D/cudaMalloc3D
顾名思义，cudaMemcpy2D/cudaMalloc3D是应对2D及3D矩阵数据的。以图像为例，我们可以用cudaMalloc来分配一维数组来存储一张图像数据，但这不是效率最快的方案，推荐的方式是使用cudaMallocPitch来分配一个二维数组来存储图像数据，存取效率更快。
__host__ cudaError_t cudaMallocPitch ( void** devPtr, size_t* pitch, size_t width, size_t height )
相比于cudaMemcpy2D对了两个参数dpitch和spitch，他们是每一行的实际字节数，是对齐分配cudaMallocPitch返回的值。
__host__ cudaError_t cudaMemcpy2D ( void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, cudaMemcpyKind kind )
cudaMallocPitch有一个非常好的特性是二维矩阵的每一行是内存对齐的，访问效率比一维数组更高。而通过cudaMallocPitch分配的内存必须配套使用cudaMemcpy2D完成数据传输。C 中二维数组内存分配是转化为一维数组，连贯紧凑，每次访问数组中的元素都必须从数组首元素开始遍历；而 cuda 中这样分配的二维数组内存保证了数组每一行首元素的地址值都按照 256 或 512 的倍数对齐，提高访问效率，但使得每行末尾元素与下一行首元素地址可能不连贯，使用指针寻址时要注意考虑尾部。
- cudaMAllocPitch() 传入存储器指针 **devPtr，偏移值的指针 *pitch，数组行字节数 widthByte，数组行数 height。函数返回后指针指向分配的内存（每行地址对齐到 AlignByte 字节，为 256B 或 512B），偏移值指针指向的值为该行实际字节数（= sizeof(datatype) * width + alignByte - 1) / alignByte）。
- cudaMemcpy2D() 传入目标存储器的指针 *dst，目标存储器行字节数 dpitch，源存储器指针 *src，源存储器行字节数 spitch，数组行字节数 width，数组行数 height，拷贝方向 kind。这里要求存储器行字节数不小于数组行字节数，多出来的部分就是每行尾部空白部分。
- 整个测试代码。
并非说cudaMemcpy2D/cudaMemcpy3D比cudaMemcpy传输更快，而是对齐内存必须使用cudaMemcpy2D/cudaMemcpy3D来配套使用。
#include <stdio.h>
#include <malloc.h>
#include <cuda_runtime_api.h>
#include "device_launch_parameters.h"
__global__ void myKernel(float* devPtr, int height, int width, int pitch)
{
int row, col;
float *rowHead;
for (row = 0; row < height; row++)
{
rowHead = (float*)((char*)devPtr + row * pitch);
for (col = 0; col < width; col++)
{
printf("\t%f", rowHead[col]);// 逐个打印并自增 1
rowHead[col]++;
}
printf("\n");
}
}
int main()
{
size_t width = 6;
size_t height = 5;
float *h_data, *d_data;
size_t pitch;
h_data = (float *)malloc(sizeof(float)*width*height);
for (int i = 0; i < width*height; i++)
h_data[i] = (float)i;
printf("\n\tAlloc memory.");
cudaMallocPitch((void **)&d_data, &pitch, sizeof(float)*width, height);
printf("\n\tPitch = %d B\n", pitch);
printf("\n\tCopy to Device.\n");
cudaMemcpy2D(d_data, pitch, h_data, sizeof(float)*width, sizeof(float)*width, height, cudaMemcpyHostToDevice);
myKernel << <1, 1 >> > (d_data, height, width, pitch);
cudaDeviceSynchronize();
printf("\n\tCopy back to Host.\n");
cudaMemcpy2D(h_data, sizeof(float)*width, d_data, pitch, sizeof(float)*width, height, cudaMemcpyDeviceToHost);
for (int i = 0; i < width*height; i++)
{
printf("\t%f", h_data[i]);
if ((i + 1) % width == 0)
printf("\n");
}
free(h_data);
cudaFree(d_data);
getchar();
return 0;
}
结果：
3D矩阵的配套API为：
__host__ cudaError_t cudaMalloc3D ( cudaPitchedPtr* pitchedDevPtr, cudaExtent extent )
__host__ cudaError_t cudaMemcpy3D ( const cudaMemcpy3DParms* p )
代码示例如下:
#include <stdio.h>
#include <malloc.h>
#include <cuda_runtime_api.h>
#include "device_launch_parameters.h"
#include <driver_functions.h>
__global__ void myKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent)
{
float * devPtr = (float *)devPitchedPtr.ptr;
float *sliceHead, *rowHead;
// 可以定义为 char * 作面、行迁移的时候直接加减字节数，取行内元素的时候再换回 float *
for (int z = 0; z < extent.depth; z++)
{
sliceHead = (float *)((char *)devPtr + z * devPitchedPtr.pitch * extent.height);
for (int y = 0; y < extent.height; y++)
{
rowHead = (float*)((char *)sliceHead + y * devPitchedPtr.pitch);
for (int x = 0; x < extent.width / sizeof(float); x++)// extent 存储的是行有效字节数，要除以元素大小
{
printf("\t%f",rowHead[x]);// 逐个打印并自增 1
rowHead[x]++;
}
printf("\n");
}
printf("\n");
}
}
int main()
{
size_t width = 2;
size_t height = 3;
size_t depth = 4;
float *h_data;
cudaPitchedPtr d_data;
cudaExtent extent;
cudaMemcpy3DParms cpyParm;
h_data = (float *)malloc(sizeof(float) * width * height * depth);
for (int i = 0; i < width * height * depth; i++)
h_data[i] = (float)i;
printf("\n\tAlloc memory.");
extent = make_cudaExtent(sizeof(float) * width, height, depth);
cudaMalloc3D(&d_data, extent);
printf("\n\tCopy to Device.\n");
cpyParm = {0};
cpyParm.srcPtr = make_cudaPitchedPtr((void*)h_data, sizeof(float) * width, width, height);
cpyParm.dstPtr = d_data;
cpyParm.extent = extent;
cpyParm.kind = cudaMemcpyHostToDevice;
cudaMemcpy3D(&cpyParm);
myKernel << <1, 1 >> > (d_data, extent);
cudaDeviceSynchronize();
printf("\n\tCopy back to Host.\n");
cpyParm = { 0 };
cpyParm.srcPtr = d_data;
cpyParm.dstPtr = make_cudaPitchedPtr((void*)h_data, sizeof(float) * width, width, height);
cpyParm.extent = extent;
cpyParm.kind = cudaMemcpyDeviceToHost;
cudaMemcpy3D(&cpyParm);
for (int i = 0; i < width*height*depth; i++)
{
printf("\t%f", h_data[i]);
if ((i + 1) % width == 0)
printf("\n");
if ((i + 1) % (width*height) == 0)
printf("\n");
}
free(h_data);
cudaFree(d_data.ptr);
getchar();
return 0;
}
结果如下：
（3）异步传输：cudaMemcpyAsync / cudaMemcpy2DAsync / cudaMemcpy3DAsync
我们知道传输是走PCIe总线的，计算和PCIe总线里的数据流通完全独立，那么某些情况下，我们可以让计算和传输异步进行，而不是等数据传输完再做计算。
举个例子：我必须一次传入两张图像，做处理运算。常规操作是使用cudaMemcpy或者cudaMemcpy2D把两张图像都传输到显存，再启动kernel运算。传输和运算是串行的，运算必须等待传输完成。
而cudaMemcpyAsync / cudaMemcpy2DAsync / cudaMemcpy3DAsync 可以让传输和运算之间异步并行。上面的例子，如果用cudaMemcpyAsync或cudaMemcpy2DAsync，可以先传输第一张影像到显存，然后启动第一张影像的运算kernel，同时启动第二张影像的传输，此时第一张影像的运算和第二张影像的传输就是异步进行的，互相独立，便可隐藏掉第二张影像的传输耗时。
三个异步传输接口如下：
__host__ __device__ cudaError_t cudaMemsetAsync ( void* devPtr, int value, size_t count, cudaStream_t stream = 0 )
__host__ __device__ cudaError_t cudaMemcpy2DAsync ( void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, cudaMemcpyKind kind, cudaStream_t stream = 0 )
__host__ __device__ cudaError_t cudaMemcpy3DAsync ( const cudaMemcpy3DParms* p, cudaStream_t stream = 0 )
异步传输是非常实用的，当你一次处理多个数据时，可以考虑是否可以用异步传输来隐藏一部分传输耗时。
（4）锁页内存（Page-locked）
锁页内存是在主机端上的内存。主机端常规方式分配的内存（用new、malloc等方式）都是可分页（pageable）的，操作系统可以将可分页内存和虚拟内存（硬盘上的一块空间）相互交换，以获得比实际内存容量更大的内存使用。
可分页内存在分配后是可能被操作系统移动的，GPU端无法获知操作系统是否正在移动对可分页内存，所以不可让GPU端直接访问。实际的情况是，当从可分页内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的主机内存，将可分页内存复制到页面锁定内存中 [copy 1]，然后再从页面锁定内存传输到设备内存 [copy 2]。显然，这里面有两次传输。
所以我们能否直接分配页面锁定的内存？让GPU端直接访问，让传输只有一次！
答案是肯定的，我们可以在主机端分配锁页内存。锁页内存是主机端一块固定的物理内存，它不能被操作系统移动，不参与虚拟内存相关的交换操作。简而言之，分配之后，地址就固定了，被释放之前不会再变化。
GPU知道锁页内存的物理地址，可以通过“直接内存访问（Direct Memory Access，DMA）”技术直接在主机和GPU之间复制数据，传输仅一次，效率更高。
CUDA提供两种方式在主机端分配锁页内存
1. cudaMallocHost
__host__ cudaError_t cudaMallocHost ( void** ptr, size_t size )
2. cudaHostAlloc
pHost为分配的锁页内存地址，size为分配的字节数，flags为内存分配类型，取值如下：
-
cudaHostAllocDefault 默认值，等同于cudaMallocHost。
-
cudaHostAllocPortable
分配所有GPU都可使用的锁页内存 -
cudaHostAllocMapped。
此标志下分配的锁页内存可实现零拷贝功能，主机端和设备端各维护一个地址，通过 地址直接访问该块内存，无需传输。 -
cudaHostAllocWriteCombined 将分配的锁页内存声明为write-combined写联合内存，此类内存不使用L1和L2cache，所以程序的其它部分就有更多的缓存可用。此外，write-combined内存通过PCIe传输数据时不会被监视，能够获得更高的传输速度。因为没有使用L1、L2cache，所以主机读取write-combined内存很慢，write-combined适用于主机端写入、设备端读取的锁页内存。
分配的锁页内存必须使用cudaFreeHost接口释放。
对于一个已存在的可分页内存，可使用cudaHostRegister() 函数将其注册为锁页内存：
__host__ cudaError_t cudaHostRegister ( void* ptr, size_t size, unsigned int flags )
flags和上面一致。
锁页内存的缺点是分配空间过多可能会降低主机系统的性能，因为它减少了用于存储虚拟内存数据的可分页内存的数量。对于图像这类小内存应用还是比较合适的。
（5）零拷贝内存（Zero-Copy）
通常来说，设备端无法直接访问主机内存，但有一个例外：零拷贝内存！顾名思义，零拷贝内存是无需拷贝就可以在主机端和设备端直接访问的内存。
零拷贝具有如下优势：
- 当设备内存不足时可以利用主机内存
- 避免主机和设备间的显式数据传输
**准确来说，零拷贝并不是无需拷贝，而是无需显式拷贝。**使用零拷贝内存时不需要cudaMemcpy之类的显式拷贝操作，直接通过指针取值，所以对调用者来说似乎是没有拷贝操作。但实际上是在引用内存中某个值时隐式走PCIe总线拷贝，这样的方式有几个优点：
- 无需所有数据一次性显式拷贝到设备端，而是引用某个数据时即时隐式拷贝
- 隐式拷贝是异步的，可以和计算并行，隐藏内存传输延时
零拷贝内存是一块主机端和设备端共享的内存区域，是锁页内存，使用cudaHostAlloc接口分配。上一小结已经介绍了零拷贝内存的分配方法。分配标志是cudaHostAllocMapped。
对于零拷贝内存，设备端和主机端分别有一个地址，主机端分配时即可获取，设备端通过函数cudaHostGetDevicePointer函数获取地址。
__host__ cudaError_t cudaHostGetDevicePointer ( void** pDevice, void* pHost, unsigned int flags )
该函数返回一个在设备端的指针pDevice，该指针可以在设备端被引用以访问映射得到的主机端锁页内存。如果设备端不支持零拷贝方式（主机内存映射），则返回失败。可以使用接口cudaGetDeviceProperties来检查设备是否支持主机内存映射：
struct cudaDeviceProp device_prop
cudaGetDeviceProperties(&device_prop,device_num);
zero_copy_supported=device_prop.canMapHostMemory;
如上所述，零拷贝不是无需拷贝，而是一种隐式异步即时拷贝策略，每次隐式拷贝还是要走PCIe总线，所以频繁的对零拷贝内存进行读写，性能也会显著降低。
以下几种情况，可建议使用零拷贝内存：
- 在一大块主机内存中你只需要使用少量数据
- 你不会频繁的对这块内存进行重复访问，频繁的重复访问建议在设备端分配内存显式拷贝。最合适的情况，该内存的数据你都只需要访问一次
- 你需要比显存容量大的内存，或许你可以通过即时交换来获得比显存更大的内存使用，但是零拷贝内存也是一个可选思路
核心代码使用：
// allocate the memory on the CPU
cudaHostAlloc((void**) &a, size * sizeof(float),
cudaHostAllocWriteCombined | cudaHostAllocMapped);
cudaHostAlloc((void**) &b, size * sizeof(float),
cudaHostAllocWriteCombined | cudaHostAllocMapped);
cudaHostAlloc((void**) &partial_c, blocksPerGrid * sizeof(float),
cudaHostAllocMapped);
// find out the GPU pointers
cudaHostGetDevicePointer(&dev_a, a, 0);
cudaHostGetDevicePointer(&dev_b, b, 0);
cudaHostGetDevicePointer(&dev_partial_c, partial_c, 0);
（6）CUDA流的使用
CUDA流在加速应用程序方面起着重要的作用。CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。我们可以在流中添加一些操作，如核函数启动，内存复制等。将这些操作添加到流的顺序也就是他们的执行顺序。你可以将每个流视为GPU上的一个任务，并且这些任务可以并行执行。
1） 首先，选择一个支持设备重叠功能的设备。支持设备重叠功能的GPU能够在执行一个CUDA C/C++核函数的同时，还能在设备与主机之间执行复制操作。
cudaDeviceProp prop;
int whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDeviceProperties(&prop, whichDevice);
if (!prop.deviceOverlap)
{
printf("Device will not handle overlaps, so no speed up from streams.\n");
return 0;
}
2） 接下来，创建在应用程序中使用的流：
cudaStream_t stream;
cudaStreamCreate(&stream);
3） 然后是数据分配操作。注意，程序将使用主机上的固定内存，即调用cudaHostAlloc()来执行内存分配：
int *host_a, *host_b, *host_c;
int *dev_a, *dev_b, *dev_c;
cudaError_t cudaStatus;
cudaStatus = cudaMalloc((void **)&dev_a, N * sizeof(int));
if (cudaStatus != cudaSuccess)
{
printf("cudaMalloc dev_a failed!\n");
}
cudaStatus = cudaHostAlloc((void **)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
if (cudaStatus != cudaSuccess)
{
printf("cudaHostAlloc host_a failed!\n");
}
4） 在执行核函数时，首先我们不会将输入缓冲区整体都复制到GPU，而是将输入缓冲区划分为更小的块，并在每个块上执行一个包含三个步骤（复制到GPU–运行核函数–复制回主机）的过程。需要这种方法的一种情形是：GPU的内存远小于主机内存，由于整个缓冲区无法一次性填充到GPU，因此需要分块进行计算：
for (int i = 0; i < FULL_DATA_SIZE; i += N)
{
cudaStatus = cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync a failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync b failed!\n");
}
kernel << <N / GPUBLOCKNUM, GPUTHREADNUM, 0, stream >> >(dev_a, dev_b, dev_c);
cudaStatus = cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync c failed!\n");
}
}
注意，这段代码中并没有使用cudaMemcpy()，而是通过cudaMemcpyAsync()在GPU与主机之间复制数据。函数差异虽小，但却很重要。cudaMemcpy()的行为类似于C库函数memcpy()。尤其是，这个函数将以同步方式执行，也就是说，当函数返回时，复制操作已经完成。
异步函数的行为与同步函数相反，在调用cudaMemcpyAsync()时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过参数stream来指定的。当函数返回时，我们无法确保复制操作是否已经启动或完成。我们能够保证的是复制操作肯定会在下一个被放入流中的操作启动之前执行。任何传递给cudaMemcpyAsync()的主机内存指针都必须已经通过cudaHostAlloc()分配好内存，也就是，只能以异步方式对页锁定内存进行复制操作。
注意，在核函数调用的尖括号中有一个流参数stream，此时核函数调用将是异步的。从技术上来说，当循环迭代完一次时，有可能不会启动任何内存复制或核函数执行。但能够确保的是，第一次放入流中的复制操作将在第二次复制操作之前执行，第二个复制操作将在核函数启动之前执行完成。这意味着，代码中for循环的完成不保证流的完成，每个流中的任务都可能处于等待状态。
5） 当for循环结束时，队列中应该包含了许多等待GPU执行的工作。如果想要确保GPU执行完了计算与内存复制等操作，那么就需要将GPU与主机同步。也就是说，主机在继续执行之前，要首先等待GPU执行完成。可以调用cudaStreamSynchronize()并指定想要等待的流：
cudaStatus = cudaStreamSynchronize(stream);
6） 当程序执行到stream与主机同步之后的代码时，所用计算与复制操作都已完成。此时需要释放缓冲区，并销毁对GPU操作进行排队的流：
cudaFreeHost(host_a);
cudaFreeHost(host_b);
cudaFreeHost(host_c);
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
cudaStreamDestroy(stream);
至此，单个流的使用已经讲完。
三、多个流的使用
（1） GPU的工作调度机制
程序员可以将流视为有序的操作序列，其中即包含内存复制操作，又包含核函数调用。然而，在硬件中没有流的概念，而是包含一个或多个引擎来执行内存复制操作，以及一个引擎来执行核函数。这些引擎彼此独立地对操作进行排队，因此将导致如下图所示的任务调度情形。
因此，在某种程度上，用户与硬件关于GPU工作的排队方式有着完全不同的理解，而CUDA驱动程序则负责对用户和硬件进行协调。首先，在操作被添加到流的顺序中包含了重要的依赖性。例如上图，第0个流对A的内存复制需要在对B的内存复制之前完成。然而，一旦这些操作放入到硬件的内存复制引擎和核函数执行引擎的队列中时，这些依赖性将丢失，因此CUDA驱动程序需要确保硬件的执行单元不破坏流内部的依赖性。也就是说，CUDA驱动程序负责安装这些操作的顺序把它们调度到硬件上执行，这就维持了流内部的依赖性。下图说明了这些依赖性。
理解了GPU的工作调度原理之后，我们可以得到关于这些操作在硬件上执行的时间线，如下图所示。
记住，硬件在处理内存复制和核函数执行时分别采用了不同的引擎。因此，将操作放入流中队列中的顺序将影响着CUDA驱动程序调用这些操作以及执行的方式。
（2） 高效的运用多个CUDA流
将操作放入流的队列时应采用宽度优先方式而非深度优先。也就是说，不是首先添加第0个流的所有四个操作，然后再添加第1个流的所有四个操作，而是将两个流交叉添加。实际代码如下：
for (int i = 0; i < FULL_DATA_SIZE; i += N * 2)
{
cudaStatus = cudaMemcpyAsync(dev0_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 a failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev1_a, host_a + N + i, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 a failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev0_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 b failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev1_b, host_b + N + i, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 b failed!\n");
}
kernel << <N / GPUBLOCKNUM, GPUTHREADNUM, 0, stream0 >> >(dev0_a, dev0_b, dev0_c);
kernel << <N / GPUBLOCKNUM, GPUTHREADNUM, 0, stream1 >> >(dev1_a, dev1_b, dev1_c);
cudaStatus = cudaMemcpyAsync(host_c + i, dev0_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 c failed!\n");
}
cudaStatus = cudaMemcpyAsync(host_c + N + i, dev1_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 c failed!\n");
}
}
此时，如果内存复制操作的时间与核函数执行的时间大致相当，那么新的执行时间线如下图所示。
代码如下：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda_profiler_api.h"
#include <stdio.h>
#include <malloc.h>
const int size = 10 * 1024 * 1024;
const int N = 1024 * 1024;
const int FULL_DATA_SIZE = 20 * N;
const int GPUTHREADNUM = 256;
const int GPUBLOCKNUM = 256;
//测试cudamalloc函数循环100次的运行时间
float cuda_malloc_test(int size, bool up)
{
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime;
cudaEventCreate(&start);
cudaEventCreate(&stop);
a = (int *)malloc(10 * 1024 * 1024 * sizeof(int));
printf("%p", a);
if (a == NULL)
{
printf("host memory malloc fail!\n");
}
cudaError_t cudaStatus = cudaMalloc((void **)&dev_a, size * sizeof(*dev_a));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!\n");
}
cudaEventRecord(start, 0);
for (int i = 0; i < 100; i++)
{
if (up)
{
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy Host to Device failed!\n");
return -1;
}
}
else
{
cudaStatus = cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy Device to Host failed!\n");
return -1;
}
}
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
free(a);
cudaFree(dev_a);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return elapsedTime;
}
//测试cudaHostAlloc函数循环100次的运行时间
float cuda_host_alloc_test(int size, bool up)
{
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaError_t cudaStatus = cudaHostAlloc((void **)&a, size * sizeof(*a), cudaHostAllocDefault);
if (cudaStatus != cudaSuccess)
{
printf("host alloc fail!\n");
return -1;
}
cudaStatus = cudaMalloc((void **)&dev_a, size * sizeof(*dev_a));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!\n");
return -1;
}
cudaEventRecord(start, 0);
for (int i = 0; i < 100; i++)
{
if (up)
{
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy Host to Device failed!\n");
return -1;
}
}
else
{
cudaStatus = cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy Device to Host failed!\n");
return -1;
}
}
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaFreeHost(a);
cudaFree(dev_a);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return elapsedTime;
}
//核函数：测试一个随便写的加法
__global__ void kernel(int *a, int *b, int *c)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N)
{
int idx1 = (idx + 1) % GPUTHREADNUM;
int idx2 = (idx + 2) % GPUTHREADNUM;
float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
c[idx] = (as + bs) / 2;
}
}
//调用核函数检测单个流的运行时间和两个流的运行时间，比较其效率
int main()
{
cudaDeviceProp prop;
int whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDeviceProperties(&prop, whichDevice);
if (!prop.deviceOverlap)
{
printf("Device will not handle overlaps, so no speed up from streams.\n");
return 0;
}
cudaEvent_t start, stop;
cudaStream_t stream0, stream1;
cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);
int *host_a, *host_b, *host_c;
int *dev0_a, *dev0_b, *dev0_c;
int *dev1_a, *dev1_b, *dev1_c;
//利用cudaMalloc函数分配GPU内存
cudaError_t cudaStatus;
cudaStatus = cudaMalloc((void **)&dev0_a, N * sizeof(int));
if (cudaStatus != cudaSuccess)
{
printf("cudaMalloc dev0_a failed!\n");
}
cudaStatus = cudaMalloc((void **)&dev0_b, N * sizeof(int));
if (cudaStatus != cudaSuccess)
{
printf("cudaMalloc dev0_b failed!\n");
}
cudaStatus = cudaMalloc((void **)&dev0_c, N * sizeof(int));
if (cudaStatus != cudaSuccess)
{
printf("cudaMalloc dev0_c failed!\n");
}
cudaStatus = cudaMalloc((void **)&dev1_a, N * sizeof(int));
if (cudaStatus != cudaSuccess)
{
printf("cudaMalloc dev1_a failed!\n");
}
cudaStatus = cudaMalloc((void **)&dev1_b, N * sizeof(int));
if (cudaStatus != cudaSuccess)
{
printf("cudaMalloc dev1_b failed!\n");
}
cudaStatus = cudaMalloc((void **)&dev1_c, N * sizeof(int));
if (cudaStatus != cudaSuccess)
{
printf("cudaMalloc dev1_c failed!\n");
}
//利用cudaHostAlloc分配主机固定内存
cudaStatus = cudaHostAlloc((void **)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
if (cudaStatus != cudaSuccess)
{
printf("cudaHostAlloc host_a failed!\n");
}
cudaStatus = cudaHostAlloc((void **)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
if (cudaStatus != cudaSuccess)
{
printf("cudaHostAlloc host_b failed!\n");
}
cudaStatus = cudaHostAlloc((void **)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
if (cudaStatus != cudaSuccess)
{
printf("cudaHostAlloc host_c failed!\n");
}
//生成主机数据
for (int i = 0; i < FULL_DATA_SIZE; i++)
{
host_a[i] = i - N;
host_b[i] = i;
}
kernel << <N / GPUBLOCKNUM, GPUTHREADNUM, 0, stream0 >> > (dev0_a, dev0_b, dev0_c);
float elapsedTime;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
//流传输
for (int i = 0; i < FULL_DATA_SIZE; i += N * 2)
{
cudaStatus = cudaMemcpyAsync(dev0_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 a failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev0_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 b failed!\n");
}
kernel << <N / GPUBLOCKNUM, GPUTHREADNUM, 0, stream0 >> >(dev0_a, dev0_b, dev0_c);
cudaStatus = cudaMemcpyAsync(dev1_a, host_a + N + i, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 a failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev1_b, host_b + N + i, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 b failed!\n");
}
kernel << <N / GPUBLOCKNUM, GPUTHREADNUM, 0, stream1 >> >(dev1_a, dev1_b, dev1_c);
cudaStatus = cudaMemcpyAsync(host_c + i, dev0_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 c failed!\n");
}
cudaStatus = cudaMemcpyAsync(host_c + N + i, dev1_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 c failed!\n");
}
}
cudaStatus = cudaStreamSynchronize(stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaStreamSynchronize0 failed!\n");
}
cudaStatus = cudaStreamSynchronize(stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaStreamSynchronize1 failed!\n");
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Time taken 11 : %3.1f ms\n", elapsedTime);
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
//高效流传输
for (int i = 0; i < FULL_DATA_SIZE; i += N * 2)
{
cudaStatus = cudaMemcpyAsync(dev0_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 a failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev1_a, host_a + N + i, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 a failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev0_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 b failed!\n");
}
cudaStatus = cudaMemcpyAsync(dev1_b, host_b + N + i, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 b failed!\n");
}
kernel << <N / GPUBLOCKNUM, GPUTHREADNUM, 0, stream0 >> > (dev0_a, dev0_b, dev0_c);
kernel << <N / GPUBLOCKNUM, GPUTHREADNUM, 0, stream1 >> > (dev1_a, dev1_b, dev1_c);
cudaStatus = cudaMemcpyAsync(host_c + i, dev0_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync0 c failed!\n");
}
cudaStatus = cudaMemcpyAsync(host_c + N + i, dev1_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaMemcpyAsync1 c failed!\n");
}
}
cudaStatus = cudaStreamSynchronize(stream0);
if (cudaStatus != cudaSuccess)
{
printf("cudaStreamSynchronize0 failed!\n");
}
cudaStatus = cudaStreamSynchronize(stream1);
if (cudaStatus != cudaSuccess)
{
printf("cudaStreamSynchronize1 failed!\n");
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Time taken : %3.1f ms\n", elapsedTime);
cudaFreeHost(host_a);
cudaFreeHost(host_b);
cudaFreeHost(host_c);
cudaFree(dev0_a);
cudaFree(dev0_b);
cudaFree(dev0_c);
cudaFree(dev1_a);
cudaFree(dev1_b);
cudaFree(dev1_c);
cudaStreamDestroy(stream0);
cudaStreamDestroy(stream1);
cudaProfilerStop();
return 0;
}
/* 主函数调用cuda_malloc_test函数与cuda_host_alloc_test函数进行固定内存的测试*/
/*
int main(int argc, char ** argv)
{
float elapsedTime;
float MB = (float)100 * size * sizeof(int) / 1024 / 1024;
elapsedTime = cuda_malloc_test(size, true);
printf("Time using cudaMalloc: %3.1f ms.", elapsedTime);
printf("\tMB/s during copy up: %3.1f.\n", MB / (elapsedTime / 1000));
elapsedTime = cuda_malloc_test(size, false);
printf("Time using cudaMalloc: %3.1f ms.", elapsedTime);
printf("\tMB/s during copy up: %3.1f.\n", MB / (elapsedTime / 1000));
elapsedTime = cuda_host_alloc_test(size, true);
printf("Time using cudaMalloc: %3.1f ms.", elapsedTime);
printf("\tMB/s during copy up: %3.1f.\n", MB / (elapsedTime / 1000));
elapsedTime = cuda_host_alloc_test(size, false);
printf("Time using cudaMalloc: %3.1f ms.", elapsedTime);
printf("\tMB/s during copy up: %3.1f.\n", MB / (elapsedTime / 1000));
return 0;
}
*/
四、CUDA C/C++中如何隐藏数据传输
核函数执行和数据传输的重叠
要实现它有几点要求：
- 设备必须可以“并发地拷贝和执行”。我们可以通过访问cudaDeviceProp结构体的deviceOverlap属性或者从CUDA SDK/Toolkit中deviceQuery示例程序的输出中获得。几乎所有计算能力1.1及以上的设备都支持设备重叠。
- 核函数执行和数据传输必须在不同的非默认流中。
- 涉及到数据传输的主机内存必须是固定主机内存。
使用多个CUDA流,看一看是否实现了数据传输的隐藏。
在这个被修改的代码中，我们将大小为N的数组分为streamSize大小的数据块。既然核函数可以独立地操作所有数据，那么每个数据块也可以被独立地处理。流(非默认流)的数量nStreams=N/streamSize。实现数据的分解处理有多种方式，一种是将对每个数据块的所有操作都放到一个循环中，代码如下所示：
for (int i = 0; i < nStreams; ++i) {
int offset = i * streamSize;
cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i]);
kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
cudaMemcpyAsync(&a[offset], &d_a[offset], streamBytes, cudaMemcpyDeviceToHost, stream[i]);
}
另一种方式是将类似的操作放在一起批处理，首先发布所有主机到设备的数据传输，之后是核函数执行，然后就是设备到主机的数据传输，代码如下所示：
for (int i = 0; i < nStreams; ++i) {
int offset = i * streamSize;
cudaMemcpyAsync(&d_a[offset], &a[offset],
streamBytes, cudaMemcpyHostToDevice, cudaMemcpyHostToDevice, stream[i]);
}
for (int i = 0; i < nStreams; ++i) {
int offset = i * streamSize;
kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
}
for (int i = 0; i < nStreams; ++i) {
int offset = i * streamSize;
cudaMemcpyAsync(&a[offset], &d_a[offset],
streamBytes, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToHost, stream[i]);
}
上述两种异步方法都会产生正确的结果，而且同一个流中相互依赖的操作都会按照需要的顺序执行。
可以发现改进过后的代码传输效率更高，应为类似这种数据打包的方式减少了不必要的传输消耗。
代码如下：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <assert.h>
using namespace std;
// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
if (result != cudaSuccess) {
fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
assert(result == cudaSuccess);
}
#endif
return result;
}
__global__ void kernel(float* a, int offset)
{
int i = offset + threadIdx.x + blockIdx.x * blockDim.x;
float x = (float)i;
float s = asinf(x);
float c = acosf(x);
a[i] = a[i] + sqrtf(s * s + c * c);
}
float maxError(float* a, int n)
{
float maxE = 0;
for (int i = 0; i < n; i++) {
float error = fabs(a[i] - 1.0f);
if (error > maxE) maxE = error;
}
return maxE;
}
int main(int argc, char** argv)
{
const int blockSize = 512, nStreams = 8;
const int n = 4 * 1024 * blockSize * nStreams * 8;
const int streamSize = n / nStreams;
const int streamBytes = streamSize * sizeof(float);
const int bytes = n * sizeof(float);
int devId = 0;
if (argc > 1) devId = atoi(argv[1]);
cudaDeviceProp prop;
checkCuda(cudaGetDeviceProperties(&prop, devId));
printf("Device : %s\n", prop.name);
checkCuda(cudaSetDevice(devId));
// allocate pinned host memory and device memory
float* a, * d_a;
checkCuda(cudaMallocHost((void**)&a, bytes)); // host pinned
checkCuda(cudaMalloc((void**)&d_a, bytes)); // device
float ms; // elapsed time in milliseconds
// create events and streams
cudaEvent_t startEvent, stopEvent, dummyEvent;
cudaStream_t stream[nStreams];
checkCuda(cudaEventCreate(&startEvent));
checkCuda(cudaEventCreate(&stopEvent));
checkCuda(cudaEventCreate(&dummyEvent));
for (int i = 0; i < nStreams; ++i)
checkCuda(cudaStreamCreate(&stream[i]));
// baseline case - sequential transfer and execute
memset(a, 0, bytes);
checkCuda(cudaEventRecord(startEvent, 0));
checkCuda(cudaMemcpy(d_a, a, bytes, cudaMemcpyHostToDevice));
kernel << <n / blockSize, blockSize >> > (d_a, 0);
checkCuda(cudaMemcpy(a, d_a, bytes, cudaMemcpyDeviceToHost));
checkCuda(cudaEventRecord(stopEvent, 0));
checkCuda(cudaEventSynchronize(stopEvent));
checkCuda(cudaEventElapsedTime(&ms, startEvent, stopEvent));
printf("Time for sequential transfer and execute (ms): %f\n", ms);
printf(" max error: %e\n", maxError(a, n));
// asynchronous version 1: loop over {copy, kernel, copy}
memset(a, 0, bytes);
checkCuda(cudaEventRecord(startEvent, 0));
for (int i = 0; i < nStreams; ++i) {
int offset = i * streamSize;
checkCuda(cudaMemcpyAsync(&d_a[offset], &a[offset],
streamBytes, cudaMemcpyHostToDevice,
stream[i]));
kernel << <streamSize / blockSize, blockSize, 0, stream[i] >> > (d_a, offset);
checkCuda(cudaMemcpyAsync(&a[offset], &d_a[offset],
streamBytes, cudaMemcpyDeviceToHost,
stream[i]));
}
checkCuda(cudaEventRecord(stopEvent, 0));
checkCuda(cudaEventSynchronize(stopEvent));
checkCuda(cudaEventElapsedTime(&ms, startEvent, stopEvent));
printf("Time for asynchronous V1 transfer and execute (ms): %f\n", ms);
printf(" max error: %e\n", maxError(a, n));
// asynchronous version 2:
// loop over copy, loop over kernel, loop over copy
memset(a, 0, bytes);
checkCuda(cudaEventRecord(startEvent, 0));
for (int i = 0; i < nStreams; ++i)
{
int offset = i * streamSize;
checkCuda(cudaMemcpyAsync(&d_a[offset], &a[offset],
streamBytes, cudaMemcpyHostToDevice,
stream[i]));
}
for (int i = 0; i < nStreams; ++i)
{
int offset = i * streamSize;
kernel << <streamSize / blockSize, blockSize, 0, stream[i] >> > (d_a, offset);
}
for (int i = 0; i < nStreams; ++i)
{
int offset = i * streamSize;
checkCuda(cudaMemcpyAsync(&a[offset], &d_a[offset],
streamBytes, cudaMemcpyDeviceToHost,
stream[i]));
}
checkCuda(cudaEventRecord(stopEvent, 0));
checkCuda(cudaEventSynchronize(stopEvent));
checkCuda(cudaEventElapsedTime(&ms, startEvent, stopEvent));
printf("Time for asynchronous V2 transfer and execute (ms): %f\n", ms);
printf(" max error: %e\n", maxError(a, n));
// cleanup
checkCuda(cudaEventDestroy(startEvent));
checkCuda(cudaEventDestroy(stopEvent));
checkCuda(cudaEventDestroy(dummyEvent));
for (int i = 0; i < nStreams; ++i)
checkCuda(cudaStreamDestroy(stream[i]));
cudaFree(d_a);
cudaFreeHost(a);
return 0;
}
850
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
