parallel-processing - openCL减少，并传递二维数组

Question

这是我要转换为 openCL 的循环。

for(n=0; n < LargeNumber; ++n) {    
    for (n2=0; n2< SmallNumber; ++n2) {
        A[n]+=B[n2][n];
    }                                                         
    Re+=A[n];       
}

这就是我到目前为止所拥有的，虽然，我知道它是不正确的并且遗漏了一些东西。

__kernel void openCL_Kernel( __global  int *A,
                         __global  int **B,  
                         __global  int *C, 
                         __global _int64 Re,
                                   int D) 
{

int i=get_global_id(0);
int ii=get_global_id(1);

A[i]+=B[ii][i];

//barrier(..); ?

Re+=A[i];

}

我是这类事情的完全初学者。首先，我知道我不能将全局双指针传递给 openCL 内核。如果可以，请在发布解决方案之前等待几天左右，我想自己解决这个问题，但如果您能帮助我指出正确的方向，我将不胜感激。

Answer 1

关于传递双指针的问题:这种问题通常通过将整个矩阵(或您正在处理的任何内容)复制到一个连续的内存块中来解决，如果这些 block 具有不同的长度，则传递另一个数组，其中包含偏移量各个行(因此您的访问权限看起来像 B[index[ii]+i] )。

现在将您减少到 Re :因为您没有提到您正在使用哪种设备，所以我将假设它的 GPU。在那种情况下，我会避免在同一个内核中进行缩减，因为它会像你发布它的方式一样慢(你必须在数千个线程上序列化对 Re 的访问(以及对 A[i] 的访问)。
相反，我会写出想要的内核，它将所有 B[*][i] 加总为 A[i] 并将从 A 减少到另一个内核中的 Re 并分几个步骤进行，也就是说，您使用一个减少内核，它在 n 元素上运行并将它们减少到类似于 n / 16 的东西(或任何其他数字)。然后你迭代地调用那个内核，直到你找到一个元素，这就是你的结果(我故意让这个描述含糊不清，因为你说你想自己想办法)。

作为旁注:您意识到原始代码并不完全具有良好的内存访问模式？假设 B 相对较大(并且由于第二维而比 A 大得多)，内部循环遍历外部索引会产生大量缓存未命中。这在移植到对连贯内存访问非常敏感的 gpu 时更糟

因此，像这样重新排序可能会大大提高性能:

for (n2=0; n2< SmallNumber; ++n2)
  for(n=0; n < LargeNumber; ++n)    
    A[n]+=B[n2][n];
for(n=0; n < LargeNumber; ++n)                                                 
  Re+=A[n];       

如果您有一个擅长自动向量化的编译器，则尤其如此，因为它可能能够对该构造进行向量化，但对于原始代码来说，它不太可能这样做(并且如果它不能证明 A 和 B[n2] 不能引用相同的内存它不能把原始代码变成这个)。