一、parallel communication patterns 并行通信模式
Map:映射,在特定的位置读取和写入。
Gather:收集,从多个不同的位置读入,写入一个位置。
Scatter:分发,写入多个位置。
Transpose转置
结构数组缩写为AOS,数组结构缩写为SOA
转置运算是指任务重新排序内存中的数据元素
Stencil模板
能够在输出结果为每一个元素生成一个结果
二、流处理器Stream Multiprocessors
所有在同一个线程块中的线程运行在同一个SM上
在一个内核中的所有块,在下个内核中任意块启动前都已完成
三、共享内存
共享内存是在一个块的线程之间共享
CUDA中的线程协作主要是通过共享内存实现的。使用关键字“__share__”声明共享变量,将使这个变量驻留在共享内存中,该变量具有以下特征:
- 位于线程块的共享存储器空间中
- 与线程块具有相同的生命周期
- 仅可通过块内的所有线程访问
对于GPU上启动的每个线程块,CUDA C编译器都将创建该变量的一个副本。 线程块中的每个线程都共享这块内存,但线程却无法看到也不能修改其他线程块的变量副本。 这就使得一个线程块中的多个线程能够在计算上进行通信和协作。而且,共享内存缓冲区驻留在物理GPU上,在访问共享内存时的延迟要远远低于访问普通缓冲区的延迟,使得共享内存的访问非常高效。
线程同步机制“__syncthreads()”
关键字“__share__”只是声明了共享变量,位于同一个线程块中的不同线程都可以访问该变量,如果没有同步机制,将会发生竞态条件 (Race Condition),导致错误的运行结果。
CUDA确保同步的方法是调用“__syncthreads()”。__syncthreads()将确保线程块中的每个线程都执行完 __syncthreads()前面的语句后,才会执行下一条语句。
一个线程可以访问局部内存、一个线程块的共享内存和所有线程都可以访问的全局内存
速度方面:local>shared>global>cpu
四、屏障
__syncthreads()的线程需要其他可以到达该点的线程,而不是等待块内所有其他线程
一个高端的GPU每秒可以进行超过3万亿次的数学计算
五、Global Memory
如果想操作全局变量必须传递一个指向该内存的指针,如:
诀窍:只可以在内核中传递局部变量,然后传入一个指针,要在内核外分配和初始化全局变量