为什么在此openCL内核代码中，localSum是本地的？

Question

不幸的是，我正在用reference的页面是法语。

有此内核：

  __kernel void sumGPU ( __global const double *input, 
                         __global double *partialSums,
                         __local double *localSums)
 {
  uint local_id = get_local_id(0);
  uint group_size = get_local_size(0);

  // Copy from global to local memory
  localSums[local_id] = input[get_global_id(0)];

  // Loop for computing localSums : divide WorkGroup into 2 parts
  for (uint stride = group_size/2; stride>0; stride /=2)
     {
      // Waiting for each 2x2 addition into given workgroup
      barrier(CLK_LOCAL_MEM_FENCE);

      // Add elements 2 by 2 between local_id and local_id + stride
      if (local_id < stride)
        localSums[local_id] += localSums[local_id + stride];
     }

  // Write result into partialSums[nWorkGroups]
  if (local_id == 0)
    partialSums[get_group_id(0)] = localSums[0];
 } 

这里的想法是执行并行求和（这将花费O（log n）而不是O（n））。我唯一的问题是为什么将localSums声明为__local？可以进一步优化此代码吗？

我只是感到困惑，因为通常我会在内核内部声明局部变量，而在此特定示例中，主机指定大小并将NULL作为clSetKernelArg的参数传递。

该行是以下（主机端）

   // Create the OpenCL kernel
   cl_kernel kernel = clCreateKernel(program, "sumGPU", &ret);

   // Set the arguments of the kernel
   clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&inputBuffer);
   clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&reductionBuffer);
   clSetKernelArg(kernel, 2, local_item_size*sizeof(double),NULL);

Answer 1

我唯一的问题是为什么将localSums声明为local？可以进一步优化此代码吗？

注意下一行如何访问数组中另一个工作项的“插槽”。

   localSums[local_id] += localSums[local_id + stride];

唯一的方法是使用本地或全局内存，并且本地内存无疑比全局更有效。

我只是感到困惑，因为通常我会在内核内部声明局部变量，而在此特定示例中，主机指定大小并传递NULL作为clSetKernelArg的参数。

如果在内核中声明local缓冲区，则其大小必须在编译时固定。通过内核参数动态设置本地缓冲区的大小意味着您可以使用具有不同组大小的相同内核代码。

必须始终从内核内部初始化本地缓冲区，因此将NULL作为主机端的参数指针传递是正确的。