2013-03-04 71 views
6

我是OpenCL的新手。但是,我了解C/C++基础知识和OOP。 我的问题如下:是否可以并行运行总计算任务?这在理论上是可能的吗?下面我将介绍我所试图做的:是否可以在OpenCL中并行运行求和计算?

的任务,例如:

double* values = new double[1000]; //let's pretend it has some random values inside 
double sum = 0.0; 

for(int i = 0; i < 1000; i++) { 
    sum += values[i]; 
} 

我试过的OpenCL内核做(我觉得这是错误的,因为或许它访问相同的“总和”变量从不同的线程/任务同时):

__kernel void calculate2dim(__global float* vectors1dim, 
          __global float output, 
          const unsigned int count) { 
    int i = get_global_id(0); 
    output += vectors1dim[i]; 
} 

此代码是错误的。如果有人回答我,如果理论上可以并行运行这些任务,并且如果是这样的话,我将非常感激。

+5

这是一个典型的下降问题。看看[这里](http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf)一步一步解释优化这个过程,核心架构(它是CUDA,但是原理完全一样,除了关于模板的部分可能)。虽然有关该主题的更多介绍性材料可能会更有帮助,但我会将其留给适当的答案。 – 2013-03-04 11:43:55

+0

非常感谢!现在我知道这是一个普遍的问题,并且会很快解决它! – Vladimir 2013-03-04 12:28:20

回答

0

如果您想以并行方式对数组的值进行求和,则应确保减少争用并确保跨线程无数据依赖关系。

数据依赖性将导致线程必须等待对方,造成争用,这是你想避免得到真正的并行化。

你可以做到这一点的一种方法是将你的数组拆分成N个数组,每个数组包含你原始数组的一些子部分,然后用每个不同的数组调用你的OpenCL内核函数。

最后,当所有内核都完成了艰苦的工作后,您可以将每个数组的结果总结为一个。该操作可以通过CPU轻松完成。

关键是在每个内核中完成的计算之间没有任何依赖关系,因此您必须拆分数据并进行相应的处理。

我不知道你的数据是否与你的问题有任何实际的依赖关系,但这是你要弄清楚的。

+0

谢谢你的回答。可能,我应该将我的阵列分成几个独立的阵列!你知道任何方式将二维数组(如double [] [])传递给内核吗?因为pointet-to-pointer不能用作函数参数。 – Vladimir 2013-03-04 12:21:32

+1

你不需要传递二维数组,只需传递你的缓冲区并像这样myarr [y * WIDTH + x]就可以了。 – alariq 2013-03-06 15:50:40

0

我提供的供参考的一段代码应该可以完成这项工作。

E.g.您有N元素,并且您的工作组的大小为WS = 64。我认为N的倍数2 * WS(这很重要,一个工作组计算2 * WS元素的总和)。然后,你需要运行内核指定:

globalSizeX = 2*WS*(N/(2*WS)); 

结果总和阵列将拥有2 * WS元素的部分和。 (例如,sum [1] - 将包含索引为2 * WS4 * WS-1)的元素之和。

如果你的globalSizeX是2 * WS或更少(这意味着你只有一个工作组),那么你就完成了。只需使用sum [0]作为结果。 如果不是 - 您需要重复此过程,此时使用总和数组作为输入数组并输出到其他数组(在它们之间创建2个数组和乒乓)。等等,直到你将只有一个工作组。

还搜索Hilli Steele/Blelloch并行算法。 This文章可能还有

下面是实际的例子有用:

__kernel void par_sum(__global unsigned int* input, __global unsigned int* sum) 
{ 
    int li = get_local_id(0); 
    int groupId = get_group_id(0); 

    __local int our_h[2 * get_group_size(0)]; 
    our_h[2*li + 0] = hist[2*get_group_size(0)*blockId + 2*li + 0]; 
    our_h[2*li + 1] = hist[2*get_group_size(0)*blockId + 2*li + 1]; 

    // sweep up 
    int width = 2; 
    int num_el = 2*get_group_size(0)/width; 
    int wby2 = width>>1; 

    for(int i = 2*BLK_SIZ>>1; i>0; i>>=1) 
    { 

     barrier(CLK_LOCL_MEM_FENCE); 

     if(li < num_el) 
     { 
      int idx = width*(li+1) - 1; 
      our_h[idx] = our_h[idx] + our_h[(idx - wby2)]; 
     } 

     width<<=1; 
     wby2 = width>>1; 
     num_el>>=1; 
    } 

     barrier(CLK_LOCL_MEM_FENCE); 

    // down-sweep 
    if(0 == li) 
     sum[groupId] = our_h[2*get_group_size(0)-1]; // save sum 
} 
相关问题