网站建设简历,怎么做淘宝推广网站,全国企业信息公示(全国),郑州短视频拍摄制作公司文章目录页锁定主机内存可分页内存函数页锁定内存函数CUDA流使用单个CUDA流使用多个CUDA流GPU的工作调度机制高效地使用多个CUDA流遇到的问题(未解决)页锁定主机内存
在之前的各个示例中#xff0c;都是通过 cudaMalloc() 在GPU上分配内存#xff0c;以及通过标准的C库函数 …
文章目录页锁定主机内存可分页内存函数页锁定内存函数CUDA流使用单个CUDA流使用多个CUDA流GPU的工作调度机制高效地使用多个CUDA流遇到的问题(未解决)页锁定主机内存
在之前的各个示例中都是通过 cudaMalloc() 在GPU上分配内存以及通过标准的C库函数 malloc() 在主机上分配内存。除此之外CUDA运行时还提供了自己独有的机制来分配主机内存cudaHostAlloc()。如果 malloc() 已经能很好地满足C程序员的需求那么为什么还要使用这个函数
事实上malloc() 分配的内存与 cudaHostAlloc() 分配的内存之间存在着一个重要差异。C库函数 malloc() 将分配标准的、可分页的(Pagable) 主机内存而 cudaHostAlloc() 将分配页锁定的主机内存。页锁定内存也称为固定内存(Pinned Memory)或者不可分页内存它有一个重要的属性操作系统将不会对这块内存分页并交换到磁盘上从而确保了该内存始终驻留在物理内存中。因此操作系统能够安全地使某个应用程序访问该内存的物理地址因为这块内存将不会被破坏或者重新定位。
由于GPU知道内存的物理地址因此可以通过 “直接内存访问(Direct Memory Access, DMA)” 技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU的介入这也就同样意味着CPU很可能在DMA的执行过程中将目标内存交换到磁盘上或者通过更新操作系统的可分页表来重新定位目标内存的物理地址。CPU可能会移动可分页的数据这就可能对DMA操作造成延迟。因此在DMA复制过程中使用固定内存是非常重要的。事实上当使用可分页内存进行复制时CUDA驱动程序仍然会通过DMA把数据传输给GPU。因此复制操作将执行两遍
第一遍从可分页内存复制到一块 ”临时的“ 页锁定内存然后再从这个页锁定内存复制到GPU上
因此每当从可分页内存中执行复制操作时复制速度将受限于PCIE传输速度和系统前端总线速度相对较低的一方。在某些系统中这些总线在带宽上有着巨大的差异。因此当在GPU和主机间复制数据时这种差异会使页锁定主机内存的性能比标准可分页内存的性能要高大约2倍。即使PCIE的速度与前端总线的速度相等由于可分页内存需要更多一次由CPU参与的复制操作因此会带来额外的开销。
然而你也不能进入另一个极端查找每个 malloc 调用并将其替换为 cudaHostAlloc() 调用。固定内存是一把双刃剑。当使用固定内存时你将失去虚拟内存的所有功能。特别是在应用程序中使用每个页锁定内存时都需要分配物理内存因为这些内存不能交换到磁盘上。这意味着与使用标准的malloc()调用相比系统将更快地耗尽内存。因此应用程序在物理内存较少的机器上会运行失败而且意味着应用程序将影响在系统上运行的其他应用程序的性能。
这些情况并不是说不使用 cudaHostAlloc()而是提醒你应该清楚页锁定内存得到隐含作用。我们建议仅对 cudaMemcpy() 调用中的源内存或者目标内存才使用页锁定内存并且在不再需要使用它们时立即施放而不是等到应用程序关闭时才施放。cudaHostAlloc() 与到目前为止学习的其他内容一样简单下面通过一个示例说明如何分配固定内存以及它对于标准可分页内存的性能优势。
这里要做的就是分配一个GPU缓冲区以及一个大小相等的主机缓冲区然后在这两个缓冲区之间执行一些复制操作。我们允许用户指定复制的方向例如为 “上”从主机到设备或者为 “下”从设备到主机。为了获得精确的时间统计我们为复制操作的起始时刻和结束时刻分别设置了CUDA事件。 可分页内存函数
首先为 size 个整数分别分配主机缓冲区和GPU缓冲区
float cuda_malloc_test(int size, bool up) {cudaEvent_t start, stop;int* a, * dev_a;float elapsedTime;HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));a (int*)malloc(size * sizeof(*a));HANDLE_NULL(a);HANDLE_ERROR(cudaMalloc((void**)dev_a, size * sizeof(*dev_a)));然后执行100次复制操作并由参数 up 来指定复制方向在完成复制操作后停止计时器。 HANDLE_ERROR(cudaEventRecord(start, 0));for (int i 0; i 100; i) {if (up)HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice));elseHANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost));}HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));在执行了 100 次复制操作后释放主机缓冲区和GPU缓冲区并且销毁计时事件。 free(a);HANDLE_ERROR(cudaFree(dev_a));HANDLE_ERROR(cudaEventDestroy(start));HANDLE_ERROR(cudaEventDestroy(stop));return elapsedTime;
}页锁定内存函数
与可分页内存函数的区别就在于使用 cudaHostAlloc() 分配内存使用 cudaFreeHost() 施放内存
float cuda_host_alloc_test(int size, bool up) {cudaEvent_t start, stop;int* a, * dev_a;float elapsedTime;HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));HANDLE_ERROR(cudaHostAlloc((void**)a, size * sizeof(*a), cudaHostAllocDefault));HANDLE_ERROR(cudaMalloc((void**)dev_a, size * sizeof(*dev_a)));HANDLE_ERROR(cudaEventRecord(start, 0));for (int i 0; i 100; i) {if (up)HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice));elseHANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost));}HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));HANDLE_ERROR(cudaFreeHost(a));HANDLE_ERROR(cudaFree(dev_a));HANDLE_ERROR(cudaEventDestroy(start));HANDLE_ERROR(cudaEventDestroy(stop));return elapsedTime;
}完整代码 #include cuda_runtime.h
#include device_launch_parameters.h
#include ../../common/book.h#include stdio.h
#include iostream#define SIZE (10*1024*1024)float cuda_malloc_test(int size, bool up) {cudaEvent_t start, stop;int* a, * dev_a;float elapsedTime;HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));a (int*)malloc(size * sizeof(*a));HANDLE_NULL(a);HANDLE_ERROR(cudaMalloc((void**)dev_a, size * sizeof(*dev_a)));HANDLE_ERROR(cudaEventRecord(start, 0));for (int i 0; i 100; i) {if (up)HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice));elseHANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost));}HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));free(a);HANDLE_ERROR(cudaFree(dev_a));HANDLE_ERROR(cudaEventDestroy(start));HANDLE_ERROR(cudaEventDestroy(stop));return elapsedTime;
}float cuda_host_alloc_test(int size, bool up) {cudaEvent_t start, stop;int* a, * dev_a;float elapsedTime;HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));HANDLE_ERROR(cudaHostAlloc((void**)a, size * sizeof(*a), cudaHostAllocDefault));HANDLE_ERROR(cudaMalloc((void**)dev_a, size * sizeof(*dev_a)));HANDLE_ERROR(cudaEventRecord(start, 0));for (int i 0; i 100; i) {if (up)HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice));elseHANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost));}HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));HANDLE_ERROR(cudaFreeHost(a));HANDLE_ERROR(cudaFree(dev_a));HANDLE_ERROR(cudaEventDestroy(start));HANDLE_ERROR(cudaEventDestroy(stop));return elapsedTime;
}int main(void) {float elapsedTime;float MB (float)100 * SIZE * sizeof(int) / 1024 / 1024;elapsedTime cuda_malloc_test(SIZE, true);std::cout Time using cudaMalloc: elapsedTime ms\n;std::cout \tMB/s during copy up: MB / (elapsedTime / 1000) std::endl;elapsedTime cuda_malloc_test(SIZE, false);std::cout Time using cudaMalloc: elapsedTime ms\n;std::cout \tMB/s during copy down: MB / (elapsedTime / 1000) std::endl;elapsedTime cuda_host_alloc_test(SIZE, true);std::cout Time using cudaHostAlloc: elapsedTime ms\n;std::cout \tMB/s during copy up: MB / (elapsedTime / 1000) std::endl;elapsedTime cuda_host_alloc_test(SIZE, false);std::cout Time using cudaHostAlloc: elapsedTime ms\n;std::cout \tMB/s during copy down: MB / (elapsedTime / 1000) std::endl;
}
运行结果 可以发现使用页锁定内存比使用可分页内存的读写速度快了2倍多。 CUDA流
在之前的文章中我们引入了CUDA事件的概念。当时并没有介绍 cudaEventRecord() 的第二个参数而只是简要地指出这个参数用于指定插入事件的流(Stream)。
cudaEvent_t start;
cudaEventCreate(start);
cudaEventRecord(start, 0);CUDA流在加速应用程序方面起着重要的作用。CUDA流表示一个GPU操作队列并且该队列中的操作将以指定的顺序执行。我们可以在流中添加一些操作例如核函数启动、内存复制、以及事件的启动和结束等。将这些操作添加到流的顺序也就是它们的执行顺序。你可以将每个流视为GPU上的一个任务并且这些任务可以并行执行。
下面将首先介绍如何使用流然后介绍如何使用流来加速应用程序。 使用单个CUDA流
下面首先通过在应用程序中使用单个流来说明流的用法。假设有一个CUDA C核函数该函数带有两个输入数据缓冲区a 和 b。核函数将对这些缓冲区中相应位置上的值执行某种计算并将生成的结果保存到输出缓冲区 c。下面这个示例中将计算 a 中三个值和 b 中三个值的平均值
__global__ void kernel(int* a, int* b, int* c) {int idx threadIdx.x blockIdx.x * blockDim.x;if (idx N) {int idx1 (idx 1) % 256;int idx2 (idx 2) % 256;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;}
}这个核函数很简单下面重要的是函数 main() 中与流相关的代码
int main(void) {cudaDeviceProp prop;int whichDevice;HANDLE_ERROR(cudaGetDevice(whichDevice));HANDLE_ERROR(cudaGetDeviceProperties(prop, whichDevice));if (!prop.deviceOverlap) {std::cout Device will not handle overlaps, so no speed up from streams std::endl;return 0;}首先选择一个支持设备重叠(Device Overlap)功能的设备。支持设备重叠功能的GPU能够在执行一个CUDA C核函数的同时还能在设备与主机之间执行复制操作。
正如前面提到的我们将使用多个流来实现这种计算与数据传输的重叠但首先来看看如何创建和使用一个流。与其他需要测量性能提升(或者降低)的示例一样首先创建和启动一个事件计时器 cudaEvent_t start, stop;float elapsedTime;// 启动计时器HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));HANDLE_ERROR(cudaEventRecord(start, 0));启动计时器之后创建在应用程序中使用的流 // 初始化流cudaStream_t stream;HANDLE_ERROR(cudaStreamCreate(stream));这就是创建流需要的全部工作接下来是数据分配操作 int* host_a, * host_b, * host_c;int* dev_a, * dev_b, * dev_c;// 在GPU上分配内存HANDLE_ERROR(cudaMalloc((void**)dev_a, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_b, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_c, N * sizeof(int)));// 分配由流使用的页锁定内存HANDLE_ERROR(cudaHostAlloc((void**)host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));for (int i 0; i FULL_DATA_SIZE; i) {host_a[i] rand();host_b[i] rand();}我们在GPU和主机上分别分配好了输入内存和输出内存。注意由于程序将使用主机上的固定内存因此调用 cudaHostAlloc() 来执行内存分配操作。
使用固定内存的原因并不只在于使复制操作执行得更快还存在另外一个好处。会在后面进行详细地分析我们将使用一种新的 cudaMemcpy() 函数并且在这个新函数中需要页锁定主机内存。在分配完输入内存后调用C的库函数rand()并用随机整数填充主机内存。
在创建了流和计时事件并且分配了设备内存和主机内存后就准备好了执行一些计算。通常我们会将这个阶段一带而过只是将两个输入缓冲区复制到GPU启动核函数然后将输出缓冲区复制回主机。我们将再次沿用这种模式只是进行了一些小修改。
首先我们不将输入缓冲区整体都复制到GPU而是将输入缓冲区划分为更小的块并在每个块上执行一个包含三个步骤的过程。我们将一部分输入缓冲区复制到GPU在这部分缓冲区上运行核函数然后将输出缓冲区中的这部分结果复制回主机。
想象一下需要使用这种方法的一种情形GPU的内存远少于主机内存由于整个缓冲区无法一次性填充到GPU因此需要分块进行计算。
执行分块计算的代码如下所示 //在整体数据上循环每个数据块的大小为Nfor (int i 0; i FULL_DATA_SIZE; i N) {// 将锁定内存以异步方式复制到设备上HANDLE_ERROR(cudaMemcpyAsync(dev_a, host_a i, N * sizeof(int), cudaMemcpyHostToDevice, stream));HANDLE_ERROR(cudaMemcpyAsync(dev_b, host_b i, N * sizeof(int), cudaMemcpyHostToDevice, stream));kernel N / 256, 256, 0, stream (dev_a, dev_b, dev_c);// 将数据从设备复制到锁定内存HANDLE_ERROR(cudaMemcpyAsync(host_c i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream));}注意到代码没有使用熟悉的 cudaMemcpy()而是通过一个新函数 cudaMemcpyAsync() 在GPU与主机之间复制数据。这些函数之间的差异虽然很小但却很重要。cudaMemcpy() 的行为类似于C库函数 memcpy()。尤其是这个函数将以同步方式执行这意味着当函数返回时复制操作就已经完成并且在输出缓冲区中包含了复制进去的内容。
异步函数的行为与同步函数相反通过名字 cudaMemcpyAsync()就可以知道。在调用 cudaMemcpyAsync() 时只是放置一个请求表示在流中执行一次内存复制操作这个流是通过参数 stream 来指定的。当函数返回时我们无法确保复制操作是否已经启动更无法保证它是否已经结束。我们能够得到的保证是复制操作肯定会当下一个被放入流中的操作之前执行。任何传递给 cudaMemcpyAsync() 的主机内存指针都必须已经通过 cudaHostAlloc() 分配好内存。也就是你只能以异步方式对页锁定内存进行复制操作。
注意在核函数调用的尖括号中还可以带有一个流参数。此时核函数调用将是异步的就像之前与GPU之间的内存复制操作一样。从技术上来说当循环迭代完一次时有可能不会启动任何内存复制或核函数执行。
这里只能确保的是第一次放入流中的复制操作将在第二次复制操作之前执行。第二个复制操作将在核函数启动之前完成而核函数将在第三次复制操作开始之前完成。流就像一个有序的工作队列GPU从该队列中依次取出工作并执行。
当 for() 循环结束时在队列中应包含了许多等待GPU执行的工作。如果想要确保GPU执行完了计算和内存复制等操作那么就需要将GPU与主机同步。也就是说主机在继续执行之前要首先等待GPU执行完成。可以调用 cudaStreamSynchronize() 并指定想要等待的流 //将计算结果从页锁定内存复制到主机内存HANDLE_ERROR(cudaStreamSynchronize(stream));当程序执行到 stream 与主机同步之后的代码时所有的计算和复制操作都已经完成因此可以停止计时器收集性能数据并释放输入缓冲区和输出缓冲区。 HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));std::cout Time taken: elapsedTime ms std::endl;//释放流和内存HANDLE_ERROR(cudaFreeHost(host_a));HANDLE_ERROR(cudaFreeHost(host_b));HANDLE_ERROR(cudaFreeHost(host_c));HANDLE_ERROR(cudaFree(dev_a));HANDLE_ERROR(cudaFree(dev_b));HANDLE_ERROR(cudaFree(dev_c));最后在退出应用程序之前记得销毁对GPU操作进行排队的流。 HANDLE_ERROR(cudaStreamDestroy(stream));return 0;
}这个示例并没有充分说明流的强大功能。当然如果当主机正在执行一些工作时GPU也正忙于处理填充到流的工作那么即使使用单个流也有助于应用程序速度的提升。但即使不需要在主机上做太多的工作我们仍然可以通过使用流来加速应用程序。 完整代码
#include ../../common/book.h
#include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostream#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)__global__ void kernel(int* a, int* b, int* c) {int idx threadIdx.x blockIdx.x * blockDim.x;if (idx N) {int idx1 (idx 1) % 256;int idx2 (idx 2) % 256;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(void) {cudaDeviceProp prop;int whichDevice;HANDLE_ERROR(cudaGetDevice(whichDevice));HANDLE_ERROR(cudaGetDeviceProperties(prop, whichDevice));if (!prop.deviceOverlap) {std::cout Device will not handle overlaps, so no speed up from streams std::endl;return 0;}cudaEvent_t start, stop;float elapsedTime;// 启动计时器HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));HANDLE_ERROR(cudaEventRecord(start, 0));// 初始化流cudaStream_t stream;HANDLE_ERROR(cudaStreamCreate(stream));int* host_a, * host_b, * host_c;int* dev_a, * dev_b, * dev_c;// 在GPU上分配内存HANDLE_ERROR(cudaMalloc((void**)dev_a, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_b, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_c, N * sizeof(int)));// 分配由流使用的页锁定内存HANDLE_ERROR(cudaHostAlloc((void**)host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));for (int i 0; i FULL_DATA_SIZE; i) {host_a[i] rand();host_b[i] rand();}//在整体数据上循环每个数据块的大小为Nfor (int i 0; i FULL_DATA_SIZE; i N) {// 将锁定内存以异步方式复制到设备上HANDLE_ERROR(cudaMemcpyAsync(dev_a, host_a i, N * sizeof(int), cudaMemcpyHostToDevice, stream));HANDLE_ERROR(cudaMemcpyAsync(dev_b, host_b i, N * sizeof(int), cudaMemcpyHostToDevice, stream));kernel N / 256, 256, 0, stream (dev_a, dev_b, dev_c);// 将数据从设备复制到锁定内存HANDLE_ERROR(cudaMemcpyAsync(host_c i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream));}//将计算结果从页锁定内存复制到主机内存HANDLE_ERROR(cudaStreamSynchronize(stream));HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));std::cout Time taken: elapsedTime ms std::endl;//释放流和内存HANDLE_ERROR(cudaFreeHost(host_a));HANDLE_ERROR(cudaFreeHost(host_b));HANDLE_ERROR(cudaFreeHost(host_c));HANDLE_ERROR(cudaFree(dev_a));HANDLE_ERROR(cudaFree(dev_b));HANDLE_ERROR(cudaFree(dev_c));HANDLE_ERROR(cudaStreamDestroy(stream));return 0;
}运行结果 使用多个CUDA流
下面将单个流的版本改为使用两个不同的流。改进这个程序的思想很简单分块计算以及内存复制和核函数执行的重叠。
即在第 0 个流执行核函数的同时第一个流将输入缓冲区复制到GPU。然后在第 0 个流将计算结果复制回主机的同时第 1 个流将执行核函数…
如下图所示这里假设内存复制操作和核函数执行的时间大致相且GPU可以同时执行一个内存复制操作和一个核函数因此空的方框表示一个流正在等待执行哦某个操作的时刻这个操作不能与其他流的操作相互重叠。 事实上实际的执行时间线可能比上图给出的更好看在一些新的 NVIDIA GPU 中同时支持核函数和两次内存复制操作一次是从主机到设备另一次是从设备到主机。在任何支持内存复制和核函数的执行相互重叠的设备上当使用多个流时应用程序的整体性能都会提升。
核函数代码保持不变
__global__ void kernel(int* a, int* b, int* c) {int idx threadIdx.x blockIdx.x * blockDim.x;if (idx N) {int idx1 (idx 1) % 256;int idx2 (idx 2) % 256;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;}
}与使用单个流的版本一样我们将判断设备是否支持计算与内存复制操作的重叠。如果设备支持重叠那么就像前面一样创建CUDA事件并对应用程序计时。
int main( void ) {cudaDeviceProp prop;int whichDevice;HANDLE_ERROR(cudaGetDevice(whichDevice));HANDLE_ERROR(cudaGetDeviceProperties(prop, whichDevice));if (!prop.deviceOverlap) {std::cout Device will not handle overlaps, so no speed up from streams std::endl;return 0;}cudaEvent_t start, stop;float elapsedTime;// 启动计时器HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));HANDLE_ERROR(cudaEventRecord(start, 0));接下来创建两个流创建方式与前面单个流的版本完全一样。 // 初始化流cudaStream_t stream0, stream1;HANDLE_ERROR(cudaStreamCreate(stream0));HANDLE_ERROR(cudaStreamCreate(stream1));假设在主机上仍然是两个输入缓冲区和一个输出缓冲区。输入缓冲区中填充的是随机数据与使用单个流的应用程序采样的方式一样。然而现在我们将使用两个流来处理数据分配两组相同的GPU缓冲区这样每个流都可以独立地在输入数据块上执行工作。 int* host_a, * host_b, * host_c;int* dev_a0, * dev_b0, * dev_c0; // 为第0个流分配的GPU内存int* dev_a1, * dev_b1, * dev_c1; // 为第1个流分配的GPU内存// 在GPU上分配内存HANDLE_ERROR(cudaMalloc((void**)dev_a0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_b0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_c0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_a1, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_b1, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_c1, N * sizeof(int)));// 分配由流使用的页锁定内存HANDLE_ERROR(cudaHostAlloc((void**)host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));for (int i 0; i FULL_DATA_SIZE; i) {host_a[i] rand();host_b[i] rand();}然后程序在输入数据块上循环。然而由于现在使用了两个流因此在 for() 循环的迭代中需要处理的数据量也是原来的两倍。在 stream() 中我们首先将 a 和 b 的异步复制操作放入GPU的队列然后将一个核函数执行放入队列接下来再将一个复制回 c 的操作放入队列
//在整体数据上循环每个数据块的大小为Nfor (int i 0; i FULL_DATA_SIZE; i N*2) {// 将锁定内存以异步方式复制到设备上HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));kernel N / 256, 256, 0, stream0 (dev_a0, dev_b0, dev_c0);// 将数据从设备复制到锁定内存HANDLE_ERROR(cudaMemcpyAsync(host_c i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0));在将这些操作放入 stream0 的队列后再把下一个数据块上的相同操作放入 stream1 的队列中。 // 将锁定内存以异步方式复制到设备上HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a i N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));HANDLE_ERROR(cudaMemcpyAsync(dev_b1, host_b i N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));kernel N / 256, 256, 0, stream1 (dev_a1, dev_b1, dev_c1);// 将数据从设备复制到锁定内存HANDLE_ERROR(cudaMemcpyAsync(host_c i N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1));}这样在 for() 循环的迭代过程中将交替地把每个数据块放入这两个流的队列直到所有待处理的输入数据都被放入队列。在结束了 for() 循环后在停止应用程序的计时器之前首先将 GPU 与 GPU进行同步。由于使用了两个流因此需要对二者都进行同步。
HANDLE_ERROR(cudaStreamSynchronize(stream0)); HANDLE_ERROR(cudaStreamSynchronize(stream1));之后停止计时器显示经历的时间并且执行清理工作。当然我们要记住现在需要销毁两个流并且需要释放两倍的GPU内存。 HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));std::cout Time taken: elapsedTime ms std::endl;//释放流和内存HANDLE_ERROR(cudaFreeHost(host_a));HANDLE_ERROR(cudaFreeHost(host_b));HANDLE_ERROR(cudaFreeHost(host_c));HANDLE_ERROR(cudaFree(dev_a0));HANDLE_ERROR(cudaFree(dev_b0));HANDLE_ERROR(cudaFree(dev_c0));HANDLE_ERROR(cudaFree(dev_a1));HANDLE_ERROR(cudaFree(dev_b1));HANDLE_ERROR(cudaFree(dev_c1));HANDLE_ERROR(cudaStreamDestroy(stream0));HANDLE_ERROR(cudaStreamDestroy(stream1));return 0;
}完整代码
#include ../../common/book.h
#include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostream#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)__global__ void kernel(int* a, int* b, int* c) {int idx threadIdx.x blockIdx.x * blockDim.x;if (idx N) {int idx1 (idx 1) % 256;int idx2 (idx 2) % 256;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(void) {cudaDeviceProp prop;int whichDevice;HANDLE_ERROR(cudaGetDevice(whichDevice));HANDLE_ERROR(cudaGetDeviceProperties(prop, whichDevice));if (!prop.deviceOverlap) {std::cout Device will not handle overlaps, so no speed up from streams std::endl;return 0;}cudaEvent_t start, stop;float elapsedTime;// 启动计时器HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));HANDLE_ERROR(cudaEventRecord(start, 0));// 初始化流cudaStream_t stream0, stream1;HANDLE_ERROR(cudaStreamCreate(stream0));HANDLE_ERROR(cudaStreamCreate(stream1));int* host_a, * host_b, * host_c;int* dev_a0, * dev_b0, * dev_c0; // 为第0个流分配的GPU内存int* dev_a1, * dev_b1, * dev_c1; // 为第1个流分配的GPU内存// 在GPU上分配内存HANDLE_ERROR(cudaMalloc((void**)dev_a0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_b0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_c0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_a1, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_b1, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_c1, N * sizeof(int)));// 分配由流使用的页锁定内存HANDLE_ERROR(cudaHostAlloc((void**)host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));for (int i 0; i FULL_DATA_SIZE; i) {host_a[i] rand();host_b[i] rand();}//在整体数据上循环每个数据块的大小为Nfor (int i 0; i FULL_DATA_SIZE; i N*2) {// 将锁定内存以异步方式复制到设备上HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));kernel N / 256, 256, 0, stream0 (dev_a0, dev_b0, dev_c0);// 将数据从设备复制到锁定内存HANDLE_ERROR(cudaMemcpyAsync(host_c i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0));// 将锁定内存以异步方式复制到设备上HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a i N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));HANDLE_ERROR(cudaMemcpyAsync(dev_b1, host_b i N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));kernel N / 256, 256, 0, stream1 (dev_a1, dev_b1, dev_c1);// 将数据从设备复制到锁定内存HANDLE_ERROR(cudaMemcpyAsync(host_c i N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1));}HANDLE_ERROR(cudaStreamSynchronize(stream0));HANDLE_ERROR(cudaStreamSynchronize(stream1));HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));std::cout Time taken: elapsedTime ms std::endl;//释放流和内存HANDLE_ERROR(cudaFreeHost(host_a));HANDLE_ERROR(cudaFreeHost(host_b));HANDLE_ERROR(cudaFreeHost(host_c));HANDLE_ERROR(cudaFree(dev_a0));HANDLE_ERROR(cudaFree(dev_b0));HANDLE_ERROR(cudaFree(dev_c0));HANDLE_ERROR(cudaFree(dev_a1));HANDLE_ERROR(cudaFree(dev_b1));HANDLE_ERROR(cudaFree(dev_c1));HANDLE_ERROR(cudaStreamDestroy(stream0));HANDLE_ERROR(cudaStreamDestroy(stream1));return 0;
}运行结果 GPU的工作调度机制
虽然从逻辑上来看不同的流之间是相互独立的但事实上这种理解并不完全符合GPU的队列机制。程序员可以将流视为有序的操作序列其中既包含内存复制操作又包含核函数调用。
然而在硬件中并没有流的概念而是包含一个或多个引擎来执行内存复制操作以及一个引擎来执行核函数。这些引擎彼此独立地对操作进行排队因此将导致如下图所示的任务调度情形。图中的箭头说明了硬件引擎如何调度流中队列的操作并实际执行。 因此在某种程度上用户与硬件关于GPU工作的排队方式有着完全不同的理解而CUDA驱动程序则负责对用户和硬件进行协调。首先在操作被添加到流的顺序中包含了重要的依赖性。
如上图中第 0 个流对 A 的内存复制需要在对 B 的内存复制之前完成而对 B 的复制又要在核函数 A 启动之前完成。然而一旦这些操作放入到硬件的内存复制引擎和核函数执行引擎的队列中这些依赖性将丢失因此CUDA驱动程序需要确保硬件的执行单元不破坏内部的依赖性。
这意味着说明之前在代码中应用程序基本上是对 a 调用一次 cudaMemcpyAsync()对 b 调用一次 cudaMemcpyAsync()然后再是执行核函数以及调用 cudaMemcpyAsync() 将 c 复制回主机。应用程序首先将对第 0 个流的所有操作放入队列然后是第 1 个流的所有操作。CUDA 驱动程序负责按照这些操作的顺序把它们调度到硬件上执行这就维持了流内部的依赖性。下图说明了这些依赖性其中从复制操作到核函数的箭头表示复制操作要等核函数执行完成之后才能开始。 假定理解了 GPU 的工作调度原理后我们可以得到关于这些操作再硬件上执行的时间线如下图所示 由于第 0 个流中将 c 复制回主机的操作要等待核函数执行完成因此第 1 个流中将 a 和 b 复制到 GPU 的操作虽然是完全独立的但却被阻塞了这是因为GPU引擎是按照指定的顺序来执行工作。这种情况也说明了为什么上面使用了两个流却没有获得很大的速度提升。
这个问题的直接原因是我们没有意识到硬件的工作方式与CUDA流编程模型的方式是不同的。
硬件在处理内存复制和核函数执行时分别采用了不同的引擎因此我们需要知道将操作放入流中队列中的顺序将影响着 CUDA 驱动程序调度这些操作以及执行的方式。下面我们将看到如何帮助硬件实现内存复制操作与核函数执行的重叠。 高效地使用多个CUDA流
如上节所看到的如果同时调度某个流的所有操作那么容易在无意中阻塞另一个流的复制操作或者核函数执行。要解决这个问题在将操作放入流的队列时应采用宽度优先方式而非深度优先方式。
也就是说不是首先添加第 0 个流的所有四个操作(即a的复制、b的复制、核函数以及c的复制)然后再添加第 1 个流的所有四个操作。而是将这两个流之间的操作交叉添加。首先将 a 的复制操作添加到第 0 个流然后将 a 的复制操作添加到第 1 个流。接着将 b 的复制操作添加到第 0 个流再将 b 的复制操作添加到第 1 个流。接下来将核函数调用添加到第 0 个流再将相同的操作添加到第 1 个流中。最后将 c 的复制操作添加到第 0 个流中然后将相同的操作添加到第 1 个流中。
下面给出具体的代码只需要修改 for() 循环内的代码。
for (int i 0; i FULL_DATA_SIZE; i N*2) {// 将复制a的操作放入stream0和stream1的队列HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a i N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));// 将复制b的操作放入stream0和stream1的队列HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));HANDLE_ERROR(cudaMemcpyAsync(dev_b1, hos i N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));// 将核函数的执行放入stream0和stream1的队列中kernel N / 256, 256, 0, stream0 (dev_a0, dev_b0, dev_c0);kernel N / 256, 256, 0, stream1 (dev_a1, dev_b1, dev_c1);// 将复制c的操作放入stream0和stream1的队列中HANDLE_ERROR(cudaMemcpyAsync(host_c i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0));HANDLE_ERROR(cudaMemcpyAsync(host_c i N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1));}如果内存复制操作的时间与核函数执行的时间大致相当那么新的执行时间线将如下图所示。引擎箭的依赖性通过箭头表示可以看到在新的调度顺序中这些依赖性仍然能得到满足。 由于采用了宽度优先方式将操作放入各个流的队列中因此第0个流对c的复制操作将不会阻塞第1个流对a和b的内存复制操作。这使得GPU能够并行地执行复制操作和核函数从而使应用程序的运行速度显著加快。
完整代码
#include ../../common/book.h
#include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostream#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)__global__ void kernel(int* a, int* b, int* c) {int idx threadIdx.x blockIdx.x * blockDim.x;if (idx N) {int idx1 (idx 1) % 256;int idx2 (idx 2) % 256;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(void) {cudaDeviceProp prop;int whichDevice;HANDLE_ERROR(cudaGetDevice(whichDevice));HANDLE_ERROR(cudaGetDeviceProperties(prop, whichDevice));if (!prop.deviceOverlap) {std::cout Device will not handle overlaps, so no speed up from streams std::endl;return 0;}cudaEvent_t start, stop;float elapsedTime;// 启动计时器HANDLE_ERROR(cudaEventCreate(start));HANDLE_ERROR(cudaEventCreate(stop));HANDLE_ERROR(cudaEventRecord(start, 0));// 初始化流cudaStream_t stream0, stream1;HANDLE_ERROR(cudaStreamCreate(stream0));HANDLE_ERROR(cudaStreamCreate(stream1));int* host_a, * host_b, * host_c;int* dev_a0, * dev_b0, * dev_c0; // 为第0个流分配的GPU内存int* dev_a1, * dev_b1, * dev_c1; // 为第1个流分配的GPU内存// 在GPU上分配内存HANDLE_ERROR(cudaMalloc((void**)dev_a0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_b0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_c0, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_a1, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_b1, N * sizeof(int)));HANDLE_ERROR(cudaMalloc((void**)dev_c1, N * sizeof(int)));// 分配由流使用的页锁定内存HANDLE_ERROR(cudaHostAlloc((void**)host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));HANDLE_ERROR(cudaHostAlloc((void**)host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));for (int i 0; i FULL_DATA_SIZE; i) {host_a[i] rand();host_b[i] rand();}//在整体数据上循环每个数据块的大小为Nfor (int i 0; i FULL_DATA_SIZE; i N*2) {// 将锁定内存以异步方式复制到设备上HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a i N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));// 将锁定内存以异步方式复制到设备上HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));HANDLE_ERROR(cudaMemcpyAsync(dev_b1, host_b i N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));kernel N / 256, 256, 0, stream0 (dev_a0, dev_b0, dev_c0);kernel N / 256, 256, 0, stream1 (dev_a1, dev_b1, dev_c1);// 将数据从设备复制到锁定内存HANDLE_ERROR(cudaMemcpyAsync(host_c i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0));// 将数据从设备复制到锁定内存HANDLE_ERROR(cudaMemcpyAsync(host_c i N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1));}HANDLE_ERROR(cudaStreamSynchronize(stream0));HANDLE_ERROR(cudaStreamSynchronize(stream1));HANDLE_ERROR(cudaEventRecord(stop, 0));HANDLE_ERROR(cudaEventSynchronize(stop));HANDLE_ERROR(cudaEventElapsedTime(elapsedTime, start, stop));std::cout Time taken: elapsedTime ms std::endl;//释放流和内存HANDLE_ERROR(cudaFreeHost(host_a));HANDLE_ERROR(cudaFreeHost(host_b));HANDLE_ERROR(cudaFreeHost(host_c));HANDLE_ERROR(cudaFree(dev_a0));HANDLE_ERROR(cudaFree(dev_b0));HANDLE_ERROR(cudaFree(dev_c0));HANDLE_ERROR(cudaFree(dev_a1));HANDLE_ERROR(cudaFree(dev_b1));HANDLE_ERROR(cudaFree(dev_c1));HANDLE_ERROR(cudaStreamDestroy(stream0));HANDLE_ERROR(cudaStreamDestroy(stream1));return 0;
}运行结果 遇到的问题(未解决)
我发现使用多个流实际并没有产生运行速度的提升我试了单个流、两个流、四个流发现消耗的时间基本没有差别暂时不知道是什么原因导致的欢迎大佬解答。
