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

Question

我是 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))

Answer 1

将评论组合成答案：

奇偶排序不能easily/readily扩展到单个线程块之外（因为它需要同步）CUDA __syncthreads() 仅在块级同步。没有同步，CUDA 没有指定线程执行的特定顺序。
对于严肃的排序工作，我推荐一个库实现，例如cub. If you want to do this from python I recommend cupy。
CUDA 有 a sample code 在块级别演示奇偶排序，但由于同步问题，它选择合并方法来合并结果
应该可以写一个只做一次交换的奇偶排序内核，然后循环调用这个内核。内核调用本身充当设备范围的同步点。
或者，应该可以使用 cooperative groups grid sync.
在单个内核启动中完成工作
none 这些方法可能比一个好的库实现更快（它不依赖于奇偶排序开始）。

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

Odd-even sort: Incorrect results when using multiple blocks in CUDA

python

cuda

pycuda