GPU并行运算与CUDA编程--优化篇

一般有三大瓶颈：内存带宽受限、指令吞吐受限、延迟受限

1.内存带宽受限

优化方式一：
用其他内存分担压力，如：TEX/Shared Memory/Constant Memory
优化方式二：
改变访问顺序，降低上一级内存的cache miss，缓解当前内存的压力
优化方式三：
用算法压缩数据/改变数据访问方式，降低不必要的数据访问

图1.1GPU内存层次结构图 上图表示GPU中内存的整体层次，由上往下访问速度依次变慢。CPU传输到GPU中的数据受限会被存储到DRAM，GPU在访问数据时会从上往下访问，如果遇到cache miss，就会向下一级访问。

Texture cache的利用

__ldg()指定只读缓存

__global__ void vecAdd(int n, float* a, float* b, float* c)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;//线程的实际ID
	if (i < n)
	{
		c[i] = __ldg(&a[i]) * __ldg(&b[i]);
		c[i] = pow(c[i],1000);
	}
}

对于只读数据，添加上__ldg()之后，GPU便可以直接从更快的texture缓存中读取。

Shared Memory的利用

SM可以让同一个block中的所有线程自由访问，如果多个线程需要访问同一组数据，可以考虑把数据存放在SM中再进行计算。
使用关键字__shared__即可在GPU中指定一定大小的SM

Constant cache的利用

使用关键字__constant__申明GPU中一个constant cache
使用cudaMemcpyToSymbol(cc名称, ……)进行拷贝

2.指令吞吐受限

1.使用更快的指令

对于立即数“1.0”在默认情况下是双精度的浮点数，加上后缀“f”后，“1.0f”就会被编译为单精度浮点数，计算时更加快速。（sm_53之后的也支持半精度half2）

2.使用intrinsic function

针对于一些比较简单的数学函数：
fun()适用于对精度要求比较高的情况，但速度慢
__fun()精度比较低，但速度快

3.减少Bank conflict

32个bank组成了shared memory，当不同的线程访问同一个bank的不同word时，就会出现bank conflict（如上图）。
当不同线程访问同一个bank的同一word时，不会出现bank conflict（如下图）。

当出现bank conflict时，就会增加GPU的replay，这是需要尽量减少的。

4.减少warp里的指令发散

一个warp执行的时间 = warp里不同分支执行时间之和
多个warp执行多个分支是对整体的性能没有影响的。

3.延迟受限型

1.增加active warp数量

通过增加warp来隐藏指令或者data I/O所带来的延迟。

Occupancy

用于衡量active warp
计算公式为 occupancy = active warp数/SM中所允许的最大warp数
提升occupancy的方法：
1、一个warp尽量少用SM/寄存器资源
2、降低每个block占用的SM
3、降低每个block占用的线程

achieved occupancy：
实际上，由于有些warp执行时间比其他warp长，某些时候只有少数warp运行，这样延迟就比较高。

2.从延迟源头解决

指令的读取/执行的相关性/同步/数据请求/texture等等

利用CUDA的分析工具，可以得到以下的一些参数，锁定延迟的源头。