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


通信方式

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

Map

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

Gather

gather
多对一的方式。

Scatter

scatter
一对多的方式。

Stencil

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

Transpose

transpose
转置操作,改变形状、顺序等。
transpose2
进行合理的顺序改变在数据读取速度上会提升速度。


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):
$$\frac{\text{math we do}}{\text{memory we access}}$$
从分子上,要最大化每个线程上的操作;从分母上,要减小对内存访问的时间。
这就要求需要把频繁读取的数据放在更快的memory上。
从速度上来说,mem的排序为:
local > shared >> global >> host mem
这是因为local存在寄存器或L1缓存中。

Coalesce

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


atomic 操作

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


divergence发散

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


总结

sum


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 定义的定值)。
    为了取得某一维度的线程号,可以使用

    1
    int i = blockIdx .x * blockDim .x + threadIdx .x;
  3. 最大化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的数目!
    在官方的示例中:

    1
    2
    3
    4
    5
    6
    7
    8
    9
    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倍。
当然在图像处理中线程也要注意访问越界的问题,得到线程负责的像素位置之后,先进行判断,越界退出即可:

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