二维数组在存储时会被扁平化处理,所以在程序中,最好是将多为数组扁平化为等价的一维数组。CUDA对于多为数组也是按行优先存储。
算法与硬件越是契合,算法的执行效率,硬件的使用率越高。所以有两个选择
- 可以为当前算法设计特殊的处理架构,DSA(Domain Specific Architecture)如Google TPU。
- 根据当前硬件精心设计算法。
向量数据类型
同时适用于Host和Device,通过
make_<type name>
来构造,如int2 i2 = make_int2(3,4)
:i2向量含有3和4两个元素。
float4 f4 = make_float4(1.0f, 2.0f, 3.0f, 4.0f)
:f4是含有4个元素的数组。访问方式如下:
int x = i2.x; int y = i2.y;
CUDA程序调试和开发
- 可以使用
ssh
登路远程含有CUDA enabled GPU的服务器。 - 通常使用双GPU的系统开发CUDA程序,一个GPU负责显示,另一个负责计算,可以使用
Nsight
等工具。 - 只有一个GPU时,在Linux系统中可以关闭桌面环境(释放桌面环境对GPU的占用),只在命令行中使用
CUDA-gdb
调试。(实际上,实验阶段不关闭桌面环境,也是可以正确执行的)
- 可以使用
CUDA开发的任务
有效的数据并行算法 + 针对GPU架构特性的优化 = 获得并行最优性能
OpenCL
OpenCL使用起来繁琐,而且运行速度远远低于CUDA运行速度。OpenCL与CUDA的主要功能有着十分相似之处,一个CUDA程序员很容易掌握OpenCL编程。
half-warp
截图
Streams-流
任务并行(Task Parallelism),不同于在大量的数据上执行相同的任务(SIMD),而是同时执行多个不同的任务。
CUDA数据并行原语
啥是原语,就是这个领域的基本操作。CUDA并行原语库(CUDA Data Parallel Primitive Library, CUDPP). 含有并行前缀和,并行排序,并行规约等。这些原语为许多数据并行算法提供了重要基础。如果你正在编写某个复杂算法,那么CUDPP很可能已经提供了这个算法。
CUDPP 下载地址:http://code.google.com/p/cudpp
for more information:CUDA By Example 178页
想清楚一个问题
如果使用shared memory,每个block对应自己的shared memory,当这个block中threads的ID更新后,这个block并没有再被分配新的shared memory. 也就是说如果使用两个blocks, 两个blocks有两段shared memory, block中threads 的ID更新后, 所计算结果会写入这个thread所在block的shared memory中. thread ID变了, 但所属的block不变.
理解SIMD,SIMT
Single Instruction Multiple Data(Thread).
- Single:相同的操作,kernel函数只有一个
- Instruction:kernel函数所做的事情
- Multiple:所处理的数据量大,要拆分为一批一批
- Data:大的数据量
- Threads:一个SP上的大量thread超快速切换,获得延时隐藏
warp中的divergence
已经知道,在一个warp中,所有的threads执行相同的指令,但是如果指令中含有条件分支语句,很大程度上会发生divergence。比如
优化并行归约
中的描述:一个宿舍的6个学生可以是一个warp,今天有的想先上厕所,后吃饭,而有的不需要上厕所,此时所有的同学都会一起先上厕所,后一起吃饭。总之,分支发散使得性能明显下降。但是,注意divergence只发生在一个warp中。
可以综合算法的上下文,将divergence的粒度变为32(warp的大小)的倍数,从而避免warp内的分支发散。比如:
1
2
3
4
5
6
7
8
9
10__global__ void kernel(float* c){
int tid = threadIdx.x + blockDim.x * blockIdx.x;
float a=0;
float b=0;
if (tid%2 == 0){
c[tid] = 100;
}else{
c[tid] = 200;
}
}偶数id的thread把100写入偶数编号的地址,奇数id的thread将200写入奇数编号的地址。会发生分支发散。如果将分支的粒度定为32,则没有了divergence [代码中tid%32 == 0有问题]:
1
2
3
4
5
6
7
8
9
10__global__ void kernel(float* c){
int tid = threadIdx.x + blockDim.x * blockIdx.x;
float a=0;
float b=0;
if ( tid%32 == 0 ){
c[tid] = 100;
}else{
c[tid] = 200;
}
}但是两者的结果的是不同的,实际中,需要根据算法上下文考虑结合此方法。
使用
brand_efficiency
指标来查看divergence的情况:$ nvprof --mereics brand_efficiency ./out
, 但是CUDA编译器会进行优化,将短的,有条件的代码段的断定指令取代了分支指令。所以,会看到虽然代码中有分支,却显示分支效率100%。
数据预读取?
在一次读取global memory的操作和使用这个数据之间,插入独立于以上数据的操作,可以隐藏访问延迟。如:
1
2
3float m = dev_a[i]; // 1. 从global memory中读取
float f = a * b; // 2. 与m无关的操作
float f2 = m*f; // 3. 使用m分析:在warp切换中,
指令优化
GPU中执行指令时很快速的,所以通常不用太在意指令的优化。优化顺序一般是存储优化,后执行配置优化,最后可以考虑指令优化。
比如:除以
2^n
,使用移位操作>>n
,以2^n
求模,使用&(2^n - 1)
;避免从double到float的自动转换,float a = 0.0
,使用float a = 0.0f
。还比如,两种运行时数学库函数的取舍,精度高的速度低,精度低的速度高。使用
-use-fast-math
编译选项后,强制将速度慢的func()
转化为速度快的__func()
。还比如循环展开。