如果share memory不够了怎么办

加速计算 CUDA
1

/****************************************************************
** 计算t1 t1 = D1.conj(H1) + D2.conj(H2); **
** 计算t2 t2 = abs(H1).^2+abs(H2).^2 + gama; **
************************************************************/
global void t1_t2(compx H1_d,compx H2_d,compx D1_d,compx D2_d,compx t1_d,double t2_d,int row,int col)
{
int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx = icol+j;
if(i<row&&j<col)
{

	t1_d[idx].real = (D1_d[idx].real * H1_d[idx].real + D1_d[idx].imag * H1_d[idx].imag)+
		D2_d[idx].real * H2_d[idx].real + D2_d[idx].imag * H2_d[idx].imag;
	t1_d[idx].imag = (H1_d[idx].real * D1_d[idx].imag - D1_d[idx].real * H1_d[idx].imag)+
		H2_d[idx].real * D2_d[idx].imag - D2_d[idx].real * H2_d[idx].imag;

	t2_d[idx] = (H1_d[idx].real * H1_d[idx].real + H1_d[idx].imag * H1_d[idx].imag) + 
		(H2_d[idx].real * H2_d[idx].real + H2_d[idx].imag * H2_d[idx].imag);
}

}
上面是我的kernel，其中D1 H1 都重复使用，想把它们放的share memory 可是share memory只有48k，上面这个kernel有没有好的方法可以优化呢？

2

楼主您好：

如果想将大于最大shared memory容量的数组强制放入shared memory,
这个真心无法放入。也是强人所难。

以及，您线程间对D1和H1中的元素都无重复使用（您的下标是[idx]), 您可能需要使用shared memory的迫切性真心不大。

您最应该解决的是您的下标问题，您的计算：
int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx = i*col+j;
的多个线程的idx之间严重的不连续，造成了巨大的读取上的效率降低。

这才是您真的应该考虑的因素。
（而不是去考虑Shared memory)。

3

版主我没看明白这里的idx为什么不是连续的，要怎么改才是连续的呢？谢谢哈

4

LZ您好：

您的写法是：

int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx = i*col+j;
一般而言此时是二维的block和grid。

CUDA中线程编号是x优先变化的，考虑最简单的情况，您的第一个block，上式化简为
int i=threadIdx.x;
int j=threadIdx.y;
考虑block一行中相邻的线程，j是相同的，i是连续变化的。
此时您的
int idx = i*col+j;
一行中相邻的两个线程的idx相差col，并且一般而言col是一个较大的数。

这样您一个warp中的各个线程的访存位置就相差col个int元素，这将引起非常严重的global 非合并访问，将极大地影响您的程序效能。

如果您修正了这个问题，（直接只考虑合并访问的话，写成j*col+i就可以，但是这么写可能和您的算法安排就不对应了，请您通盘考虑），那么按照您的写法，无需再使用shared memory。因为您每个线程都只是访问自己对应的数组中的那个元素，这个值会自动使用该线程的寄存器保存，以及这里并没有线程之间需要“shared data”的情况，因此您直接用即可。

祝您编码顺利~

5

谢谢版主哈，我忘记线程编号是按照x优先变化了。

6

不客气的，欢迎您常来~

祝您编码顺利~

7

我代码中的col=row=256，我的block大小为（16,16），grid大小为（16,16）我现在用idx=j*col+i之后DRAM Utilization 为52.8%(96.81Gb/s) 请问一下还能否提高呢？

8

请记录当前的运行时间，

然后在您的if(i < row && j < col)里面的第一句加上：
compx a = D1_d[idx];
compx b = H1_d[idx];
然后将下文的所有D1_d[idx]改成a,将所有的H1_d[idx]改成b，
重新编译并运行。

再次继续这次的运行时间。

然后请您比较下时间变化。
（感谢ICE提供建议）

9

确实提高了不少，现在的DRAM Utilization为67.8%（124.3GB），是不是这样修改后每个数据都是在register中，所以速度提高了呢？

10

嗯嗯。以及，建议您以时间为准，较好。

以及，您还可以考虑使用use_fast_math, 看看时间有无变化。

谢谢。