Global Load Efficiency 低下

我的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.x
blockDim.x+threadIdx.x;
int j=blockIdx.y
blockDim.y+threadIdx.y;
int idx=j
col+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);
}

}

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

LZ您好:

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

祝您好运~

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

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

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

感谢来访。

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;

}

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;

}

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中?

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

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

一个问题一个问题的来。

使用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;
这样才可以的。

您说呢?

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

这里的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%

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

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

(以及,您确定是global load efficiency而不是global store efficiency? 不要打错!)

以及,需要为您指出的是:

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

(只是建议,您保留着也可以的)

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

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

不过,您先尝试一下?

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

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

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

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

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

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

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

没有更多的建议了。。。

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

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

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

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