openCL减少,并传递2d数组

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

关于你传递双指针的问题:这种问题通常是通过将整个矩阵(或你正在处理的任何东西)复制到一个连续的内存块中来解决的,如果这些块有不同的长度通过另一个数组,其中包含偏移量各行(所以你的访问看起来像B[index[ii]+i]).

现在你的减少到Re:因为你没有提到你正在研究什么样的设备我会假设它的GPU.在这种情况下,我会避免在同一个内核中进行减少,因为它会像你发布它的方式一样慢(你必须序列化Re对数千个线程的访问(以及对它的访问A[i]).相反,我会写想内核,总结所有B[*][i]进A[i],并把减少从A成Re另一种内核和做几个步骤,这是您使用它运行在一个下降的内核n元素,并将它们降低到类似n / 16(或任何其他号码).然后,你迭代地调用那个内核,直到你归结为一个元素,这是你的结果(我说这个描述故意模糊,因为你说你想要自己想出来).

作为旁注:您意识到原始代码并不完全具有良好的内存访问模式？假设B相对较大(并且A由于第二维而大得多)内部循环遍及外部索引会产生大量的高速缓存.移植到gpu时更是如此,因为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]能够指的是它不能将原始代码转换成相同的内存.