cuda 学习 | GPU硬件与并行通信模式

---

通信方式

通信方式主要以课程截图为主……

Map

这是一种一一对应的方式。

Gather

多对一的方式。

Scatter

一对多的方式。

Stencil

模板，多对多的方式。
图中左中为输入，左下为输出，不同颜色为不同线程的读取、输出位置。

Transpose

转置操作，改变形状、顺序等。

进行合理的顺序改变在数据读取速度上会提升速度。

---

GPU结构

从大到小来说，结构为：
Kernel -》 Block -》 Thread
硬件上GPU有很多流处理器（streaming multiprocessors），GPU给这些sm分配block，同一个sm可能运行多个block。
GPU的优势在于多个线程同时工作，但缺点是无法知道这些线程执行的先后顺序。
可以确定的有：

1. 线程同时在同一个SM运行
2. 所有kernel中的blocks在下个kernel启动前结束

内存关系

每个Thread都有自己的local mem。
每个Block有sheared mem，其中的Thread均可以访问。声明时需要加上__shared__前缀。
global mem可以被不同Block的访问。

---

Barrier同步

为了避免因Thread执行不同步带来的数据存储问题，有时候需要等待所有的Thread都执行结束再继续执行。
这时需要在kernel中加入__syncthreads();进行同步。
比如在分线程读取同一数组其他位置并修改当前位置时，需要在读取后同步并存入临时变量。在写操作后也要同步。
对于两个kernel间可以不加同步，默认存在隐式同步。

---

优化

GPU的优化目标是吞吐量，及最大化计算强度（Maximize arithmetic intensity）：

math we domemory we access

从分子上，要最大化每个线程上的操作；从分母上，要减小对内存访问的时间。
这就要求需要把频繁读取的数据放在更快的memory上。
从速度上来说，mem的排序为：
local > shared >> global >> host mem
这是因为local存在寄存器或L1缓存中。

Coalesce

在读取数据时，GPU对于连续的数据有高效的访问效率。

---

atomic 操作

为了避免内存访问读取冲突，也可以使用cuda中的atomic操作。
比如对于相加的操作，可使用：
atomicAdd(&g[i],1)代替g[i]=g[i]+1。
不过这种操作存在局限，只支持部分运算类型和部分数据类型，并且会降低运算速度。

---

divergence发散

发散出现在内核中非顺序的操作，如判断、循环。这会导致不同线程运行的时间明显不同。加入同步会让所有线程完全结束后继续执行，耗费时间较多。

---

总结

---

homework

这次作业比上次麻烦点，不过也有很多收获。

1. 图像处理中统一的行列定义
   作业里recombineChannels函数是已经给定的，需要从这里得到有关的行列信息。
   而我一开始定义的gridSize把行列的顺序定义反了。
2. grid、block、thread的区别
   官网上有相关的介绍：Kernel。
   另外Value of threadidx.x (.y, .z), blockidx.x etc. in CUDA也给出了blockDim、gridDim、threadIdx、blockIdx之间的关系。
   Dim是定值，Idx是变量，一般来说0 <= Idx < Dim = Dim_define(是通过dim3 定义的定值)。
   为了取得某一维度的线程号，可以使用

int i = blockIdx .x * blockDim .x + threadIdx .x;

1. 最大化thread以提高速度
   同样的总线程数运行的速度一样吗？
   之前我是认为一样的，是因为我觉得不同块（block）间的线程（thread）也是同步执行的，只要<<< numBlocks，threadsPerBlock >>>中两者之积相同即可。但经过实验，发现其实差距很大。
   在dim3 numBlocks(numCols, numRows)，dim3 threadsPerBlock(1)的情况下，速度大约为58ms；
   在dim3 numBlocks(numCols, numRows)，dim3 threadsPerBlock(2)的情况下，速度大约为31ms；
   在dim3 numBlocks(numCols, numRows)，dim3 threadsPerBlock(4)的情况下，速度大约为17ms；
   在dim3 numBlocks(numCols, numRows)，dim3 threadsPerBlock(4，4)的情况下，速度大约为8ms；
   在dim3 numBlocks(numCols/4+1, numRows/4+1)，dim3 threadsPerBlock(4，4)的情况下，速度大约为5ms；
   在dim3 numBlocks(numCols/32+1, numRows/32+1)，dim3 threadsPerBlock(32，32)的情况下，速度大约为1.4ms。
   由此可见，不同block之间并不是一种完全的并行关系，加快速度一定要加大thread的数目！
   在官方的示例中：

int main() {     ...     // Kernel invocation     dim3 threadsPerBlock(16, 16);    dim3 numBlocks(N / threadsPerBlock.x + 1, N / threadsPerBlock.y + 1);//官方没有+1操作    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);    ...}

线程数是固定的256（16*16），而块数是可变的，根据数据大小不同而进行变化。既然加大线程可以加快速度，把所有的都用上不就可以了吗？实际上试验中最下面的dim3 threadsPerBlock(32，32)已经达到了目前线程上限1024，速度的确较一开始的58ms快了40倍。
当然在图像处理中线程也要注意访问越界的问题，得到线程负责的像素位置之后，先进行判断，越界退出即可：

if (thread_2D_pos.x >= numCols || thread_2D_pos.y >= numRows)    return;