【CUDA 基础】5.3 减少全局内存访问

发表于 | 分类于 ， | 评论数： | 阅读次数：

Abstract: 本文介绍使用共享内存进行归约，并比较全局内存归约与共享内存归约之间的性能差距
Keywords: 共享内存，归约

减少全局内存访问

逻辑是非常重要的，一旦你学会了逻辑，很多假的东西你可以轻松的识别出来，这会使你更加强大而不会被任何人或者组织洗脑。
废话少说，开始今天的博客。
使用共享内存的主要原因就是减少对全局内存的访问，来减少不必要的延迟，第三章我们学过了归约，可以参考：

这两篇博客包含我们前面使用全局内存进行归约的各种技术，我们几天也要用其中一部分代码作为比较，来体现我们共享内存的优势。
我们要集中解决下面两个问题：

  1. 如何重新安排数据访问模式以避免线程束分化
  2. 如何展开循环以保证有足够的操作使指令和内存带宽饱和

本文我们通过对比研究前面的部分代码，来分析为何要使用共享内存，以及如何使用共享内存。

使用共享内存的并行归约

我们首先来回忆全局内存下的，完全展开的归约计算：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
__global__ void reduceGmem(int * g_idata,int * g_odata,unsigned int n)
{
	//set thread ID
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockDim.x*blockIdx.x+threadIdx.x;
	//boundary check
	if (tid >= n) return;
	//convert global data pointer to the
	int *idata = g_idata + blockIdx.x*blockDim.x;

	//in-place reduction in global memory
	if(blockDim.x>=1024 && tid <512)
		idata[tid]+=idata[tid+512];
	__syncthreads();
	if(blockDim.x>=512 && tid <256)
		idata[tid]+=idata[tid+256];
	__syncthreads();
	if(blockDim.x>=256 && tid <128)
		idata[tid]+=idata[tid+128];
	__syncthreads();
	if(blockDim.x>=128 && tid <64)
		idata[tid]+=idata[tid+64];
	__syncthreads();
	//write result for this block to global mem
	if(tid<32)
	{
		volatile int *vmem = idata;
		vmem[tid]+=vmem[tid+32];
		vmem[tid]+=vmem[tid+16];
		vmem[tid]+=vmem[tid+8];
		vmem[tid]+=vmem[tid+4];
		vmem[tid]+=vmem[tid+2];
		vmem[tid]+=vmem[tid+1];

	}

	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];

}

下面这步是计算当前线程的索引位置：

1
unsigned int idx = blockDim.x*blockIdx.x+threadIdx.x;

当前线程块对应的数据块首地址

1
int *idata = g_idata + blockIdx.x*blockDim.x;

然后是展开循环的部分，tid是当前线程块中线程的标号，主要区别于全局编号idx：

1
2
3
4
5
6
7
8
9
10
11
12
if(blockDim.x>=1024 && tid <512)
    idata[tid]+=idata[tid+512];
__syncthreads();
if(blockDim.x>=512 && tid <256)
    idata[tid]+=idata[tid+256];
__syncthreads();
if(blockDim.x>=256 && tid <128)
    idata[tid]+=idata[tid+128];
__syncthreads();
if(blockDim.x>=128 && tid <64)
    idata[tid]+=idata[tid+64];
__syncthreads();

这一步把是当前线程块中的所有数据归约到前64个元素中，接着使用如下代码，将最后64个元素归约成一个

1
2
3
4
5
6
7
8
9
10
if(tid<32)
{
    volatile int *vmem = idata;
    vmem[tid]+=vmem[tid+32];
    vmem[tid]+=vmem[tid+16];
    vmem[tid]+=vmem[tid+8];
    vmem[tid]+=vmem[tid+4];
    vmem[tid]+=vmem[tid+2];
    vmem[tid]+=vmem[tid+1];
}

注意这里声明了一个volatile变量，如果我们不这么做，编译器不能保证这些数据读写操作按照代码中的顺序执行（参考5.1中关于编译器数据传输部分的说明），所以必须要这么做。
然后我们执行以下这段代码，虽然前面执行过了，我们还是执行以下，观察下结果：
完整的可执行代码依旧在GitHub上可以找到
github:https://github.com/Tony-Tan/CUDA_Freshman点个星星也不会很累，拜托啦😆

re-1

这里我们假装看不到别的，只看我们的核函数，可见器质性时间是4.25ms左右
然后我们对上面的代码进行改写，改写成共享内存的版本，来看代码：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
__global__ void reduceSmem(int * g_idata,int * g_odata,unsigned int n)
{
	//set thread ID
    __shared__ int smem[DIM];
	unsigned int tid = threadIdx.x;
	//unsigned int idx = blockDim.x*blockIdx.x+threadIdx.x;
	//boundary check
	if (tid >= n) return;
	//convert global data pointer to the
	int *idata = g_idata + blockIdx.x*blockDim.x;

    smem[tid]=idata[tid];
	__syncthreads();
	//in-place reduction in global memory
	if(blockDim.x>=1024 && tid <512)
		smem[tid]+=smem[tid+512];
	__syncthreads();
	if(blockDim.x>=512 && tid <256)
		smem[tid]+=smem[tid+256];
	__syncthreads();
	if(blockDim.x>=256 && tid <128)
		smem[tid]+=smem[tid+128];
	__syncthreads();
	if(blockDim.x>=128 && tid <64)
		smem[tid]+=smem[tid+64];
	__syncthreads();
	//write result for this block to global mem
	if(tid<32)
	{
		volatile int *vsmem = smem;
		vsmem[tid]+=vsmem[tid+32];
		vsmem[tid]+=vsmem[tid+16];
		vsmem[tid]+=vsmem[tid+8];
		vsmem[tid]+=vsmem[tid+4];
		vsmem[tid]+=vsmem[tid+2];
		vsmem[tid]+=vsmem[tid+1];

	}

	if (tid == 0)
		g_odata[blockIdx.x] = smem[0];

}

