如何优化这个 OpenCL 内核？

Question

我正在做一个项目，但我在使用这个 OpenCL 内核时遇到了一些问题 :-(

__kernel void gemm_fast_5(
    __global double *ar, __global double *br, __global double *cr,
    __global double *pr, __global double *ur,

    unsigned long c, unsigned long c2,
    unsigned long c3, unsigned long c4,
    unsigned long c5, unsigned long m,
    unsigned char com
){
    unsigned long i = get_global_id(0);
    unsigned long j = get_global_id(1);

    unsigned long x = get_local_id(0);
    unsigned long y = get_local_id(1);

    unsigned long cur = i*c3 + j, rl, rl2, rl3;

    #if ks == 1 || ks == 2 || ks == 3 || ks == 4
    unsigned long rl4;
    #endif


    #if ks == 2
    rl = (i << 1)*c;
    #elif ks == 3
    rl = ((i << 1) + 1)*c;
    #else
    rl = i*c;
    #endif

    __local double ut, pt;

    if (x == 0) pt = pr[i*c4 + ks];
    if (y == 0) ut = ur[j*c5 + ks];

    double aa = 0.0;

    double bb, cc;
    double dd, ee;

    for (unsigned long k=0; k<m; k++){
        #if ks == 1 || ks == 4
        rl3 = (k << 1) + 1; rl4 = (k << 2) + 3;

        bb = ar[rl + rl3 - 1]; cc = ar[rl + rl3];
        dd = br[rl2 + rl4 - 1]; ee = br[rl2 + rl4 - 3];
        #elif ks == 2 || ks == 3
        rl3 = (k << 2) + 3; rl4 = (k << 1) + 1;

        bb = ar[rl + rl3 - 3]; cc = ar[rl + rl3 - 2];
        dd = br[rl2 + rl4]; ee = br[rl2 + rl4 - 1];
        #else
        rl3 = (k << 1) + 1;

        bb = ar[rl + rl3 - 1]; cc = ar[rl + rl3];
        dd = br[rl2 + rl3]; ee = br[rl2 + rl3 - 1];
        #endif

        aa += (bb + dd)*(cc + ee);
    }
    cr[cur] = aa - pt - ut;
}

在工作时，我注意到如果我删除最后一行，内核运行时间会减少 6 倍，即使使用 cr[cur] = 5.0 - pt - ut; 更改最后一行也是如此。

不应该采取相同的方式，或者至少类似的东西吗？ 即使在寻找答案，利用我拥有 CPU 和 GPU 的事实，我已经尝试了几个运行时（PoCL 和 opencl-amd）并且发生了同样的事情：-/

如果有人能帮助我理解为什么会发生这种情况，我将不胜感激。我不明白：“v

Answer 1

循环内的所有操作都没有副作用，您只需从那些__global 指针中读取，然后计算一些临时值，这些临时值最终会通过最终的aa += ... 累积到aa 中。换句话说，该循环的唯一目的是计算aa 的值。

因此，如果你从最后一行（循环外）删除aa，循环内的所有操作都完全没用，你最终会得到一个除了读取一些值和更新局部变量之外什么都不做的循环将在函数返回时被丢弃。在启用优化的情况下编译上述代码（我假设您正在这样做，否则您的问题没有多大意义），编译器很可能只是摆脱整个循环。因此，没有最后 aa 的代码运行得更快。

这里是a GCC example（已修改为删除 CUDA 注释），您可以在其中看到即使是最低级别的优化 (-O1) 也会删除整个循环体，只留下比较和 i 的增量。与-O2、the whole loop is removed。