CUDA中归并排序的共享变量传递及线程映射实现问询

我写了一段CUDA代码，先把全局内存的数据加载到寄存器，完成16个数值的排序后写入共享数组：

__global__ void sort16(u32* arr, u32 n) {
__shared__ u32 tmp[16*32]; // 512元素，2KB共享内存，应该能放下？
__shared__ u32 tmp2[16*32]; // 512元素，2KB共享内存，应该能放下？

u32 tid = threadIdx.x;
u32 loc = blockIdx.x * 16;
u32 tmp_loc = tid * 16;
u32 a0 = arr[loc+0], a1 = arr[loc+1], a2 = arr[loc+2], a3 = arr[loc+3],
   a4 = arr[loc+4], a5 = arr[loc+5], a6 = arr[loc+6], a7 = arr[loc+7];
u32 a8 = arr[loc+8], a9 = arr[loc+9], a10 = arr[loc+10], a11 = arr[loc+11],
   a12 = arr[loc+12], a13 = arr[loc+13], a14 = arr[loc+14], a15 = arr[loc+15];
// 此处省略16元素排序逻辑
tmp[tmp_loc] = a0; tmp[tmp_loc+1] = a1; tmp[tmp_loc+2] = a2; tmp[tmp_loc+3] = a3;
tmp[tmp_loc+4] = a4; tmp[tmp_loc+5] = a5; tmp[tmp_loc+6] = a6; tmp[tmp_loc+7] = a7;
tmp[tmp_loc+8] = a8; tmp[tmp_loc+9] = a9; tmp[tmp_loc+10] = a10; tmp[tmp_loc+11] = a11;
tmp[tmp_loc+12] = a12; tmp[tmp_loc+13] = a13; tmp[tmp_loc+14] = a14; tmp[tmp_loc+15] = a15;
__syncthreads();

现在我想基于这两个512元素的共享数组，实现**16元素块→32元素块→64元素块……**的归并，直到共享内存无法容纳，最后把结果写回全局内存。

我知道串行归并的循环写法，但不知道怎么在CUDA里用线程实现，也搞不清merge2核函数的线程和块配置参数：

// 串行归并逻辑示例
__global__ merge2(u32* tmp, u32*tmp2, u32 size) {
  for (u32 a = 0; a < 512; a += 2*size) {
    u32 b = a + size;
    u32 enda = a + size, endb = b + size;
    u32 c = a; // 补充缺失的c初始化
    while (a < enda && b < endb) {
      if (tmp[a] < tmp[b])
        tmp2[c++] = tmp[a++];
      else
        tmp2[c++] = tmp[b++];
    }
    while (a < enda)
      tmp2[c++] = tmp[a++];
    while (b < endb)
      tmp2[c++] = tmp[b++];
  }
}

我预想的调用方式如下，但不确定其中的线程与块参数该怎么填：

u32* t = tmp;
u32* t2 = tmp2;
for (u32 i = 16; i < 512; i *= 2) {
    merge2<<<??, ??>>>(t, t2, i);
    u32* ptrtmp = t;
    t = t2;
    t2 = ptrtmp;
}

核心问题是：随着归并块的大小翻倍，可用的线程数也应该增加，但我不知道该怎么合理利用这些线程来并行归并。

解决方案 #

1. 共享内存归并的线程分工逻辑

每个线程负责归并一组独立的2×size大小的块，而非单个线程串行处理所有块。共享内存总大小为512元素，当归并块大小为size时，总共有512/(2*size)组待归并的块对，让每个线程处理一组即可实现并行化。

2. merge2核函数的修正

修正后的核函数基于线程ID分配任务，同时做边界检查避免越界：

__global__ void merge2(u32* tmp, u32* tmp2, u32 size) {
    // 每个线程负责一组2*size的块
    u32 thread_id = threadIdx.x;
    u32 block_start = thread_id * 2 * size;
    
    // 边界检查：确保不越界共享内存
    if (block_start >= 512) return;
    
    u32 a = block_start;
    u32 b = block_start + size;
    u32 enda = a + size;
    u32 endb = min(b + size, (u32)512); // 处理最后一组可能不足2*size的情况
    u32 c = block_start;
    
    // 归并逻辑
    while (a < enda && b < endb) {
        tmp2[c++] = (tmp[a] < tmp[b]) ? tmp[a++] : tmp[b++];
    }
    while (a < enda) {
        tmp2[c++] = tmp[a++];
    }
    while (b < endb) {
        tmp2[c++] = tmp[b++];
    }
}

3. 线程与块的配置规则

归并阶段的线程数等于待归并的块对数量，即512/(2*size)：

• 当size=16时，块对数量为512/(32)=16，启动16个线程（<<<1, 16>>>）
• 当size=32时，块对数量为512/(64)=8，启动8个线程（<<<1, 8>>>）
• 当size=64时，块对数量为512/(128)=4，启动4个线程（<<<1, 4>>>）
• 以此类推，直到size=256时，块对数量为512/(512)=1，启动1个线程（<<<1, 1>>>）

对应的调用循环写法：

u32* t = tmp;
u32* t2 = tmp2;
for (u32 i = 16; i < 512; i *= 2) {
    u32 num_threads = 512 / (2 * i);
    merge2<<<1, num_threads>>>(t, t2, i);
    __syncthreads(); // 归并完成后必须同步，确保所有线程写完共享内存
    u32* ptrtmp = t;
    t = t2;
    t2 = ptrtmp;
}

4. 关键注意点

• 每次归并后必须调用__syncthreads()，确保所有线程完成当前归并操作，再交换读写数组
• 线程处理连续块的方式能保证共享内存合并访问，提升带宽效率
• 当size=256时，归并后得到512个有序元素，此时可将整个共享数组写回全局内存

内容的提问来源于stack exchange，提问作者Dov