Global Load Efficiency 低下

system · 2013 年10 月 31 日 01:27

我的kernel如下，profiler结果显示Global Load Efficiency 只有50% 我的理解是在对P1_d赋值的时候，把实部和虚部分开了，所以没能合并内存访问，是这样的吗？如果要提高效率该怎么做呢？
global void computer_P1_with_P2(compx *P1_d,compx *P2_d,double *pupil_d,const double zercoef_d,double wf65_d,double theta_d,int row,int col,int high)
{
int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx=jcol+i;

if(i<row&&j<col)
{
	float pupil = pupil_d[idx];
	float wavefront1 = 0;

	//#pragma unroll
	for(int k=2;k<Zs;k++)
	{
		wavefront1 += wf65[k*col*row+idx] * zercoef[k];
	}

	float wavefront2 = wavefront1 + theta_d[idx];

	P1_d[idx].real = pupil * __cosf(wavefront1);
	P1_d[idx].imag = pupil * __sinf(wavefront1);
	P2_d[idx].real = pupil * __cosf(wavefront2);
	P2_d[idx].imag = pupil * __sinf(wavefront2);
}

}

system · 2013 年10 月 31 日 03:18

是不是要重新改造一下compx结构体呢？原来的结构体
typedef struct
{
double real; //实部
double imag;//虚部
}compx;
我的应用是256X256的图像所以要定compx data[256][256]大小的数据，这种方式应该是AoS，如果把结构体改成
typedef struct
{
double real[256]; //实部
double imag[256];//虚部
}compx;
SoA这种结构体，是不是符合合并内存访问呢？

system · 2013 年10 月 31 日 04:15

LZ您好：

将复数的实部和虚部分开存放的话，确实可以立即实现合并访问，但是这样需要修改以前的数据结构。
您前面那种“交错式”的存储方法，有时可以通过一些技巧实现合并访问，正在寻找相关资料，稍后为您继续回答。

祝您好运~

system · 2013 年10 月 31 日 05:15

楼主也不用考虑神马SOA, AOS的，

您可以直接用double2来读取，然后这样即可简单的充分合并了。原本的compx的结构无需做任何更改。
（请注意的是，因为对齐要求，只对您cudaMalloc出来的显存有效。如果您是自定义的分配器，则可能不能用。但一般无问题。适用99.9%的人）

读取完毕后，请直接在2个分量上运算（相当于您原来的real和imag分量）

感谢来访。

system · 2013 年10 月 31 日 06:11

if(i<row&&j<col)
{
float pupil = pupil_d[idx];
float wavefront1 = 0;
double2 p1,p2;

//#pragma unroll
for(int k=2;k<Zs;k++)
{
wavefront1 += wf65[kcolrow+idx] * zercoef[k];
}

float wavefront2 = wavefront1 + theta_d[idx];

P1_d[idx].real = pupil * __cosf(wavefront1);
P1_d[idx].imag = pupil * __sinf(wavefront1);
P2_d[idx].real = pupil * __cosf(wavefront2);
P2_d[idx].imag = pupil * __sinf(wavefront2);

P1_d[idx] = p1;
P2_d[idx] = p2;

}

system · 2013 年10 月 31 日 06:11

if(i<row&&j<col)
{
float pupil = pupil_d[idx];
float wavefront1 = 0;
double2 p1,p2;

//#pragma unroll
for(int k=2;k<Zs;k++)
{
wavefront1 += wf65[kcolrow+idx] * zercoef[k];
}

float wavefront2 = wavefront1 + theta_d[idx];

P1_d[idx].real = pupil * __cosf(wavefront1);
P1_d[idx].imag = pupil * __sinf(wavefront1);
P2_d[idx].real = pupil * __cosf(wavefront2);
P2_d[idx].imag = pupil * __sinf(wavefront2);

P1_d[idx] = p1;
P2_d[idx] = p2;

}

system · 2013 年10 月 31 日 06:13

if(i<row&&j<col)
{
float pupil = pupil_d[idx];
float wavefront1 = 0;
double2 p1,p2;

//#pragma unroll
for(int k=2;k<Zs;k++)
{
wavefront1 += wf65[kcolrow+idx] * zercoef[k];
}

float wavefront2 = wavefront1 + theta_d[idx];

p1.real = pupil * __cosf(wavefront1);
p1.imag = pupil * __sinf(wavefront1);
p2.real = pupil * __cosf(wavefront2);
p2.imag = pupil * __sinf(wavefront2);

P1_d[idx] = p1;
P2_d[idx] = p2;

}
请问一下版主，这样是不是就合并访问了呢？p1 p2 是不是在register中呢？如果不是怎么可以保证其在register中？

