system
1
我的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
2
是不是要重新改造一下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
3
LZ您好:
将复数的实部和虚部分开存放的话,确实可以立即实现合并访问,但是这样需要修改以前的数据结构。
您前面那种“交错式”的存储方法,有时可以通过一些技巧实现合并访问,正在寻找相关资料,稍后为您继续回答。
祝您好运~
system
4
楼主也不用考虑神马SOA, AOS的,
您可以直接用double2来读取,然后这样即可简单的充分合并了。原本的compx的结构无需做任何更改。
(请注意的是,因为对齐要求,只对您cudaMalloc出来的显存有效。如果您是自定义的分配器,则可能不能用。但一般无问题。适用99.9%的人)
读取完毕后,请直接在2个分量上运算(相当于您原来的real和imag分量)
感谢来访。
system
5
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
6
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
7
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
8
p1.x= pupil * __cosf(wavefront1);
p1.y= pupil * __sinf(wavefront1);
p2.x= pupil * __cosf(wavefront2);
p2.y= pupil * __sinf(wavefront2);
上面这里有点小问题 改正一下
system
9
一个问题一个问题的来。
使用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
10
以及,您这是Store效率的问题所在,而不是load的。
system
11
这里的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
12
首先楼主您不应该怀疑我,
让您使用double2,而不是使用您看起来一模一样的您的compx自然有道理的。
(以及,您确定是global load efficiency而不是global store efficiency? 不要打错!)
system
13
以及,需要为您指出的是:
你的wf65_d读取到wf65的行为(以及zercoef的读取)纯粹是多此一举,
建议去掉。
(只是建议,您保留着也可以的)
system
14
我查了一下double2是 做了对齐处理的,这算不算你推荐我用double2的原因呢?
上面是我打错了,不好意思 是global store efficiency。
system
15
double2的确有对齐要求,可能是这个要求导致了它比你的compx好,让我推荐了它。
不过,您先尝试一下?
system
16
那我需要把我程序所有的compx 改成 double2 结构了
system
17
您可以把您前面compx定义加以修改,改成double typedef得到,这样就不用修改后面的代码了,您可以快速搞定修改。
祝您好运~
——————————————————————————————————————
修正一下,刚才没能注意到您的compx里面的元素是real和imag,而不是x和y,这样就不行了,您不能直接typedef了,请修改您的代码。
刚才一时没有看仔细,深表歉意。
system
18
既然楼主不接受前文我们给给出的使用double2的建议,
那么楼主能否重新定义下贵Struct?
至少改成例如:
struct align(16) compx
{
double real;
double imag;
};
这样您只需要改动一下定义即可,可能会消除您的50%问题。
不知道您是否愿意尝试下?
没有更多的建议了。。。
system
19
出现了如下问题
error C2719: “xxx”: 具有 __declspec(align(‘16’)) 的形参将不被对齐
正在尝试解决…
system
20
这是最后一次警告了,您已经多次无视了我们的意见。
请您仔细看看18#的发文,里面是让你将__align__(16)写到参数列表吗!!!
是吗!!!!看仔细了!!别上来责备论坛!!
请按提示操作!
请你保持起码的一点尊重!!!