openCL减少，并传递二维数组

Posted 2023-05-08

【中文标题】openCL减少，并传递二维数组

【英文标题】：openCL reduction, and passing 2d array

【发布时间】：2012-02-04 11:12:28

【问题描述】：

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

【问题讨论】：

“我不能将全局双指针传递给 openCL 内核” 你的措辞让我很困惑。您可以传递一个双指针（例如“__global double *A”）。您不能传递 2D 指针（例如“__global int **B”）。

您是否考虑过将程序分成两个独立的内核（按顺序执行），一个用于内循环，一个用于外循环？

【参考方案1】：

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

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

【讨论】：

谢谢！这给了我很多思考。