如何优化此OpenCL内核？

如何解决如何优化此OpenCL内核？

我正在一个项目上，这个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;
}

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

是否应该相同或至少相似？ 即使利用我拥有CPU和GPU的事实寻找答案，我也在多个运行时（PoCL和opencl-amd）中进行了尝试，并且发生了相同的事情：-/

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

解决方法

循环内的所有操作都没有副作用，您只能从那些__global指针中读取，并计算一些临时值，这些临时值最终将通过最后一个{{ 1}}。换句话说，该循环的唯一目的是计算aa的值。

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

这里是a GCC example（适用于删除CUDA注释），您可以看到即使是最低级别的优化（aa）也会删除整个循环，仅留下比较和{ {1}}。使用aa，the whole loop is removed。