Cuda 累积图像的线条

Question

我必须有效地累积存储在数组中的图像的行。

我想出了一个真正天真的解决方案，但我很确定有更好的方法可以用 cuda 来实现。

__global__
void Accumulate(double *x, double *y, int height, int width)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;

    if (i >= width)
        return;
    for (int j = 0; j < height; j++)
    {
        y[i] += x[j*width+ i];
    }

}

这个问题的优雅解决方案是什么？

Answer 1

一般提示

__global__
void Accumulate(float *x, float *y, int height, int width)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i >= width) return;
    float sum = 0;
    for (int j = 0; j < height; j++)
    {
        sum += x[i*width + j];
    }
    y[i] = sum;
}

这是对应用一些改进后代码的外观的建议。

已将 double 替换为 float。这在大多数显卡上要快得多。如果您真的需要双精度，这是不可能的。但通常情况并非如此。
在将总和写入位于全局内存中的数组 y 之前，将总和保存在临时变量 sum 中。访问全局内存是非常昂贵的，尽可能避免它们。

这些提示可能会帮助您在 GPU 上加速代码。

寻找最佳块大小

同样对性能有影响的是块大小。这里没有最佳推荐。但是，一个块中的线程总数应始终能被 32 整除。这是一个 warp 的大小。

Answer 2

对于普通的 CPU 优化解决方案，我会使用这个：

void Accumulate(double *x, double *y, int height, int width)
{
    constexpr int line_chunk_size = 16;
    double sum[line_chunk_size];    // local stack memory for summing lines
    double* line_0_begin = x, * line_0_end = x + width;
    // for K chunks (until the less-than-chunk_size part of line is left)
    while (line_0_begin+line_chunk_size <= line_0_end) {
        double* line_b = line_0_begin;
        // copy first line[chunk] to sum[chunk] (to init local sum)
        for (int i = 0; i < line_chunk_size; ++i) sum[i] = line_b[i];
        // sum remaining lines to sum[chunk]
        for (int j = 1; j < height; ++j) {
            line_b += width;
            for (int i = 0; i < line_chunk_size; ++i) sum[i] += line_b[i];
        }
        // add sum[chunk] to global y[chunk]
        // or assign it by "=", if y is not initialized
        for (int i = 0; i < line_chunk_size; ++i) y[i] += sum[i];
        y += line_chunk_size;
        // process next chunk of lines
        line_0_begin += line_chunk_size;
    }

    // process remaining less-than-chunk values of line
    // (this part can be removed, if you have fixed widths divisible by chunk-size)
    int remaining_chunk = line_0_end - line_0_begin;
    if (remaining_chunk <= 0) return;
    // copy first line[remaining] to sum[remaining] (to init local sum)
    for (int i = 0; i < remaining_chunk; ++i) sum[i] = line_0_begin[i];
    // sum remaining lines to sum[remaining]
    for (int j = 1; j < height; ++j) {
        line_0_begin += width;
        for (int i = 0; i < remaining_chunk; ++i) sum[i] += line_0_begin[i];
    }
    // add sum[remaining] to global y[remaining]
    // or assign it by "=", if y is not initialized
    for (int i = 0; i < remaining_chunk; ++i) y[i] += sum[i];
}

gcc6.3 with -O3 -std=c++14 -march=skylake-avx512 命令行选项将生成 this code（如果 link 不再起作用，只需将上面的例程复制到 godbolt.org并自行设置编译器 + 选项）。这在我看来非常理想，我也尝试了其他块大小，例如 32 和 64，但它更像常规循环，而 16 耗尽了 xmm 寄存器以通过完全展开来矢量化块。

（当然，将 -march 调整为您的实际 CPU，这样可执行文件就可以工作了。我只是想试试高端矢量化是否按预期工作）。

这是我用来测试上面代码有效性的代码：

int main()
{
    constexpr int width = 40, height = 20;
    double testX[width*height], testY[width] = {0};
    double* tXptr = testX;
    for (int y = 0; y < height; ++y) {
        for (int x = 0; x < width; ++x) {
            *tXptr = (y <= x) ? 1+x : 0;
            ++tXptr;
        }
    }
    // print testX for verification
    for (int y = 0; y < height; ++y) {
        for (int x = 0; x < width; ++x) std::cout << testX[y*width + x] << " ";
        std::cout << "\n";
    }
    std::cout << "Y before:\n";
    for (int x = 0; x < width; ++x) std::cout << testY[x] << " ";
    std::cout << "\n";
    Accumulate(testX, testY, height, width);
    std::cout << "Y after first Accumulate:\n";
    for (int x = 0; x < width; ++x) std::cout << testY[x] << " ";
    std::cout << "\n";
    Accumulate(testX, testY, height, width);
    std::cout << "Y after second Accumulate:\n";
    for (int x = 0; x < width; ++x) std::cout << testY[x] << " ";
    std::cout << "\n";
}

拜托，我将不胜感激任何 profiling/measures 在现实生活中，它与您最初的 CUDA 实现的对比（这是单个 CPU 核心，但可以按块传播到多个线程很容易，虽然我不确定 CPU 会在多大程度上识别出内存访问模式没有冲突，也许我会在每个核心上使用几个块，所以 core0 会先处理 10 个列块，core1 块 11。 .20 等。因此每个线程将彼此分开）。我很想听听 CPU（根据我的说法 "good code" ;)）在你的原始版本中对抗 CUDA 有多大，只是想知道这些 GPU 有多强大。谢谢（如果做一些分析不需要太多时间的话）。

编辑：我所做的只是同时处理来自同一行的多个值 ("chunk") 到局部总和变量中，最后我将其添加到 y。我避免在寻址 x/y 内存时使用任何 multiplication/division/modulus，总体上让编译器有机会了解正在发生的事情，因此它可以积极优化。

如果"chunk"大小是4，你的数据是10x4，这是x访问内存的顺序：

Cuda 累积图像的线条

Cuda accumulate lines of an image

c++

cuda

image

reduction