一个使用share memory进行性能优化的实例

下面是这段代码是我的一个算法中用来求和以及求平方和的kernel函数：

__global__ static void CompSumAndSquare(int rate,int I_n,int size,int width,int wsize,int *image,float *sum,float *sumOfSquare)

{

    int x,y,p,i,j,ix,k;

    const int pos = threadIdx.x+blockIdx.x*blockDim.x;

    const int max_dimension=blockDim.x*gridDim.x;

 

    k=wsize/2;

    for(i=0;i<rate;i++)

    {

        p=pos+i*max_dimension;     

        if(p<size){

            x=p%I_n+wsize/2;

            y=p/I_n;

            ix=y*width+x;

            sum[ix]=0;

            sumOfSquare[ix]=0;

            for(j=-k;j<=k;j++)

            {

                sum[ix]+=(float)image[ix+j];               

                sumOfSquare[ix]+=pow((float)image[ix+j],2);

            }          

        }

    }

}

这个函数的运行时间约为12ms，其中窗口大小为15，测试平台为nGeforce 9800GT，由于这个函数占据了我整个算法时间的一半以上。所以本人在无聊之余将累加和平方和运算屏蔽一个，看单独运算的时间，结果这个函数运行的时间为3ms多一点。于是我觉得这个函数的性能瓶颈应该是global memory访问的方式而非计算指令的吞吐量，然后做了下面一点改动。

 

改动部分：

            for(j=-k;j<=k;j++)

            {

                sum[ix]+=(float)image[ix+j];               

            //  sumOfSquare[ix]+=pow((float)image[ix+j],2);

            }

            for(j=-k;j<=k;j++)

            {

            //  sum[ix]+=(float)image[ix+j];               

                sumOfSquare[ix]+=pow((float)image[ix+j],2);

            }

这个版本函数运行的时间结果约为7ms，相比较第一版本有些提高，但还是很慢。随后想到了share memory作为中间运算的载体，使用后的代码如下，改动之后的函数时间约为1.65ms，还不错哈！改动的代码如下：

__global__ static void CompSumAndSquare(int rate,int I_n,int size,int width,int wsize,int *image,float *sum,float *sumOfSquare)

{

    int x,y,p,i,j,ix,k,t;

    const int pos = threadIdx.x+blockIdx.x*blockDim.x;

    const int tid = threadIdx.x;

    const int max_dimension=blockDim.x*gridDim.x;

    k=wsize/2;

    extern __shared__ float s_m[];

    float *s_q=(float*)&s_m[wsize*THREADNO_VAR];

    int N;

 

    for(i=0;i<rate;i++)

    {

        p=pos+i*max_dimension;     

        if(p<size){

            x=p%I_n+k;

            y=p/I_n;

            ix=y*width+x;

            N=tid*wsize;

            s_m[N]=0;       

            for(j=-k,t=0;j<=k;j++,t++)

            {

                s_m[N+t]=(float)image[ix+j];           

            }

            s_q[N]=pow(s_m[N],2);

            for(j=1;j<wsize;j++)

            {

                s_m[N]+=s_m[N+j];              

                s_q[N]+=pow(s_m[N+j],2);

            }

            sum[ix]=s_m[N];            

            sumOfSquare[ix]=s_q[N];

        }

    }

}

 

我在做这个简单的性能优化过程时，觉得最让自己满意的是这种思考方式，首先找到算法的瓶颈，然后猜测性能不好的原因，起初我以为计算速度是瓶颈，想用快速函数，当通过一个小的测试发现存储器的访问方式才是问题所在（因为计算规模减少一半，但时间却减少四分之三），最后想到使用高速存储器share memory作为中间计算载体，得到比较好的程序性能。