c-如何优化这个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的代码运行速度要快得多。

这里有一个GCC示例(适用于删除CUDA注释(，您可以看到，即使是最低级别的优化(-O1(也会删除整个循环，只留下比较和i的增量。有了-O2，整个循环就被去除了。