system · 2013 年10 月 31 日 06:15

p1.x= pupil * __cosf(wavefront1);
p1.y= pupil * __sinf(wavefront1);
p2.x= pupil * __cosf(wavefront2);
p2.y= pupil * __sinf(wavefront2);

上面这里有点小问题改正一下

system · 2013 年10 月 31 日 06:30

一个问题一个问题的来。

使用double2的确可以解决global load/store efficiency的问题。

但是你下文需要做一些修改：
P1_d[idx].real => p1.x
P1_d[idx].imag => p1.y
P2_d[idx].real => p2.x
P2_d[idx].imag => p2.y

最后还得再合并写入，改成：
((double2 *)P1_d)[idx] = p1;
((double2 *)P2_d)[idx] = p2;
这样才可以的。

您说呢？

system · 2013 年10 月 31 日 06:32

以及，您这是Store效率的问题所在，而不是load的。

system · 2013 年10 月 31 日 06:47

这里的double2 和我自己定义的compx 其实是一样的。我把修改后的代码附上
if(i<row&&j<col)
{
float pupil = pupil_d[idx];
float wavefront1 = 0;

	double wf65[15];
	double zercoef[15];
	compx P1;
	compx P2;

	//#pragma unroll
	for(int k=2;k<Zs;k++)
	{
		wf65[k]=wf65_d[k*col*row+idx];
		zercoef[k]=zercoef_d[k];
		wavefront1 += wf65[k] * zercoef[k];
	}

	float wavefront2 = wavefront1 + theta_d[idx];

	P1.real = pupil * __cosf(wavefront1);
	P1.imag = pupil * __sinf(wavefront1);
	P2.real = pupil * __cosf(wavefront2);
	P2.imag = pupil * __sinf(wavefront2);

	((compx *)P1_d)[idx] = P1;
	((compx *)P2_d)[idx] = P2;
}

可是Global Load Efficiency 还是50%

system · 2013 年10 月 31 日 06:55

首先楼主您不应该怀疑我，

让您使用double2，而不是使用您看起来一模一样的您的compx自然有道理的。

（以及，您确定是global load efficiency而不是global store efficiency? 不要打错！）

system · 2013 年10 月 31 日 06:57

以及，需要为您指出的是：

你的wf65_d读取到wf65的行为（以及zercoef的读取）纯粹是多此一举，
建议去掉。

（只是建议，您保留着也可以的）

system · 2013 年10 月 31 日 07:13

我查了一下double2是做了对齐处理的，这算不算你推荐我用double2的原因呢？
上面是我打错了，不好意思是global store efficiency。

system · 2013 年10 月 31 日 07:16

double2的确有对齐要求，可能是这个要求导致了它比你的compx好，让我推荐了它。

不过，您先尝试一下？

system · 2013 年10 月 31 日 07:24

那我需要把我程序所有的compx 改成 double2 结构了

system · 2013 年10 月 31 日 08:00

您可以把您前面compx定义加以修改，改成double typedef得到，这样就不用修改后面的代码了，您可以快速搞定修改。

祝您好运~
——————————————————————————————————————
修正一下，刚才没能注意到您的compx里面的元素是real和imag，而不是x和y，这样就不行了，您不能直接typedef了，请修改您的代码。

刚才一时没有看仔细，深表歉意。

system · 2013 年10 月 31 日 08:26

既然楼主不接受前文我们给给出的使用double2的建议，

那么楼主能否重新定义下贵Struct？
至少改成例如：
struct align(16) compx
{
double real;
double imag;
};
这样您只需要改动一下定义即可，可能会消除您的50%问题。

不知道您是否愿意尝试下？

没有更多的建议了。。。

system · 2013 年10 月 31 日 10:17

出现了如下问题
error C2719: “xxx”: 具有 __declspec(align(‘16’)) 的形参将不被对齐
正在尝试解决…

system · 2013 年10 月 31 日 11:29

这是最后一次警告了，您已经多次无视了我们的意见。

请您仔细看看18#的发文，里面是让你将__align__(16)写到参数列表吗！！！
是吗！！！！看仔细了！！别上来责备论坛！！

请按提示操作！
请你保持起码的一点尊重！！！