非 空 流 中 的 并 发 内 核
下面的示例演示了如何使用多个流并发运行多个核函数，也介绍了并发内核执行的几个基本问题，包括以下几个方面：使用深度优先或广度优先方法的调度工作；调整硬件工作队列；在Kepler设备和Fermi设备上避免虚假的依赖关系；检查默认流的阻塞行为；在非默认流之间添加依赖关系；检查资源使用是如何影响并发的。
#include <stdio.h>
#include <cuda_runtime.h>
#include <stdlib.h>
#define N 300000
#define NSTREAM 4
#define CHECK(call) \
{ \
const cudaError_t error = call; \
if (error != cudaSuccess) \
{ \
fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__); \
fprintf(stderr, "code: %d, reason: %s\n", error, \
cudaGetErrorString(error)); \
} \
}
//使用的核函数包括在设备上仿真有用工作的虚拟计算，这确保了内核驻留在GPU上的事件足够长
//以使重叠在可视化性能分析器中更加明显，主要是为了在nvvp中更容易将不同内核的执行进行可视化
__global__ void kernel_1()
{
double sum = 0.0;
for(int i = 0; i < N; i++)
{
sum = sum + tan(0.1) * tan(0.1);
}
}
__global__ void kernel_2()
{
double sum = 0.0;
for(int i = 0; i < N; i++)
{
sum = sum + tan(0.1) * tan(0.1);
}
}
__global__ void kernel_3()
{
double sum = 0.0;
for(int i = 0; i < N; i++)
{
sum = sum + tan(0.1) * tan(0.1);
}
}
__global__ void kernel_4()
{
double sum = 0.0;
for(int i = 0; i < N; i++)
{
sum = sum + tan(0.1) * tan(0.1);
}
}
int main(int argc, char **argv)
{
int n_streams = NSTREAM;
int isize = 1;
int iblock = 1;
int bigcase = 0;
// get argument from command line
if (argc > 1) n_streams = atoi(argv[1]);
if (argc > 2) bigcase = atoi(argv[2]);
float elapsed_time;
// set up max connectioin
char* iname = "CUDA_DEVICE_MAX_CONNECTIONS";
setenv (iname, "32", 1);
char *ivalue = getenv (iname);
printf ("%s = %s\n", iname, ivalue);
int dev = 0;
cudaDeviceProp deviceProp;
CHECK(cudaGetDeviceProperties(&deviceProp, dev));
printf("> Using Device %d: %s with num_streams=%d\n", dev, deviceProp.name,
n_streams);
CHECK(cudaSetDevice(dev));
// check if device support hyper-q
if (deviceProp.major < 3 || (deviceProp.major == 3 && deviceProp.minor < 5))
{
if (deviceProp.concurrentKernels == 0)
{
printf("> GPU does not support concurrent kernel execution (SM 3.5 "
"or higher required)\n");
printf("> CUDA kernel runs will be serialized\n");
}
else
{
printf("> GPU does not support HyperQ\n");
printf("> CUDA kernel runs will have limited concurrency\n");
}
}
printf("> Compute Capability %d.%d hardware with %d multi-processors\n",
deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount);
// 首先必须要创建一组非空流。在这组非空流中，发布每个流中的内核启动应该在GPU上同时运行
// 但是应不存在由于硬件资源限制而导致的虚假依赖关系
cudaStream_t *streams = (cudaStream_t *) malloc(n_streams * sizeof(
cudaStream_t));
for (int i = 0 ; i < n_streams ; i++)
{
CHECK(cudaStreamCreate(&(streams[i])));
}
// run kernel with more threads
if (bigcase == 1)
{
iblock = 512;
isize = 1 << 12;
}
// set up execution configuration
dim3 block (iblock);
dim3 grid (isize / iblock);
printf("> grid %d block %d\n", grid.x, block.x);
// 为了计算运行事件，创建了两个事件
cudaEvent_t start, stop;
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
// 记录start事件
CHECK(cudaEventRecord(start, 0));
// 使用一个循环遍历所有的流，这样内核在每个流中都可以被调度
// 这些内核启动的执行配置被指定为单一线程块中的单一线程，以保证有足够的GPU资源能并发运行所有的内核
//因为每个内核启动相对于主机来说都是异步的，所以可以通过使用单一主机线程同时调度多个内核到不同的流中
for (int i = 0; i < n_streams; i++)
{
kernel_1<<<grid, block, 0, streams[i]>>>();
kernel_2<<<grid, block, 0, streams[i]>>>();
kernel_3<<<grid, block, 0, streams[i]>>>();
kernel_4<<<grid, block, 0, streams[i]>>>();
}
// 记录stop事件
CHECK(cudaEventRecord(stop, 0));
CHECK(cudaEventSynchronize(stop));
// 计算运行时间
CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
printf("Measured time for parallel execution = %.3fs\n",
elapsed_time / 1000.0f);
// 释放所有流
for (int i = 0 ; i < n_streams ; i++)
{
CHECK(cudaStreamDestroy(streams[i]));
}
free(streams);
// 销毁事件
CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));
// reset device
CHECK(cudaDeviceReset());
return 0;
}
下图显示了在Tesla K40通过nvvp生成的时间轴，随着时间进度向右移动，每种颜色对应不同内核的执行，并且每行对应不同的流。正如期望的，在K40上可以看到4个并发内核在4个不同的流中执行。
Fermi GPU 上 的 虚 假 依 赖 关 系
为了掩饰虚假的依赖关系，可以在Fermi设备上运行相同的代码，simpleHyperq告诉我们，Fermi设备不支持Hyper-Q,而且内核最终会限制并发一起郧西你个。下图显示了和上图相同应用程序的时间轴，但不同的是它运行在Fermi GPU上。因为在Fermi设备上有虚假的依赖关系，所以4个流不能同时启动，这是由共享硬件工作队列造成的。为什么流i+1能够在流i开始其最后任务时开始它的第一个任务呢？因为两个任务是在不同的流中，所以它们之间没有依赖关系。当流i的最后一个任务被启动时，CUDA运行时从工作队列中调度下一个任务，这是流i+1的第一个任务。因为每个流的第一个任务不依赖于之前的任何任务，并且有可用的SM，所以它可以立即启动。之后，调度流i+1的第二个任务，然而他对第一个任务的依赖却阻止它被执行，这就会导致任务执行再次被阻塞。
这种虚假的依赖关系是由主机调度内核的顺序引起的。该应用程序使用深度有限的方法，在下一个流启动前，在该流中启动全系列的操作。利用深度优先方法得到的工作队列中的任务顺序如下图所示，由于所有流被多路复用到一个硬件工作队列中，所以前面的流就连续阻塞了后面的流。
在Fermi GPU上，为了避免虚假的依赖关系，可以用广度优先的方法从主机中调度工作。
for(int i = 0; i< n_streams; i++)
{
kernel_1<<<grid,block,0,streams[i]>>>();
}
for(int i = 0; i< n_streams; i++)
{
kernel_2<<<grid,block,0,streams[i]>>>();
}
for(int i = 0; i< n_streams; i++)
{
kernel_3<<<grid,block,0,streams[i]>>>();
}
for(int i = 0; i< n_streams; i++)
{
kernel_4<<<grid,block,0,streams[i]>>>();
}
采用广度优先顺序可以确保工作队列中相邻的任务来自于不同的流(如下图所示)。因此，任何相邻的任务对之间都不会再有虚假的依赖关系，从而得以实现并发内核执行。
此处的执行事件相比采用深度优先的方法提高了3倍。内核启动调用可以用nvvp来证实。下图展示了用广度优先方法的内核执行时间周：所有流同步启动。
使 用 OpenMP 的 调 度 操 作
在前面的例子中，使用单一的主机线程将异步CUDA操作调度到多个流中，接下来将使用多个主机线程将操作调度到多个流中，并使用一个线程来管理每一个流。OpenMP是CPU的并行编程模型，它使用编译器指令来识别并行区域。支持OpenMP指令的编译器可以将它们用作如何并行化应用程序的提示。用很少的操作，在主机上就可以实现多核并行。在使用OpenMP的同时使用CUDA，不仅可以提高便携性和生产效率，而且还可以提高主机代码的性能。具体实现如下：
omp_set_num_threads(n_streams);
#pragma omp parallel
{
int i = omp_get_thread_num();
kernel_1<<<grid,block,0,stream[i]>>>();
kernel_2<<<grid,block,0,stream[i]>>>();
kernel_3<<<grid,block,0,stream[i]>>>();
kernel_4<<<grid,block,0,stream[i]>>>();
}
OpenMP函数omp_set_num_threads用来指定在OpenMP并行区域里要用到的CPU核心的数量。编译器指令#pragma omp parallel将花括号之间的代码标记为并行部分。omp_get_thread_num函数为每个主机线程返回唯一一个线程ID，将该ID用作streams数组中的索引，用来创建OpenMP线程和CUDA流间的一对一映射。
什么时候从OpenMP中调度并行CUDA操作是有用的呢？在一般情况下，如果每个流在内核执行之前、期间或之后有额外的工作待完成，那么它可以包含在同一个OpenMP并行区域里，并且跨流和线程进行重叠。这样做更明显地说明了每个OpenMP线程的主机工作与同一个线程中启动的流CUDA操作是相关的，并且可以为了优化性能简化代码的书写。
用环境变量调整流行为
支持Hyper-Q的GPU在主机和每个GPU之间维护硬件工作队列，清除虚假的依赖关系。Kepler设备支持的硬件工作队列的最大数量是32。然而，默认情况下并发硬件连接的数量被限制为8。由于每个连接都需要额外的内存和资源，所以设置默认的限制为8，减少了不需要全部32个工作队列的应用程序的资源消耗。可以使用CUDA_DEVICE_MAX_CONNECTIONS环境变量来调整并行硬件连接的数量，对于Kepler设备而言，其上限是32.
有几种设置该环境变量的方法。在Linux中，可以根据shell的版本，通过以下代码进行设置，对于Bash和Bourne Shell，其代码如下：export CUDA_DEVICE_MAX_CONNECTIONS=32。对于C-Shell，其代码如下：setenv CUDA_DEVICE_MAX_CONNECTIONS 32。这个环境变量也可以直接在C主机程序中进行设定：setenv("CUDA_DEVICES_MAX_CONNECTIONS","32",1);
每个CUDA流都会映射到单一的CUDA设备连接中。如果流的数量超过了硬件连接的数量，多个流将共享一个连接。当多个流共享相同的硬件工作队列时，可能会产生虚假的依赖关系。在支持Hyper-Q技术但是没有足够硬件连接的平台上，要将simpleHyperqDepth示例修改为使用8个CUDA流：#define NSTREAM 8。并将CUDA设备连接的数量设置为4：char *iname = "CUDA_DEVICE_MAX_CONNECTIONS"; setenv(iname,"4",1);nvvp显示如下，展示了8个流，但是只有4路并发。因为现在只有4个设备连接，两个流共享一个队列。采用深度优先的方法调度内核，导致了分配载同一工作队列中的两个流之间出现了虚假的依赖关系，这与在Fermi GPU上使用深度优先顺序时的结果类似。
下一步，使用相同的设置，检查使用广度优先的行为，如下图所示，现在8个流都是同步运行。用广度优先顺序调度内核去除了虚假的依赖关系。
GPU 资 源 的 并 发 限 制
有限的内核资源可以抑制应用程序中可能出现的内核并发的数量。在之前的例子中，启动内核时只有一个线程，以避免并发时任何的硬件限制。因此，每个内核执行只需要少量的设备计算资源。kernel_1<<<1,1,0，streams[i]>>>();在实际应用中，内核启动时通常会创建多个线程。通常，会创建数百或数千个线程。有了这么多线程，可用的硬件资源可能会成为并发的主要限制因素，因为它们阻止启动符合条件的内核。为了在活动中观察到这个行为，可以在simpleHyperqBreath例子中改变配置，在每个块中使用多个线程，在每个网格中使用更多的块：dim3 block(128); dim3 grid(32);然后将使用的CUDA流的数量增加到16：#define NSTREAM 16。nvvp上结果如下图所示，图中只实现了8路并发，即使CUDA设备连接的数量被设置为32。因为GPU无法分配足够的资源来执行所有符合条件的内核，所以并发性是有限的。
默 认 流 的 阻 塞 行 为
为了说明默认流在非空流中是如何阻塞操作的，在simpleHyperqDepth中，将深度优先调度循环改为在默认流中调用kernel_3。
for(int i = 0; i < n_streams; i++)
{
kernel_1<<<grid,block,0,streams[i]>>>();
kernel_2<<<grid,block,0,streams[i]>>>();
kernel_3<<<grid,block>>>();
kernel_4<<<grid,block,0,streams[i]>>>();
}
因为第三个内核在默认流中被启动，所以在非空流上所有之后的操作都会被阻塞，知道默认流中的操作完成。下图显示了这段代码运行的时间轴，它是在Tesla K40上使用nvvp得到的。这个时间轴显示了每个kernel_3启动是如何阻止其他阻塞流中进一步执行的。
创 建 流 间 依 赖 关 系
在理想情况下，流之间不应该有非计划之内的依赖关系(即虚假的依赖关系)。然而，在复杂的应用程序中，引入流间依赖关系是很有用的，它可以在一个流中阻塞操作直到另一个流中的操作完成。事件可以用来添加流间依赖关系。
假如我们想让一个流中的工作在其他所有流中的工作都完成后才开始执行，那么就可以使用事件来创建流之间的依赖关系。首先，将标志设置为cudaEventDisableTiming，创建同步事件，代码如下：
cudaEvent_t *kernelEvent = (cudaEvent_t *)malloc(n_streams * sizeof(cudaEvent_t));
for(int i = 0; i < nstreams; i++)
{
cudaEventCreateWithFlags(&kernelEvent[i],cudaEventDisableTiming);
}
接下来，使用cudaEventRecord函数，在每个流完成时记录不同的事件。然后，使用cudaStreamWaitEvent使最后一个流(即streams[n_streams - 1])等待其他所有流。
for(int i = 0; i < n_streams; i++)
{
kernel_1<<<grid,block,0,streams[i]>>>();
kernel_2<<<grid,block,0,streams[i]>>>();
kernel_3<<<grid,block,0,streams[i]>>>();
kernel_4<<<grid,block,0,streams[i]>>>();
cudaEventRecord(kernelEvent[i],streams[i]);
cudaStreamWaitEvent(streams[n_streams - 1],kernelEvent[i],0);
}
nvvp结果图如下，要注意，第四个流，在其他所有流完成后才能开始启动工作。
1101
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
