【CUDA 基础】6.2 并发内核执行

---

title: 【CUDA 基础】6.2 并发内核执行
categories:
- CUDA
- Freshman
tags:
- 流
- 事件
- 深度优先
- 广度优先
- 硬件工作队列
- 默认流阻塞行为
toc: true
date: 2018-06-18 22:04:08

Abstract: 本文介绍内核的并发执行，以及相关的知识
Keywords: 流，事件，深度优先，广度优先，硬件工作队列，默认流阻塞行为

开篇废话

没有废话，继续前面的内容，上文中我们说到了流，事件和同步等的概念，以及一些函数的用法，接下来的几个例子，介绍并发内核的几个基本问题，包括不限于以下几个方面：

• 使用深度优先或者广度优先方法的调度工作
• 调整硬件工作队列
• 在Kepler设备和Fermi设备上避免虚假的依赖关系
• 检查默认流的阻塞行为
• 在非默认流之间添加依赖关系
• 检查资源使用是如何影响并发的

非空流中的并发内核

本文我们开始使用NVIDIA提供的另一个可视化工具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);
}

四个核函数，N是100，tan计算在GPU中应该有优化过的高速版本，但是就算优化，这个也是相对耗时的，足够我们进行观察了。
接着我们按照上节课的套路，创建流，把不同的核函数或者指令放到不同的流中，然后看一下他们的表现。
本文完整的代码在github:https://github.com/Tony-Tan/CUDA_Freshman（欢迎随手star? ）

我们本章主要关注主机代码，下面是创建流的代码：

cudaStream_t *stream=(cudaStream_t*)malloc(n_stream*sizeof(cudaStream_t));
for(int i=0;i<n_stream;i++)
{
    cudaStreamCreate(&stream[i]);
}

首先声明一个流的头结构，是malloc的注意后面要free掉
然后为每个流的头结构分配资源，也就是Create的过程，这样我们就有n_stream个流可以使用了，接着，我们添加核函数到流，并观察运行效果

dim3 block(1);
dim3 grid(1);
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
for(int i=0;i<n_stream;i++)
{
    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]>>>();
}
cudaEventRecord(stop);
CHECK(cudaEventSynchronize(stop));
float elapsed_time;
cudaEventElapsedTime(&elapsed_time,start,stop);
printf("elapsed time:%f ms\n",elapsed_time);

这不是完整的代码，这个循环是将每个核函数都放入不同的流之中，也就是假设我们有10个流，那么这10个流中每个流都要按照上面的顺序执行这4个核函数。
注意如果没有

cudaEventSynchronize(stop)

nvvp将会无法运行，因为所有这些都是异步操作，不会等到操作完再返回，而是启动后自动把控制权返回主机，如果没有一个阻塞指令，主机进程就会执行完毕推出，这样就跟设备失联了，nvvp也会相应的报错。
然后我们创建两个事件，然后记录事件之间的时间间隔。这个间隔是不太准确的，因为是异步的。
运行结果如下：

使用nvvp检测，结果如下：

Fermi GPU 上的虚假依赖关系

虚假依赖我们在上文中讲到过了，这种情况通常出现在只有在比较古老的Fermi架构上出现，原因是其只有一个硬件工作队列，由于我们现在很难找到Fermi架构的GPU了，所以，只能看看书上给出的nvvp结果图了：

虚假依赖的问题我们在流和事件概述已经描述了引起此问题的理论原因，这里就不再解释了。
如果你手头只有老机器，这种虚假依赖关系也是可以解决的，原理就是使用广度优先的方法，组织各任务的方式如下：

// dispatch job with breadth first way
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]>>>();

这样逻辑图就不是:

而是

这样了，这就可以从抽象模型层面避免问题。
广度优先的nvvp结果是：

注意，以上结论都是我从书上原封不动弄下来的。

使用OpenMP的调度操作

完整内容 https://face2ai.com/CUDA-F-6-2-并发内核执行/