OpenCL - 计算期间的递增求和

Question

我是OpenCL编程中的绝对新手。为我的应用程序。 （分子模拟）我写了一个计算lennard-jones液体分子间电位的核心。在这个内核我需要计算所有粒子的潜在的累积值与一个：

__kernel void Molsim(__global const float* inmatrix, __global float* fi, const int c, const float r1, const float r2, const float r3, const float rc, const float epsilon, const float sigma, const float h1, const float h23) 
{ 
    float fi0; 
    float fi1; 
    float d; 

    unsigned int i = get_global_id(0); //number of particles (typically 2000) 

    if(c!=i) { 
     // potential before particle movement 
     d=sqrt(pow((0.5*h1-fabs(0.5*h1-fabs(inmatrix[c*3]-inmatrix[i*3]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(inmatrix[c*3+1]-inmatrix[i*3+1]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(inmatrix[c*3+2]-inmatrix[i*3+2]))),2.0)); 
     if(d<rc) { 
     fi0=4.0*epsilon*(pow(sigma/d,12.0)-pow(sigma/d,6.0)); 
     } 
     else { 
     fi0=0; 
     } 
     // potential after particle movement 
     d=sqrt(pow((0.5*h1-fabs(0.5*h1-fabs(r1-inmatrix[i*3]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(r2-inmatrix[i*3+1]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(r3-inmatrix[i*3+2]))),2.0)); 
     if(d<rc) { 
     fi1=4.0*epsilon*(pow(sigma/d,12.0)-pow(sigma/d,6.0)); 
     } 
     else { 
      fi1=0; 
     } 
     // cumulative difference of potentials 
     // fi[0]+=fi1-fi0; changed to full size vector 
     fi[get_global_id(0)]=fi1-fi0; 
     } 
}   

我的问题是在该行：FI [0] + = FI1-fi0 ;.在一元向量fi [0]中是错误的结果。 我读了一些关于减少和数的内容，但是我不知道在计算过程中如何去做。

存在我的问题的任何简单的解决方案？

通知： 我试图添加下一个内核为向量分量之和（见下面的代码），但有一个比当我使用CPU和向量更大放缓。

__kernel void Arrsum(__global const float* inmatrix, __global float* outsum, const int inmatrixsize, __local float* resultScratch) 
{ 
     // načtení indexu 
     int gid = get_global_id(0); 
     int wid = get_local_id(0); 
     int wsize = get_local_size(0); 
     int grid = get_group_id(0); 
     int grcount = get_num_groups(0); 

     int i; 
     int workAmount = inmatrixsize/grcount; 
     int startOffest = workAmount * grid + wid; 
     int maxOffest = workAmount * (grid + 1); 
     if(maxOffest > inmatrixsize){ 
     maxOffest = inmatrixsize; 
    } 

    resultScratch[wid] = 0.0; 
    for(i=startOffest;i<maxOffest;i+=wsize){ 
      resultScratch[wid] += inmatrix[i]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 

    if(gid == 0){ 
      for(i=1;i<wsize;i++){ 
        resultScratch[0] += resultScratch[i]; 
      } 
      outsum[grid] = resultScratch[0]; 
    } 
} 

Answer 1

我认为你需要fi [0]的atomic_add原子函数+ = fi1-fi0;

警告：使用原子函数会降低性能。

这里用增量原子函数的两个例子。

实施例而不原子功能和2的工作项：

__kernel void inc(global int * num){ 
    num[0]++; //num[0] = 0 
} 

1. 工作项1读取NUM [0]：0
2. 工作项2读取NUM [0]：0
3. 工作工作项目1增加num [0]：0 + 1
4. 工作项目2增量num [0]：0 + 1
5. 工作项目1写入num [0] ：NUM [0] = 1
6. 工作项2写入NUM [0]：NUM [0] = 1

结果：NUM [0] = 1个

实施例与原子功能和2工作项：

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable 

__kernel void inc(global int * num){ 
    atom_inc(&num[0]); 
} 

1. 工作项1读取NUM [0]：0
2. 工作项1倍的增量NUM [0]：0 + 1
3. 工作项1写入NUM [0]：NUM [0] = 1
4. 工作项2读取NUM [0]：1
5. 工作项2倍的增量NUM [0]：1个+ 1
6. 工作第2项写入NUM [0]：NUM [0] = 2

结果：NUM [0] = 2

Answer 2

原子补充的是一个解决方案，但你可以得到的性能问题，因为原子部分将连载你的工作项目。

我认为一个更好的解决方案是，对于每个工作项，在自己的变量来写，如：

音响[get_global_id（0）] + = FI1-fi0;

然后，您可以将数组传送到CPU并对所有元素进行求和，也可以在GPU上使用算法并行执行。

Answer 3

所有的线程都由“组”执行。您可以使用get_local_id（dim）函数确定组中的线程ID。每个组内的线程都可以使用共享内存（在OpenCL中称为“本地内存”）并同步它们的执行，但不同组中的线程不能直接通信。

所以，对于减少tipical溶液以下：

1. 添加临时数组part_sum（全局）和tmp_reduce（本地）到内核ARGS：

   __kernel void Molsim(..., __global float *part_sum, __local float *tmp_reduce) 
   

2. 浮点数分配的阵列大小等于内核的组数（= global_size/local_size），并设置参数part_sum。

3. 设置参数tmp_reduce你的内核 “局部尺寸” X size_of（浮动）和NULL：

   clSetKernelArg(kernel,<par number>,sizeof(float)*<local_size>,NULL); 
   

4. 在内核中，用下面的代码替换代码：

   int loc_id=get_local_id(0); 
   
   ... 
   
   //  fi[0]+=fi1-fi0; 
        tmp_reduce[loc_id]=fi1-fi0; 
        } 
       barrier(CLK_LOCAL_MEM_FENCE); 
       if(loc_id==0) { 
       int i; 
       float s=tmp_reduce[0]; 
       for(i=1;i<get_local_size(0);i++) 
        s+=tmp_reduce[i]; 
       part_sum[get_group_id(0)]=s; 
       } 
   } 
   

5. 后内核执行完成后，只需在主机上总结part_sum [array]的内容，该内容比global_size小得多。

这不是完全“平行减少”，因为你可以总结使用更复杂的算法并行使用LOG 2（local_size）操作tmp_reduce阵列，但是这必须比原子操作快得多。

此外，看看http://developer.amd.com/Resources/documentation/articles/pages/OpenCL-Optimization-Case-Study-Simple-Reductions_2.aspx的并行还原方法。