for 循环中断会减慢 OpenCL 内核

Question

学习OpenCL实现程序，为每个像素在10个像素半径内找到8个最相似的像素：

//-DBEST=8

__kernel void best_pixel_matches(
    __global const uint * image,
    const ushort width,
    const ushort height,
    // [ width ] [ height ] [ BEST ] [ 2 ]
    __global short * best,
    __global int * errors
)
{
    const int x = get_global_id(0);
    const int y = get_global_id(1);

    const uint pixel = image[ y * width + x ];

    const short SURROUND = 10;

    //local best pixels coordinates:
    short s[ BEST ][ 2 ];
    //... corresponding local best pixels computed deltas (storage instead of recompute speeds program about twice): 
    uint d[ BEST ];
    //... init with none:
    for ( short i = 0; i < BEST; ++ i ) d[ i ] = -1;

    for ( short sw = - SURROUND; sw < SURROUND; ++ sw ) {
    for ( short sh = - SURROUND; sh < SURROUND; ++ sh ) {
        //avoid inexistent pixels:
        const short tw = x + sw; if ( tw < 0 || tw >= width  ) continue;
        const short th = y + sh; if ( th < 0 || th >= height ) continue;
        //not with itself:
        if ( tw == x && th == y ) continue;
        const uint diff = (uint) abs( (long)pixel - (long)image[ th * width + tw ] );

        for ( short si = 0; si < BEST; ++ si ) {
            if ( d[si] == -1 || diff < d[si] ) {
                d[si] = diff;
                s[si][0] = tw;
                s[si][1] = th;
                break;//<-- this line causes 4 (!) times slow down
            }
        }
    } }

    //copy results from private memory to global:
    const long p = ( x * height + y ) * BEST;
    for ( short b = 0; b < BEST; ++ b ) {
        const long pb = ( p + b ) * 2;
        best[ pb     ] = s[b][0];
        best[ pb + 1 ] = s[b][1];
    }
}

问题是 GPU 3593ms 处理 2560*1440 图像，这几乎与我的普通 C++ CPU 代码完全相同之前（~8500ms）。然后我试着到处调整它，偶尔删除 break; 行和执行时间 -> 900ms!

如此惊人的加速背后的原因是什么？ break; 行只是告诉程序它不需要进一步检查任何其他像素所以应该减少执行时间 - 而不是减慢速度？也许还有其他一些方法可以加快这个程序的速度？ :)

Answer 1

您没有说明您运行使用哪种设备，但是：

在 CPU 上，优化包括自动向量化，这通常对分支非常敏感。所以因子 4 并不奇怪；您的循环迭代计数在构建时是已知的，因此如果没有提前结束，它们可以完全展开和矢量化。对于早期输出，代码可能需要序列化，因此您失去了矢量化带来的 4 倍、8 倍或 16 倍的加速（取决于 CPU 类型），并且只能从早期输出中恢复一部分。
GPU 工作项通常运行在锁步线程中，但如果发生在此处，因子 4 的变化似乎不太可能；通常情况下 break; 不会产生任何改进，除非整个波前很有可能走捷径。然而，某些 GPU 模型更喜欢 SIMD 代码，因此您可能运行遇到与 CPU.

请注意，删除 break 可能会用相同的值填充 d 和 s 中的所有 BEST 项目（如果我没有正确阅读您的代码），所以它不会' t 产生与 break;.

相同的输出

Answer 2

尝试维护内存访问顺序，就在全局写入之前

 // here
 barrier(CLK_GLOBAL_MEM_FENCE);

 //copy results from private memory to global:
 const long p = ( x * height + y ) * BEST;

如果您的性能问题源于内存访问模式，这应该可以缓解它。

如果问题仅来自线程之间的分支，那么您也可以尝试在最内层循环之后的第二个循环中保持顺序：

for ( short si = 0; si < BEST; ++ si ) {
            if ( d[si] == -1 || diff < d[si] ) {
                d[si] = diff;
                s[si][0] = tw;
                s[si][1] = th;
                break;
            }
        }
barrier(CLK_LOCAL_MEM_FENCE); // local => should be faster for many devices

以便在每次迭代中，所有相邻线程重新组合在一起并继续执行相同的指令，直到新的中断打破顺序。

最后，算法似乎通过递减值来对 SURROUND 像素之间的像素值差异进行排序。因此，最内层的循环实际上可以移动到最外层并使所有分支最小化，就在最内层的两个新循环（之前是最外层的）之后。但这会增加 image[ th * width + tw ] 读取的最佳时间，因此它也可能更慢（可能不会在此之前将数据移动到本地内存）。但是，d 现在不需要是一个数组，因此可以保存 BEST 数量的私有寄存器，并且可以通过降低寄存器压力来提高它的性能。

for 循环中断会减慢 OpenCL 内核

for-loop interruption slows down OpenCL kernel

performance

opencl