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]+=fi1-fi0; 的atomic_add原子函数；

\n\n

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

\n\n

这里有两个带有增量原子函数的示例。

\n\n

没有原子函数和 2 个工作项的示例：

\n\n

__kernel\xc2\xa0void\xc2\xa0inc(global\xc2\xa0int\xc2\xa0* num){\n    num[0]++; //num[0] = 0\n}\n

\n\n

\n
1. 工作项 1 读取 num[0]: 0
\n
2. 工作项 2 读取 num[0]: 0
\n
3. 工作项 1 增量 num[0]: 0 + 1
\n
4. 工作项 2 增量 num[0]: 0 + 1
\n
5. 工作项 1 写入 num[0]：num[0] = 1
\n
6. 工作项 2 写入\xc2\xa0num[0]：num[0] = 1
\n

\n\n

结果：数字[0] = 1

\n\n

具有原子函数和 2 个工作项的示例：

\n\n

#pragma\xc2\xa0OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n\n__kernel\xc2\xa0void\xc2\xa0inc(global\xc2\xa0int\xc2\xa0* num){\n    atom_inc(&num[0]);\n}\n

\n\n

\n
1. 工作项 1 读取 num[0]: 0
\n
2. 工作项 1 增量 num[0]: 0 + 1
\n
3. 工作项 1 写入 num[0]：num[0] = 1
\n
4. 工作项 2 读取 num[0]: 1
\n
5. 工作项 2 增量 num[0]: 1 + 1
\n
6. 工作项 2 写入\xc2\xa0num[0]：num[0] = 2
\n

\n\n

结果：数字[0] = 2

\n