奇偶排序：在 CUDA 中使用多个块时结果不正确

Posted 2023-03-23

【中文标题】奇偶排序：在 CUDA 中使用多个块时结果不正确

【英文标题】：Odd-even sort: Incorrect results when using multiple blocks in CUDA

【发布时间】：2021-10-08 02:01:17

【问题描述】：

我是 PyCUDA 的新手，正在尝试使用 PyCUDA 实现奇偶排序。

我设法在大小限制为 2048 的数组上成功运行它（使用一个线程块），但是当我尝试使用多个线程块时，结果不再正确。我怀疑这可能是同步问题，但不知道如何解决。

bricksort_src = """
__global__
void bricksort(int *in, int *out, int n)
  int tid = threadIdx.x + (blockIdx.x * blockDim.x);

  if((tid * 2) < n) out[tid * 2] = in[tid *2];
  if((tid * 2 + 1) < n) out[tid * 2 + 1] = in[tid * 2 + 1];
  __syncthreads();


  // odd and even are used for adjusting the index
  // to avoid out-of-index exception
  int odd, even, alter;
  odd = ((n + 2) % 2) != 0;
  even = ((n + 2) % 2) == 0;
  // alter is used for alternating between the odd and even phases
  alter = 0;

  for(int i = 0; i < n; i++)  

    int idx = tid * 2 + alter;
    int adjust = alter == 0 ? odd : even;

    if(idx < (n - adjust))
      int f, s;
      f = out[idx]; 
      s = out[idx + 1];
      if (f > s)
        out[idx] = s; 
        out[idx + 1] = f;
      
    
    __syncthreads();
    alter = 1 - alter;
  

"""
bricksort_ker = SourceModule(source=bricksort_src)
bricksort = bricksort_ker.get_function("bricksort")


np.random.seed(0)
arr = np.random.randint(0,10,2**11).astype('int32')

iar = gpuarray.to_gpu(arr)
oar = gpuarray.empty_like(iar)
n = iar.size
num_threads = np.ceil(n/2)

if (num_threads < 1024):
  blocksize = int(num_threads)
  gridsize = 1
else: 
  blocksize = 1024
  gridsize = int(np.ceil(num_threads / blocksize))

bricksort(iar, oar, np.int32(n),
          block=(blocksize,1,1),
          grid=(gridsize,1,1))

【问题讨论】：

奇偶排序不能轻易/轻易地扩展到单个线程块之外。如果您查看 CUDA 示例代码，您将看到它们以块级别的奇偶排序开始，然后使用合并来组合块级别的结果。如果您对排序性能很认真，您可能应该使用库实现，而不是尝试自己编写。将奇偶排序扩展到线程块之外的直接途径是使用设备范围的同步，但这很昂贵，而且从性能角度来看可能不会很有趣。

@RobertCrovella 感谢您的评论。我写这个只是为了自我练习，不知道这是不可行的。我确实考虑过跨块同步，但你是对的；在现实中这样做可能太昂贵了。

如果你想向自己证明它是可行的，那么编写你的内核使其只进行一次交换。 （内核中没有 for 循环）然后在循环中调用该内核，以进行所有奇偶交换。内核调用本身将充当设备范围的同步。

@RobertCrovella 我尝试了您推荐的方法，它适用于多个块。但是，正如您所说，它似乎效率低下，我猜这可能是由于启动多个内核的开销。再次感谢您的帮助。

【参考方案1】：

将 cmets 组装成答案：

奇偶排序不能轻易/轻易地扩展到单个线程块之外（因为它需要同步） CUDA __syncthreads() 仅在块级别同步。如果没有同步，CUDA 不会指定线程执行的特定顺序。

对于严肃的排序工作，我推荐一个库实现，例如cub。如果您想从 python 执行此操作，我推荐cupy。

CUDA 有 a sample code 演示了块级别的奇偶排序，但由于同步问题，它选择了一种合并方法来组合结果

应该可以编写一个只进行一次交换的奇偶排序内核，然后在循环中调用这个内核。内核调用本身充当设备范围的同步点。

或者，应该可以使用cooperative groups grid sync 在单个内核启动中完成这项工作。

这些方法都不会比一个好的库实现更快（它不依赖于奇偶排序开始）。