OpenCL - 计算过程中的增量求和

Question

我在 OpenCL 编程方面绝对是新手。对于我的应用程序。 (分子模拟)我编写了一个内核来计算伦纳德琼斯液体的分子间势。在这个内核中，我需要计算所有粒子势能的累积值:

__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原子函数；

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

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

没有原子函数和 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