cuda 代码的优化技巧答案

【问题标题】：Optimization tips for a cuda codecuda 代码的优化技巧
【发布时间】：2012-08-30 23:45:05
【问题描述】：

我编写了一段代码，用于在 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 编程新手。

【问题讨论】：

双精度在大多数支持 CUDA 的 GPU 上效率非常低，因为 DP 指令会被序列化 - 你可以使用单精度吗？
是的，我也测试了单精度，但结果没有改变！
另一个瓶颈可能是 if 语句。尽量避免 if 并只写 M(i,j)=M(i,j)+X(i,j)>=means 等...在不同的注释上，而不是嵌套的 for 循环，可以'你只需写： function M=return_step(X, mean) M(find(X>=means))=1; ?
我认为您的提示是针对顺序代码的，对吧？
请尝试将if (xval>=mean) 从您的内核移动到主机端。在 CUDA 中，一组线程（通常是 32 个线程）并行工作，如果 32 个线程中的一个不满足if，其他线程就会停止。尝试仅将m[idx]=filt[idx]*xval 放入您的内核中，并从中删除条件。此外，使用 CUDA 分析器可能会有所帮助，因为您将获得有关代码可能出错的提示和图形提示。

标签： c++ matlab cuda gpu

【解决方案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 内核。否则你可以忘记任何加速。

【讨论】：

我删除了 ghater 行，但运行时间只减少了 5 秒！您如何解释 cuda 代码的 10 秒运行时间？
您的问题有多大？如果您在循环中启动 CUDA 内核，您每次都将支付内核启动的成本。如果您的问题规模太小，您将无法从并行性中受益，原因有两个：1）如果所有或大部分线程都在工作，GPU 效率最高，即占用率接近 1。和 2）启动开销将占主导地位。实际上，最好将整个嵌套循环放入内核中。
这些当然只是一般的想法。您可以查看以下论文，以很好地总结高效的 GPU 编程指南。 V. Volkov，JW Demmel，基准 GPU 来调整密集线性代数，X. Cui，Y. Chen，H. Mei，提高矩阵乘法和 gpu 上 fft 的性能，JM Cohen，J. Molemake，快速双精度 cfd 代码使用 cuda，

【解决方案2】：

我在 Sobel 过滤器下的评估显示 CPU 在小图像上的表现优于 GPU。我认为您的图像大小对于比较 CPU-GPU 性能而言是如此之小。计算应该足够大以隐藏内核和通信启动开销。

【讨论】：

我不同意你的看法。因为当我在 GPU 上实现整个循环时，我得到了 30 倍的加速。