唯一的不同就是多了一个共享内存的声明，以及各线程将全局写入共享内存，以及后面的同步指令：

1
2
smem[tid]=idata[tid];
__syncthreads();

这一步过后同步保证该线程块内的所有线程，都执行到此处后继续向下进行，这是可以理解的，因为我们的归约只针对本块内，当然如果想跨几个块执行，可能同步这里就有问题了，这个是上一节课要讨论的，这里就不过多解释了，我们接着就看到一个volatile类型的指针，指向共享内存，对最后64个归约结果进行归约，整个过程和全局内存一毛一样，只不过一个在全局内存操作，一个在共享内存操作，得到相同的结果，我们也来看一下运行结果。

re-2
还是那张图，对比来看，速度提高了不少了。看一下

1
2
gld_transactions
gst_transactions

这两个指标的结果

re-3

可以看出使用共享内存的，比使用全局内存的高到不知道哪里去了

使用展开的并行归约

可能看到上面的截图你已经知道我接下来要并行4块了，对于前面说的，使用共享内存不能并行四块，是因为没办法同步读四个块，这里我们还是用老方法进行并行四个块，就是在写入共享内存之前进行归约，4个块变成一个，然后把这一个存入共享内存，进行常规的共享内存归约:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
__global__ void reduceUnroll4Smem(int * g_idata,int * g_odata,unsigned int n)
{
	//set thread ID
    __shared__ int smem[DIM];
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockDim.x*blockIdx.x*4+threadIdx.x;
	//boundary check
	if (tid >= n) return;
	//convert global data pointer to the
    int tempSum=0;
	if(idx+3 * blockDim.x<=n)
	{
		int a1=g_idata[idx];
		int a2=g_idata[idx+blockDim.x];
		int a3=g_idata[idx+2*blockDim.x];
		int a4=g_idata[idx+3*blockDim.x];
		tempSum=a1+a2+a3+a4;

	}
    smem[tid]=tempSum;
	__syncthreads();
	//in-place reduction in global memory
	if(blockDim.x>=1024 && tid <512)
		smem[tid]+=smem[tid+512];
	__syncthreads();
	if(blockDim.x>=512 && tid <256)
		smem[tid]+=smem[tid+256];
	__syncthreads();
	if(blockDim.x>=256 && tid <128)
		smem[tid]+=smem[tid+128];
	__syncthreads();
	if(blockDim.x>=128 && tid <64)
		smem[tid]+=smem[tid+64];
	__syncthreads();
	//write result for this block to global mem
	if(tid<32)
	{
		volatile int *vsmem = smem;
		vsmem[tid]+=vsmem[tid+32];
		vsmem[tid]+=vsmem[tid+16];
		vsmem[tid]+=vsmem[tid+8];
		vsmem[tid]+=vsmem[tid+4];
		vsmem[tid]+=vsmem[tid+2];
		vsmem[tid]+=vsmem[tid+1];

	}

	if (tid == 0)
		g_odata[blockIdx.x] = smem[0];

}

这段代码就是多了其他三块的求和：

1
2
3
4
5
6
7
8
9
10
11
12
13
unsigned int idx = blockDim.x*blockIdx.x*4+threadIdx.x;
//boundary check
if (tid >= n) return;
//convert global data pointer to the
int tempSum=0;
if(idx+3 * blockDim.x<=n)
{
    int a1=g_idata[idx];
    int a2=g_idata[idx+blockDim.x];
    int a3=g_idata[idx+2*blockDim.x];
    int a4=g_idata[idx+3*blockDim.x];
    tempSum=a1+a2+a3+a4;
}

这一步在3.5中已经介绍过了为什么能加速了，因为可以通过增加三步计算而减少之前的3个线程块的计算，这是非常大的减少。同时多步内存加载也可以使内存带宽达到更好的使用。
结果可想而知：

re-4

re-5

吞吐量指标：

1
nvprof  --metrics  dram_read_throughput  ./reduce_integer_shared_memory

re-6

无论是指标还是运行速度，都有非常显著的提升。
我们这里总结下展开的优势：

  • I/O得到了更多的并行，就是我上面说的更好的利用带宽，增加了吞吐量
  • 全局内存存储事务减少到 $\frac{1}{4}$ ，这个主要针对最后一步，将结果存入全局内存
  • 整体性能巨幅提升

使用动态共享内存的并行归约

然后我们看一下动态版本，其实动态版本没啥可看的，只是写法上有点不同，把宏改成核函数配置参数，注意其单位是字节就好，也就是不要忘了sizeof()就行了。
这里不啰嗦了。

有效带宽

对比一下数据，回顾一下我们的有效带宽，其计算公式(4.4中有详细介绍)：
$$
有效带宽=\frac{(读字节数 + 写字节数)\times 10^{-9}}{运行时间}\tag{1}
$$
可以研究三个核函数的有效带宽，这里就不再一个个计算了，因为这个在4.4中已经教大家做了，我们可以自己算一下，并得出结论

总结

本文主要高速大家如何使用共享内存加速归约，以及结合了共享内存的展开能更高的提高效率，注意线程块内的同步，这个是重要的。

相关文章
0%