6

我绝对是 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];
    }
}
4

3 回答 3

2

我认为您需要用于 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

于 2012-10-05T07:32:26.420 回答
0

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

因此，减少的典型解决方案如下：

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

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

  2. 分配大小等于内核组数 (=global_size/local_size) 的浮点数组并设置参数 part_sum。

  3. 使用内核“本地大小”x size_of(float) 和 NULL 设置参数 tmp_reduce：

    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 小得多。

这不是完全“并行缩减”，因为您可以使用更复杂的算法使用 Log2(local_size) 操作并行求和 tmp_reduce 数组，但这必须比原子操作快得多。

此外，请查看http://developer.amd.com/Resources/documentation/articles/pages/OpenCL-Optimization-Case-Study-Simple-Reductions_2.aspx了解更好的并行缩减方法。

于 2012-10-06T09:18:56.853 回答
0

原子添加是一种解决方案，但您可能会遇到性能问题，因为原子部分会序列化您的工作项。

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

fi[get_global_id(0)] +=fi1-fi0;

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

于 2012-10-05T07:40:24.900 回答