本小节笔记大纲：

• 1.Communication patterns
  • gather,scatter,stencil,transpose
• 2.GPU hardware & Programming Model
  • SMs,threads,blocks,ordering
  • Synchronization
  • Memory model: local, shared, global
  • Atomic Operation
• 3.Efficient GPU Programming
  • Access memory faster
    • coalescing global memory
    • use faster memory
  • Avoid thread divergence

一、Communication Patterns

1.Patterns

• Map

map很好理解，其实就是映射，也就是输入和输出一一对应，一个萝卜一个坑

• Gather

Gather中文名为收集，是将若干个输入数据经过计算后得到一个输出值，如图左示。很典型的应用就是比如说对于一个图像，我们需要每一个像素值是其四周像素的平均值。

• Scatter
  scatter的特点是每个线程一次会向内存输出多个值，也可能多个线程向一个内存输出值。

• Stencil
  Stencil表示模板的意思，所以也就是计算的时候用模子来选择输入数据，看下图就清楚了

• Transpose
  其实就是转置啦~

具体应用实例如下：

在C语言中，加入我们定义了如上图示的一个结构体，包含float和int两种变量，然后我们又定义了一个该结构体的变量数组，一般来说其在内存中是像上面那样排列的，强迫症看起来是不是不舒服，而且这种排列方式比较浪费空间，所以通过转置后形成下面的排列方式后既美观又使运算加速了，岂不美哉？

2.练习题

• 第一个很简单就是map，不仔细解释了
• 第二个个表达式我之前脑袋一热就选了C。。但是要注意，scatter的特点是每个线程一次会向内存输出多个值，这显然不符合该特点，而应该是Transpose。
• 第三个就是scatter了，原因如上
• 最后一个很容易选stencil，但是你要注意if条件语句的限制，所以应该是Gather。

3.总结神图

二、GPU Hardware

1.问题导向

• 线程是如何有效地一致访问内存
  • 子话题：如何利用数据重用
• 线程如何通过共享内存通信部分结果

2.硬件组成

如图示，GPU由若干个SM(Stream Multiprocessor流多处理器)组成，而每个SM又包含若干个SP(教材上是Stream Processor流处理器，改视频中是simple processor),anyway...开心就好，管他叫什么名字~

GPU的作用是负责分配线程块在硬件SM上运行，所有SM都以并行独立的方式运行。

下面做一下题目吧：

解析：

• 1正确.一个线程块包含许多线程
• 2正确.一个SM可能会运行多个多个线程块
• 3错误，因为一个线程块无法在一个以上的SM上运行
• 4正确，在一个线程块上所有线程有可能配合起来解决某个子问题
• 5错误，一个SM上可能有多个线程块，但是根据定义，线程和不同的线程块不应该存在协作关系。

3.程序员与GPU分工

另外需要注意的是程序员负责定义线程块，而GPU则负责管理硬件，因此程序员不能指定线程块的执行顺序，也不能指定线程块在某一特定的
SM上运行。

这样设计的好处如下：

• 硬件可以运行的更加有效率
• 运行切换不需要等待，一旦一个线程块运行完毕，SM可以自动的将另一个线程块加载进来
• 最大的优势：可扩展性，因为可以自动分配硬件资源，所以向下到单个SM，上到超级计算机的大量SM，均可以很好的适应。

有如上好处的同时，自然也就有局限性：

• 对于哪个块在哪个SM上运行无法进行任何假设
• 无法获得块之间的明确的通信

4.GPU Memory Model

如图示

• 每个线程都有它自己的本地内存(local memory)
• 线程块有一个共享内存(shared memory)，块中所有线程都可以访问该内存中的数据
• GPU中的全局内存(global memory)是所有线程块中的线程都能访问的内存，也是CPU进行数据传递的地方。

访问速度：

local memory > shared memory > global

例题：

解析：

s,t,u是本地内存中的变量，所以t=s最先运行，同理可以排除其他代码运行顺序。

注意：这只是为了说明访问速度出的例题，实际情况中，编译器可能会做出相应的调整来达到我们的目的

5.Sychronization

说道线程，很自然我们就需要考虑同步。GPU中的同步有如下几种：

Barrier(屏障)

顾名思义，就是所有线程运行到这个点都需要停下来。

如图示，红色、蓝色、绿色代表的线程先后到达barrier这个时间点后都停下来进行同步操作，完成之后线程的执行顺序是不一定的，可能如图示蓝色线程先执行，绿色，红色紧随其后。

另外其实还有一种隐式的barrier，比如说先后启动kernel A和kernel B，一般来说kernel B执行之前kernel A肯定是执行完毕了的。

说了这么多来做下题吧~233

题目：如下图示，现在需要实现一个数组前移的操作，即后面一个往前面挪，共享数组大小是128，问为实现这个功能，需要设置几次同步操作(或者说需要设置几个barrier?)

解析：
最开始的时候没想明白，写了127,128，但是都不对。后来听解释才明白。前移操作可以分为三步：

• 为每个数组元素赋值，即

array[idx] = threadIdx.x;
__syncthreads(); # 128个线程都执行完赋值语句后才能进行下一步

• 读取后面一个元素的值，存在临时变量里

int temp = array[idx+1];
__syncthreads();

• 将后一元素的值往前移

array[idx] = temp;
__syncthreads();

6.Atomic Memory Operation

在cuda编程中经常会碰到这样的情况，即大量的线程同时都需要对某一个内存地址进行读写操作，很自然这会发生冲突，如下图示:

下面是发生冲突的具体的代码示例：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>#define NUM_THREADS 10000
#define ARRAY_SIZE 10
#define BLOCK_WIDTH 100void printDevice();__global__ void increment_naive(int *g)
{int i = blockIdx.x * blockDim.x + threadIdx.x;i = i % ARRAY_SIZE;g[i] = g[i] + 1;
}int main(int argc, char **argv)
{printDevice();printf("\n");int h_array[ARRAY_SIZE];const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);int *d_array;// 分配内存cudaMalloc((void **) &d_array, ARRAY_BYTES);cudaMemset((void *) d_array, 0, ARRAY_BYTES);increment_naive<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);cudaMemcpy(h_array, d_array, ARRAY_BYTES, cudaMemcpyDeviceToHost);for(int i=0; i<ARRAY_SIZE; i++){printf("%d:%d\n",i,h_array[i]);}// 释放内存cudaFree(d_array);getchar();//CUDA_SAFE_CALL(cudaGetDeviceCount(&deviceCount));return 0;
}

运行结果：(每次运行的结果是不确定的)

__global__ void increment_atomicNaive(int *g)
{int i = blockIdx.x * blockDim.x + threadIdx.x;i = i % ARRAY_SIZE;atomicAdd(&g[i], 1);
}