map很好理解,其实就是映射,也就是输入和输出一一对应,一个萝卜一个坑
Gather中文名为收集,是将若干个输入数据经过计算后得到一个输出值,如图左示。很典型的应用就是比如说对于一个图像,我们需要每一个像素值是其四周像素的平均值。
具体应用实例如下:
在C语言中,加入我们定义了如上图示的一个结构体,包含float和int两种变量,然后我们又定义了一个该结构体的变量数组,一般来说其在内存中是像上面那样排列的,强迫症看起来是不是不舒服,而且这种排列方式比较浪费空间,所以通过转置后形成下面的排列方式后既美观又使运算加速了,岂不美哉?
如图示,GPU由若干个SM(Stream Multiprocessor流多处理器)组成,而每个SM又包含若干个SP(教材上是Stream Processor流处理器,改视频中是simple processor),anyway...开心就好,管他叫什么名字~
GPU的作用是负责分配线程块在硬件SM上运行,所有SM都以并行独立的方式运行。
下面做一下题目吧:
解析:
另外需要注意的是程序员负责定义线程块,而GPU则负责管理硬件,因此程序员不能指定线程块的执行顺序,也不能指定线程块在某一特定的 SM上运行。
这样设计的好处如下:
有如上好处的同时,自然也就有局限性:
如图示
访问速度:
local memory > shared memory > global
例题:
解析:
s,t,u是本地内存中的变量,所以t=s最先运行,同理可以排除其他代码运行顺序。
注意:这只是为了说明访问速度出的例题,实际情况中,编译器可能会做出相应的调整来达到我们的目的
说道线程,很自然我们就需要考虑同步。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();
在cuda编程中经常会碰到这样的情况,即大量的线程同时都需要对某一个内存地址进行读写操作,很自然这会发生冲突,如下图示:
下面是发生冲突的具体的代码示例:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#define NUM_THREADS 10000
#define ARRAY_SIZE 10
#define BLOCK_WIDTH 100
void 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);
}
运行结果:
使用原子操作也是有一定限制的,如下:
提高CUDA编程效率策略
前面提到了很多优化策略是集中在memory上的,把数据尽可能放到更快地内存上去,其中内存速度是 local > share > global
如图是线程发散的主要场景,即if else语句,上图右边非常生动的展现了线程发散的情形,可以看到各个线程在碰到if条件句后开始发散,最后聚合,但是最后各个线程之间的编号还是保持原来的不变的,这就是线程发散。
下面举一个更加极端的例子,就是循环语句,如下图示:
可以看到有蓝、红、绿、紫四个线程同时运行,蓝线程只循环了一次,其他线程循环次数都多于蓝线程,当蓝线程退出循环后就不得不一直等着其他线程,上图左下角的示意图可以很直观的看到这大大降低了运行效率,这也是为什么我们需要避免线程发散。