CUDA——SM中warp调度器调度机制&&访存延迟隐藏

SM中warp调度器调度机制&&访存延迟隐藏

核函数中并不是所有线程一起启动执行的，核函数的执行是以线程束(warps)作为单位，warps的执行由warp调度器进行调度，一个调度器只能调度一个warp去执行指令，一个warp里的所有线程几乎是同时执行的。
以一个warps调度器为例子：
假设一个核函数开启了128个线程，那么其被划分成4个warps。

1、调度器调度warp0执行指令。

2、warp0挂起(在进行访存)，调度器调度warp1执行指令

3、以此类推，只要有处于Ready Queue的warps,且当前的运算单元没被占用时，warp调度器就会调用Ready Queue的warps去执行指令。直到warps都挂起。

4、若此时无warp在进行运算，且所有的warp都处于挂起状态。(假设SM上其它的块也挂起)，那么就会产生访存延迟。

就是出现了都在访存，而没有进行运算的情况，这就浪费了很多时间。因此，如果一个warp的访存周期T里刚好是warp处理指令时间周期t的k倍，那么只需要k个warp即可隐藏访存带来的延迟。当第k个warp运算完之后，第一个warp也就刚好访存完毕，又处于Ready Queue状态，可随时被warp调度器调度去干活，没有浪费一丝运算的时间，这样就能使得程序的性能得到提升。

5、如何解决访存延迟？

通过以上几点，我们明确了访存延迟出现的原因。访存是有先后顺序的(SM、调度器是有限的)。
（1）应该让先访存完毕的warp去执行尽可能多的指令(不然运算单元空着也是浪费啊)，去隐藏其它warp的访存时间。
（2）增加active warps的数量，让尽可能多的warp去隐藏访存延迟。（这个有局限性，warps是有限的）

6、现实中如何编写？

a、可以先写一个核函数，测试一下数据访存所花的时间，记下时间T1。

__global__ void obj_norm(float *input, float *output, size_t data_size)
{
	unsigned int ix = threadIdx.x + blockIdx.x*blockDim.x;
	unsigned int iy = threadIdx.y + blockIdx.y*blockDim.y;
	unsigned int i = ix + iy * blockDim.x*gridDim.x;

	if (i < data_size)
	{
		output[i] = input[i];
		//float value = input[i];
	}
}

b、将核函数具体执行的代码写下来，记下时间T2，若指令执行的时间可以隐藏访存时间。
那么，T2≈T1。此时，可以看作访存的时间全用来进行指令执行操作了。

__global__ void obj(float *input, float *output, size_t data_size)
{
	unsigned int ix = threadIdx.x + blockIdx.x*blockDim.x;
	unsigned int iy = threadIdx.y + blockIdx.y*blockDim.y;
	unsigned int i = ix + iy * blockDim.x*gridDim.x;

	if (i < data_size)
	{
		output[i] = __cosf((__logf(input[i])));
		//float value = (__logf(input[i])) / 24.5f;
	}
}

c、编写核函数应该尽量满足T1≈T2，最大化去隐藏访存时间。若T2远大于T1，则考虑将指令挪到其它“未饱和”的核函数(不能够完全隐藏访存时间的核函数)内执行。

d、当访存时间很大，且无法优化时，要用足够多的指令或者增加active warps的数量去隐藏访存时间。（一般写核函数肯定是需要对数据进行运算处理的，一定会有加减乘除以及一些其它的指令的）。

最后

GPU在硬件上并不是严格的并行，但通过warp的调度，块在SM上的调度等等流水线模式，使得在软件层面看起来是并行的。