share memory问题

大家好,我想对如下核函数
global void computer_P1_with_P2(compx *P1,compx *P2,double *pupil,const double zercoef,double wf65,double theta,int row,int col,int high)
{
int i=blockIdx.x
blockDim.x+threadIdx.x;
int j=blockIdx.y
blockDim.y+threadIdx.y;
int idx=i
col+j;

if(i<row&&j<col)
{
double wavefront1 = 0;
for(int k=2;k<Zs;k++)
{
wavefront1 += wf65[kcolrow+idx] * zercoef[k];
}
double wavefront2 = wavefront1 + theta[idx];

P1[idx].real = pupil[idx] * cos(wavefront1);
P1[idx].imag = pupil[idx] * sin(wavefront1);
P2[idx].real = pupil[idx] * cos(wavefront2);
P2[idx].imag = pupil[idx] * sin(wavefront2);
}
}

优化为
global void computer_P1_with_P2(compx *P1,compx *P2,double *pupil,const double zercoef,double wf65,double theta,int row,int col,int high)
{
int i=blockIdx.x
blockDim.x+threadIdx.x;
int j=blockIdx.y
blockDim.y+threadIdx.y;
int idx=i
col+j;

shared double zercoef_s[Zs];
if(idx<Zs)
zercoef_s[idx] = zercoef[idx];
__syncthreads();

if(i<row&&j<col)
{
double wavefront1 = 0;
for(int k=2;k<Zs;k++)
{
wavefront1 += wf65[kcolrow+idx] * zercoef_s[k];
}
double wavefront2 = wavefront1 + theta[idx];

P1[idx].real = pupil[idx] * cos(wavefront1);
P1[idx].imag = pupil[idx] * sin(wavefront1);
P2[idx].real = pupil[idx] * cos(wavefront2);
P2[idx].imag = pupil[idx] * sin(wavefront2);
}
}

可是结果不对了,请问问题出在什么地方呢?

楼主您的问题可能出现在这里:

int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.y
blockDim.y+threadIdx.y;
int idx=i*col+j;

这样的idx计算,在一个block里,idx是不连续的。
例如您用(100,100) * (16,16)启动了kernel, 每个block这样有16 * 16 = 256个线程,

但是他们的编号不是N, N+1, N+2… N + 255这样的顺序的。
而是N, N + col, N + 2 *col, N + 3 * col ,… N + 255 * col,
这导致他们是不连续的(临近2个线程的编号相差了很多)。

所以您导致下文您的:
zercoef_s[idx] = zercoef[idx];
也将在zercoef_s上不连续的写入一批数据(中间有没赋值的空洞!)

自然您下文的对:
wavefront1 += wf65[kcolrow+idx] * zercoef_s[k];
(k是连续的从2到Zs)
会读取到一些没有初始化过的zercoef_s[k]值。

没有初始化的值参与运算,自然结果是错误的了。

请您修正您的错误的下标计算。

谢谢。

谢谢版主哈,我把idx=icol+j 改成 idx=jcol+i速度明显提高了哈,我还想问一下我这里的sin cos这些函数怎么快速的sin cos函数呢? 是不是在项目属性里面 选择Use Fast Math选项呢?

是这样的,但是这样也同时带来了些许的计算结果精度的下降。您可以评估下下降是否可以忍受,才能决定是否要用他们。