通信方式
通信方式主要以课程截图为主……
Map
这是一种一一对应的方式。
Gather
多对一的方式。
Scatter
一对多的方式。
Stencil
模板,多对多的方式。
图中左中为输入,左下为输出,不同颜色为不同线程的读取、输出位置。
Transpose
转置操作,改变形状、顺序等。
进行合理的顺序改变在数据读取速度上会提升速度。
GPU结构
从大到小来说,结构为:
Kernel -》 Block -》 Thread
硬件上GPU有很多流处理器(streaming multiprocessors),GPU给这些sm分配block,同一个sm可能运行多个block。
GPU的优势在于多个线程同时工作,但缺点是无法知道这些线程执行的先后顺序。
可以确定的有:
- 线程同时在同一个SM运行
- 所有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
这次作业比上次麻烦点,不过也有很多收获。
- 图像处理中统一的行列定义
作业里recombineChannels
函数是已经给定的,需要从这里得到有关的行列信息。
而我一开始定义的gridSize
把行列的顺序定义反了。 - 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;
- 最大化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;