为什么要使用CUDA流
前面章节实现的GPU并行性都是把同一个任务分成可以并行执行的子任务，利用GPU中数量众多的cuda核来实现单个任务的并行性。而如果要实现任务并行性，就需要使用CUDA流来实现。
页锁存主机内存
1. 页锁存内存的概念及用途
- 页所存内存也称为固定内存或者不可分页内存，它有一个重要属性：
操作系统将不会对这块内存分页并交换到磁盘上，从而保证了该内存始终驻留在物理内存中。 - 由于GPU知道主机内存的物理地址，因此可以通过DMA(Direct Memory Access)不需要CPU参与地在GPU和主机之间复制数据。正是因为CPU不参与数据复制，CPU很可能在DMA执行过程中移动可分页的数据，这就可能对DMA操作造成延迟。
事实上，当使用可分页内存进行复制时，复制操作将执行两遍：第一遍从可分页内存复制到一块"临时的"页锁定内存，然后再从这个页锁定内存复制到GPU上。 - 页锁定内存的使用将会在特定场合提高运行效率。
- 注意：当使用固定内存时，将会失去虚拟内存的所有功能，即系统将会更快地耗尽内存资源，因为没有办法将内存交换到磁盘中。
2. 页锁存内存的分配
通过cudaHostAlloc()分配页锁定的主机内存，通过cudaFreeHost()释放页锁定内存
//cudaHostAlloc()前两个参数是页锁定内存的地址和分配的大小
//最后一个参数的取值范围是一组标志，可以通过这些标志来修改cudaHostAlloc()的行为cudaHostAllocDefault将获得默认行为
cudaHostAlloc((void** )&a, size*sizeof(*a), cudaHostAllocDefault);
//...执行一些复制操作
//释放页锁定内存
cudaFreeHost(a);
- 页锁定内存的测试
// 代码10.2页锁存主机内存
//时间：2019.07.30
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
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(size*sizeof(*a));
cudaMalloc((void **)&dev_a, size*sizeof(*dev_a));
cudaEventRecord(start, 0);
for (int i = 0; i < 100; i++)
{
if (up)
{
cudaMemcpy(dev_a, a, size*sizeof(*dev_a), cudaMemcpyHostToDevice);
}
else
{
cudaMemcpy(a, dev_a, size*sizeof(*a), cudaMemcpyDeviceToHost);
}
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
free(a);
cudaFree(dev_a);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return elapsedTime;
}
float cuda_host_alloc_test(int size, bool up)
{
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaHostAlloc((void **)&a,size*sizeof(*a),cudaHostAllocDefault);
cudaMalloc((void **)&dev_a, size*sizeof(*dev_a));
cudaEventRecord(start, 0);
for (int i = 0; i < 100; i++)
{
if (up)
{
cudaMemcpy(dev_a, a, size*sizeof(*dev_a), cudaMemcpyHostToDevice);
}
else
{
cudaMemcpy(a, dev_a, size*sizeof(*a), cudaMemcpyDeviceToHost);
}
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaFreeHost(a);
cudaFree(dev_a);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return elapsedTime;
}
#define SIZE 10*1024*1024
int main()
{
float elapsedTime;
float MB = (float)100 * SIZE*sizeof(int) / 1024 / 1024;
elapsedTime = cuda_malloc_test(SIZE, true);
printf("Time using cudaMalloc:%3.1f ms\n", elapsedTime);
printf("\t MB/s during copy up:%3.1f M/s\n", MB / (elapsedTime / 1000));
elapsedTime = cuda_malloc_test(SIZE, false);
printf("Time using cudaMalloc:%3.1f ms\n", elapsedTime);
printf("\t MB/s during copy down:%3.1f M/s\n", MB / (elapsedTime / 1000));
elapsedTime = cuda_host_alloc_test(SIZE, true);
printf("Time using cudaHostAlloc:%3.1f ms\n", elapsedTime);
printf("\t MB/s during copy up:%3.1f M/s\n", MB / (elapsedTime / 1000));
elapsedTime = cuda_host_alloc_test(SIZE, false);
printf("Time using cudaHostAlloc:%3.1f ms\n", elapsedTime);
printf("\t MB/s during copy down:%3.1f M/s\n", MB / (elapsedTime / 1000));
system("pause");
}
发现使用页锁存内存对性能有明显的提升！
CUDA流
CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。可以在流中添加一些操作，例如核函数的启动、内存复制，以及事件的启动和结束等。将这些操作添加到流的顺序也就是它们的执行顺序。可以将每个流视为GPU上的一个任务，并且这些任务可以并行执行。
注意到，在第六章介绍cuda事件时，cudaEventRecord()一共有两个参数，第一个参数是要记录的事件，第二个参数用于指定插入事件的流。
cudaEvent_t start;
cudaEventCreate(&start);
cudaEventRecord(start,0);
使用单个CUDA流
1. 设备硬件要求
GPU必须支持设备重叠功能才能使用CUDA流。支持设备重叠功能的GPU能够在执行一个CUDA C核函数的同时，还能在设备与主机之间执行复制操作。我们最终将使用多个流来实现这种计算一数据传输的重叠，以提高运行效率。
cudaDeviceProp prop;
int whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDeviceProperties(&prop, whichDevice);
if (!prop.deviceOverlap)
{
printf("Device will not handle overlap,so no speed up from streams\n");
return 0;
}
2. 初始化流
//初始化流
cudaStream_t stream;
cudaStreamCreate(&stream);//创建一个流
3. 用异步复制取代同步复制
cudaMencpy()将以同步的方式执行，当函数返回时，复制操作就已经完成。
而cudaMemcpyAsync()是CUDA提供的异步复制函数。在调用cudaMemcpyAsync()时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过参数stream来指定的。当函数返回时，我们无法确定复制操作是否已经启动，更无法保证它是否已经结束。任何传递给cudaMemcpyAsync()的主机内存指针都必须是通过cudaHostAlloc()分配的也锁定内存，即只能以异步方式对页锁定内存进行复制操作。
cudaMemcpyAsync(dev_a, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream);
cudaMemcpyAsync(dev_b, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream);
4. 用异步核函数调用取代同步调用
kernal << <N / 256, 256, 0, stream >> >(dev_a, dev_b, dev_c);
5.主机与GPU同步
主机在继续执行之前，要首先等待GPU执行完成。此时，可以调用cudaStreamSynchronize()并指定想要等待的流
6. 销毁对GPU操作进行排队的流
在退出应用程序之前，要调用cudaStreamDestroy()销毁对GPU操作进行排队的流。
cudaStreamDestroy(stream);
7. 完整代码
示例代码假设在这样的情形下运行：GPU内存远远少于主机内存，由于整个缓冲区无法一次性填充到GPU，因此需要分块进行填充，分块计算，分块取回。
// 代码10.4使用单个cuda流
//时间：2019.07.30
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#define N 1024*1024
#define FULL_DATA_SIZE 20*N
__global__ void kernal(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()
{
cudaDeviceProp prop;
int whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDeviceProperties(&prop, whichDevice);
if (!prop.deviceOverlap)
{
printf("Device will not handle overlap,so no speed up from streams\n");
return 0;
}
cudaEvent_t start, stop;
float elapsedTime;
//启动计时器
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
//初始化流
cudaStream_t stream;
cudaStreamCreate(&stream);
int *host_a, *host_b, *host_c;
int *dev_a, *dev_b, *dev_c;
//在GPU上分配内存
cudaMalloc((void **)&dev_a, N*sizeof(int));
cudaMalloc((void **)&dev_b, N*sizeof(int));
cudaMalloc((void **)&dev_c, N*sizeof(int));
//分配由流使用的页锁定内存
cudaHostAlloc((void **)&host_a, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void **)&host_b, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault);
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();
}
//分块计算，在整体数据上循环，每个数据块大小为N
for (int i = 0; i < FULL_DATA_SIZE; i+=N)
{
cudaMemcpyAsync(dev_a, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream);
cudaMemcpyAsync(dev_b, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream);
//执行核函数
kernal << <N / 256, 256, 0, stream >> >(dev_a, dev_b, dev_c);
//将数据从设备复制到锁定内存
cudaMemcpyAsync(host_c + i, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost, stream);
}
//将GPU与主机同步
cudaStreamSynchronize(stream);
//停止计时器，收集性能数据
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(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
cudaStreamDestroy(stream);
cudaEventDestroy(start);
cudaEventDestroy(stop);
system("pause");
return 0;
}
使用多个CUDA流
我们用两个流**“分块”计算并且实现内存复制与核函数的重叠**。
在一些新的NVIDIA GPU中同时支持核函数和两次内存复制操作，一次是从主机到设备，另一次是从设备到主机。这样能使应用程序的整体性能获得再一步提升。
// 代码10.6使用多个cuda流
//时间：2019.07.30
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#define N 1024*1024
#define FULL_DATA_SIZE 20*N
__global__ void kernal(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()
{
cudaDeviceProp prop;
int whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDeviceProperties(&prop, whichDevice);
if (!prop.deviceOverlap)
{
printf("Device will not handle overlap,so no speed up from streams\n");
return 0;
}
cudaEvent_t start, stop;
float elapsedTime;
//启动计时器
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
//初始化流
cudaStream_t stream0,stream1;
cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);
int *host_a, *host_b, *host_c;
int *dev_a0, *dev_b0, *dev_c0;
int *dev_a1, *dev_b1, *dev_c1;
//在GPU上分配内存
cudaMalloc((void **)&dev_a0, N*sizeof(int));
cudaMalloc((void **)&dev_b0, N*sizeof(int));
cudaMalloc((void **)&dev_c0, N*sizeof(int));
cudaMalloc((void **)&dev_a1, N*sizeof(int));
cudaMalloc((void **)&dev_b1, N*sizeof(int));
cudaMalloc((void **)&dev_c1, N*sizeof(int));
//分配由流使用的页锁定内存
cudaHostAlloc((void **)&host_a, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void **)&host_b, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault);
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();
}
//分块计算，在整体数据上循环，每个数据块大小为N
for (int i = 0; i < FULL_DATA_SIZE; i += 2*N)
{
//第一个流
cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(dev_b0, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
//执行核函数
kernal << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
//将数据从设备复制到锁定内存
cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);
//第二个流
cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
//执行核函数
kernal << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
//将数据从设备复制到锁定内存
cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);
}
//将GPU与主机同步
cudaStreamSynchronize(stream0);
cudaStreamSynchronize(stream1);
//停止计时器，收集性能数据
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(dev_a0);
cudaFree(dev_b0);
cudaFree(dev_c0);
cudaFree(dev_a1);
cudaFree(dev_b1);
cudaFree(dev_c1);
cudaStreamDestroy(stream0);
cudaStreamDestroy(stream1);
cudaEventDestroy(start);
cudaEventDestroy(stop);
system("pause");
return 0;
}
实践表明，使用两个流确实比使用一个流效率更高。
GPU工作调度机制
程序员编程时可以将流视为有序的操作序列，其中既包含内存复制操作，又包含核函数调用。然而在硬件中并没有流的概念，而是包含一个或多个引擎来执行内存复制操作，以及一个引擎来执行核函数。这些引擎彼此独立地对操作进行排队。
将上一节两个流映射到GPU引擎：
注意到，在两个流的示例中，流队列的加入顺序为：
//第一个流
cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(dev_b0, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
//执行核函数
kernal << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
//将数据从设备复制到锁定内存
cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);
//第二个流
cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
//执行核函数
kernal << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
//将数据从设备复制到锁定内存
cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);
应用程序首先将第0个流的所有操作放入队列，然后是第1个流的所有操作。下图说明了流内部的依赖性，其中从复制操作到核函数的箭头表示：复制操作要等核函数执行完成之后才能开始。
在硬件上执行的时间线如下：
由于第0个流中将c复制回主机的操作要等待核函数执行完成，因此第一个流中将a和b复制到GPU的操作虽然是完全独立的，但却被阻塞了。
高效地使用多个CUDA流
如果同时调度某个流的所有操作(即按照流深度优先的方式调用)，那么很容易在无意中阻塞另一个流的复制操作或者核函数执行。要解决这个问题，在将操作放入流的队列时应该采用宽度优先方式，而非深度优先方式。
流的深度优先是指：先添加第0个流的所有操作(a复制，b复制，核函数，c复制)，再添加第1个流的所有操作。
流的广度优先是指：将a复制添加到第0个流，将a复制添加到第1个流；将b复制添加到第0个流，将b复制添加到第1个流；将核函数添加到第0个流，将核函数添加到第1个流；将c复制添加到第0个流，将c复制添加到第1个流。
代码修改为：
cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(dev_b0, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
//执行核函数
kernal << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
kernal << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
//将数据从设备复制到锁定内存
cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);
cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);
新的硬件上执行时间线如下：
完整代码：
// 代码10.7高效地使用多个cuda流
//时间：2019.07.30
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#define N 1024*1024
#define FULL_DATA_SIZE 20*N
__global__ void kernal(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()
{
cudaDeviceProp prop;
int whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDeviceProperties(&prop, whichDevice);
if (!prop.deviceOverlap)
{
printf("Device will not handle overlap,so no speed up from streams\n");
return 0;
}
cudaEvent_t start, stop;
float elapsedTime;
//启动计时器
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
//初始化流
cudaStream_t stream0, stream1;
cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);
int *host_a, *host_b, *host_c;
int *dev_a0, *dev_b0, *dev_c0;
int *dev_a1, *dev_b1, *dev_c1;
//在GPU上分配内存
cudaMalloc((void **)&dev_a0, N*sizeof(int));
cudaMalloc((void **)&dev_b0, N*sizeof(int));
cudaMalloc((void **)&dev_c0, N*sizeof(int));
cudaMalloc((void **)&dev_a1, N*sizeof(int));
cudaMalloc((void **)&dev_b1, N*sizeof(int));
cudaMalloc((void **)&dev_c1, N*sizeof(int));
//分配由流使用的页锁定内存
cudaHostAlloc((void **)&host_a, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void **)&host_b, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault);
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();
}
//分块计算，在整体数据上循环，每个数据块大小为N
for (int i = 0; i < FULL_DATA_SIZE; i += 2 * N)
{
cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(dev_b0, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
//执行核函数
kernal << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
kernal << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
//将数据从设备复制到锁定内存
cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);
cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);
}
//将GPU与主机同步
cudaStreamSynchronize(stream0);
cudaStreamSynchronize(stream1);
//停止计时器，收集性能数据
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(dev_a0);
cudaFree(dev_b0);
cudaFree(dev_c0);
cudaFree(dev_a1);
cudaFree(dev_b1);
cudaFree(dev_c1);
cudaStreamDestroy(stream0);
cudaStreamDestroy(stream1);
cudaEventDestroy(start);
cudaEventDestroy(stop);
system("pause");
return 0;
}
实际代码测试注意到，宽度优先的效率不及深度优先，个人猜想这可能是新版本的CUDA对其做了相关优化。
335
扫一扫
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
