c++ - cuda代码的优化技巧

Question

我在 MATLAB 中编写了一段用于计算自商形象 (SQI) 的代码。现在我想并行重写它的一部分以加快速度。这部分代码是:

siz=15;
X=normalize8(X);
[a,b]=size(X);
filt = fspecial('gaussian',[siz siz],sigma);
padsize = floor(siz/2);
padX = padarray(X,[padsize, padsize],'symmetric','both');

t0 = tic; % -------------------------------------------------------------
Z=zeros(a,b);
for i=padsize+1:a+padsize
    for j=padsize+1:b+padsize
        region = padX(i-padsize:i+padsize, j-padsize:j+padsize);
        means= mean(region(:));
        M=return_step(region, means);
        filt1=filt.*M;

        summ=sum(sum(filt1));        

        filt1=(filt1/summ);
        Z(i-padsize,j-padsize)=(sum(sum(filt1.*region))/(siz*siz));
    end
end
toc(t0) % -------------------------------------------------------------

和 return_step 函数:

function M=return_step(X, means)

[a,b]=size(X);
for i=1:a
    for j=1:b
        if X(i,j)>=means
            M(i,j)=1;
        end
    end
end

我在下面写了内核函数:

__global__ void returnstep(const double* x, double* m, double* filt, int leng, double mean, int i, int j, int width)
{
    int idx=threadIdx.y*blockDim.x+threadIdx.x;
    if(idx>=leng) return;

    int ridx= (j+threadIdx.y)*width+threadIdx.x+i;
    double xval= x[ridx];
    if (xval>=mean) m[idx]=filt[idx]*xval;
    else            m[idx]=0;
}

然后将 MATLAB 代码更改如下:

kernel= parallel.gpu.CUDAKernel('returnstep.ptx', 'returnstep.cu');
kernel.ThreadBlockSize= [double(siz) double(siz) 1];
GM = gpuArray(zeros(siz,siz));
GpadX = gpuArray(padX);
Gfilt = gpuArray(filt);

%% Process image
t0 = tic; % -------------------------------------------------------------
Z=zeros(a,b);
for i=padsize+1:a+padsize
    for j=padsize+1:b+padsize
        means= mean(region(:));
        GM= feval(kernel, GpadX, GM, Gfilt, siz*siz, means, i-padsize-1, j-padsize-1, padXwidth);
        filt1=  gather(GM);

        summ=sum(sum(filt1));        

        filt1=(filt1/summ);
        Z(i-padsize,j-padsize)=(sum(sum(filt1))/(siz*siz));
    end
end
toc(t0) % -------------------------------------------------------------

对于 330X200 图像，我的顺序代码运行时间为 2.5 秒，但新并行代码的运行时间为 15 秒。我不知道为什么？？？？我需要一些建议来改进它。我是 CUDA 编程新手。

Answer 1

> help gather
...
X = GATHER(A) when A is a GPUArray, X is an array in the local workspace
with the data transferred from the GPU device.
....

filt1 = gather(GM) 每一步都是将 GM 从 GPU 复制到 CPU，效率很低。您应该将整个计算移动到循环嵌套内，或者最好将整个循环嵌套移动到 GPU 内核。否则你会忘记任何加速。