7

这是我要转换为 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 内核。如果可以的话，请在发布解决方案之前等待几天左右，我想自己解决这个问题，但如果你能帮助我指出正确的方向，我将不胜感激。

4

1 回答 1

12

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

现在将您简化为Re：由于您没有提到您正在使用哪种设备，我将假设它的 GPU。在那种情况下，我会避免在同一个内核中进行缩减，因为它会像你发布它的方式一样慢（你必须序列化对Re数千个线程的访问（以及对的访问A[i]）。相反，我会写出想要的内核，它将所有内容汇总并放入B[*][i]另一个内核中，并分几个步骤完成，也就是说，您使用一个对元素进行操作并将它们简化为类似的归约内核A[i]ARenn / 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]可以'不引用相同的内存它不能把原始代码变成这个）。

于 2012-01-07T17:27:43.720 回答