opencl 本地和全局内存中的多个读写同步问题

Question

我有一个 opencl 内核，可以在字符串中找到最大的 ASCII 字符。 问题是我无法将多个读写同步到全局和本地内存。 我正在尝试通过将其与 local_maximum 进行比较来更新共享内存中的 local_maximum 字符，并在工作组（最后一个线程）的末尾更新 global_maximum 字符。我猜，线程是一个接一个地写。

例如：输入字符串：“加勒比海盗”。

输出字符串：'r'（但应该是's'）。

请查看代码并给出一个解决方案，说明我可以做些什么来同步所有内容。我相信有扎实知识的人可以理解代码。欢迎提供优化提示。

代码如下：

__kernel void find_highest_ascii( __global const char* data, __global char* result, unsigned int size,  __local char* localMaxC )
{
//creating variables and initialising..
unsigned int i, localSize, globalSize, j;
char privateMaxC,temp,temp1;

i = get_global_id(0);
localSize = get_local_size(0);
globalSize = get_global_size(0);

privateMaxC = '\0';

if(i<size){
if(i == 0)
read_mem_fence( CLK_LOCAL_MEM_FENCE );
*localMaxC = '\0';
mem_fence( CLK_LOCAL_MEM_FENCE);

////////////////////////////////////////////////////
/////UPDATING PRIVATE MAX CHARACTER/////////////////
////////////////////////////////////////////////////

for( j = i; j<size; j+=globalSize )
{
    if( data[j] > privateMaxC )
    {
        privateMaxC = data[j];
    }
}

///////////////////////////////////////////////////


///////////////////////////////////////////////////
////UPDATING SHARED MAX CHARACTER//////////////////
///////////////////////////////////////////////////

temp = *localMaxC;
read_mem_fence( CLK_LOCAL_MEM_FENCE );

if(privateMaxC>temp)
{
    *localMaxC = privateMaxC;
    write_mem_fence( CLK_LOCAL_MEM_FENCE );
    temp = privateMaxC;
}

//////////////////////////////////////////////////


//UPDATING GLOBAL MAX CHARACTER.

temp1 = *result;

if(( (i+1)%localSize == 0 || i==size-1) && (temp > temp1 ))
{
            read_mem_fence( CLK_GLOBAL_MEM_FENCE );
    *result = temp;
    write_mem_fence( CLK_GLOBAL_MEM_FENCE );
}


 }
}

Answer 1

线程将覆盖彼此的值是正确的，因为您的代码充满了race conditions。在 OpenCL 中，无法在不同工作组中的工作项之间进行同步。您可以通过使用内置的原子函数来使您的代码更加更简单，而不是尝试使用显式栅栏来实现这种同步。特别是，有一个内置的atomic_max 可以完美地解决您的问题。

因此，您当前必须更新本地和全局内存最大值，而不是代码，只需执行以下操作：

kernel void ascii_max(global int *input, global int *output, int size,
                      local int *localMax)
{
  int i = get_global_id(0);
  int l = get_local_id(0);

  // Private reduction                                                          
  int privateMax = '\0';
  for (int idx = i; idx < size; idx+=get_global_size(0))
  {
    privateMax = max(privateMax, input[idx]);
  }

  // Local reduction                                                            
  atomic_max(localMax, privateMax);
  barrier(CLK_LOCAL_MEM_FENCE);

  // Global reduction                                                           
  if (l == 0)
  {
    atomic_max(output, *localMax);
  }
}

这将要求您更新本地内存暂存空间和最终结果以使用 32 位整数值，但总体而言，这是解决此问题的一种更简洁的方法（更不用说它确实有效）。

---

非原子解决方案

如果您真的不想使用原子，那么您可以使用本地内存和工作组屏障来实现沼泽标准减少。这是一个例子：

kernel void ascii_max(global int *input, global int *output, int size,
                      local int *localMax)
{
  int i = get_global_id(0);
  int l = get_local_id(0);

  // Private reduction                                                          
  int privateMax = '\0';
  for (int idx = i; idx < size; idx+=get_global_size(0))
  {
    privateMax = max(privateMax, input[idx]);
  }

  // Local reduction                                                            
  localMax[l] = privateMax;
  for (int offset = get_local_size(0)/2; offset > 1; offset>>=1)
  {
    barrier(CLK_LOCAL_MEM_FENCE);
    if (l < offset)
    {
      localMax[l] = max(localMax[l], localMax[l+offset]);
    }
  }

  // Store work-group result in global memory                                   
  if (l == 0)
  {
    output[get_group_id(0)] = max(localMax[0], localMax[1]);
  }
}

这使用本地内存作为暂存空间一次比较元素对。每个工作组将产生一个结果，该结果存储在全局内存中。如果您的数据集很小，您可以使用单个工作组运行它（即使全局和本地大小相同），这将正常工作。如果它更大，您可以通过运行此内核两次来运行两阶段归约，例如：

size_t N = ...; // something big

size_t local  = 128;
size_t global = local*local; // Must result in at most 'local' number of work-groups

// First pass - run many work-groups using temporary buffer as output
clSetKernelArg(kernel, 1, sizeof(cl_mem), d_temp);
clEnqueueNDRangeKernel(..., &global, &local, ...);

// Second pass - run one work-group with temporary buffer as input
global = local;
clSetKernelArg(kernel, 0, sizeof(cl_mem), d_temp);
clSetKernelArg(kernel, 1, sizeof(cl_mem), d_output);
clEnqueueNDRangeKernel(..., &global, &local, ...);

我会留给您运行它们并决定哪种方法最适合您自己的数据集。