share memory问题

加速计算 CUDA
1

大家好，我想对如下核函数
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.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx=icol+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.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx=icol+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);
}
}

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

2

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

int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.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]值。

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

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

谢谢。

3

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

4

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