用户
 找回密码
 立即注册
quanzhang100 该用户已被删除
发表于 2013-10-31 09:27:31
3632430
我的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);
        }

}

使用道具 举报 回复
发表于 2013-10-31 11:18:47
是不是要重新改造一下compx结构体呢?原来的结构体
typedef struct
{
        double real; //实部
        double imag;//虚部
}compx;
我的应用是256X256的图像 所以要定compx data[256][256]大小的数据,这种方式应该是AoS,如果把结构体改成
typedef struct
{
        double real[256]; //实部
        double imag[256];//虚部
}compx;
SoA这种结构体,是不是符合合并内存访问呢?
使用道具 举报 回复 支持 反对
发表于 2013-10-31 12:15:39
quanzhang100 发表于 2013-10-31 11:18
是不是要重新改造一下compx结构体呢?原来的结构体
typedef struct
{

LZ您好:

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

祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-10-31 13:15:29
ice 发表于 2013-10-31 12:15
LZ您好:

将复数的实部和虚部分开存放的话,确实可以立即实现合并访问,但是这样需要修改以前的数据结构 ...

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

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

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

感谢来访。
使用道具 举报 回复 支持 反对
发表于 2013-10-31 14:11:36
横扫千军 发表于 2013-10-31 13:15
楼主也不用考虑神马SOA, AOS的,

您可以直接用double2来读取,然后这样即可简单的充分合并了。原本的com ...

      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[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);

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

        }
使用道具 举报 回复 支持 反对
发表于 2013-10-31 14:11:36
横扫千军 发表于 2013-10-31 13:15
楼主也不用考虑神马SOA, AOS的,

您可以直接用double2来读取,然后这样即可简单的充分合并了。原本的com ...

      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[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);

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

        }
使用道具 举报 回复 支持 反对
发表于 2013-10-31 14:13:48
横扫千军 发表于 2013-10-31 13:15
楼主也不用考虑神马SOA, AOS的,

您可以直接用double2来读取,然后这样即可简单的充分合并了。原本的com ...

      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[k*col*row+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中?
使用道具 举报 回复 支持 反对
发表于 2013-10-31 14:15:03
                p1.x= pupil * __cosf(wavefront1);
                p1.y= pupil * __sinf(wavefront1);
                p2.x= pupil * __cosf(wavefront2);
                p2.y= pupil * __sinf(wavefront2);

上面这里有点小问题 改正一下
使用道具 举报 回复 支持 反对
发表于 2013-10-31 14:30:42
quanzhang100 发表于 2013-10-31 14:13
if(i

一个问题一个问题的来。

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

您说呢?


使用道具 举报 回复 支持 反对
发表于 2013-10-31 14:32:15
玫瑰幻想 发表于 2013-10-31 14:30
一个问题一个问题的来。

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

以及,您这是Store效率的问题所在,而不是load的。
使用道具 举报 回复 支持 反对
1234下一页
发新帖
您需要登录后才可以回帖 登录 | 立即注册