VS2019下CUDA 11.3多线程奇偶排序代码运行正常但无法正确排序求助

核心错误点 #

• 核函数传参错误：核函数仅能直接访问显存空间的指针，当前调用Even、Odd时传入的是主机内存指针a，并非设备端通过cudaMalloc分配的ptr。这种传参方式会触发静默未定义行为，核函数根本无法对目标数组执行交换操作，相当于所有排序逻辑都没有实际生效。
• 主机内存分配错误：a = (int*)malloc(n * Size); 存在参数重复计算的问题，常量Size本身已经是sizeof(int)*n，此处再乘n会造成内存冗余分配，正确写法应为a = (int*)malloc(Size);
• 核函数启动配置冗余：当前启动核函数时传入了n个线程，但偶数阶段仅需要(n+1)/2个线程，奇数阶段仅需要n/2个线程即可覆盖所有索引，多余线程只会空转。虽然n=10时不会触发单块最大线程数限制，但属于无意义的资源浪费。
• 缺少错误检查机制：CUDA核函数是异步启动的，当前代码没有添加启动错误检查，一旦核函数启动失败（比如传参错误、线程数超限），程序不会抛出任何提示，直接往下执行，导致问题很难被定位。

逻辑层面的正确性说明 #

实现的奇偶排序交换逻辑本身没有问题：

• Even核函数中计算的index = threadIdx.x *2，判断index <n-1的边界，刚好覆盖所有偶起始的相邻对（0-1、2-3...）
• Odd核函数中计算的index = threadIdx.x*2 +1，判断index <=n-2的边界，刚好覆盖所有奇起始的相邻对（1-2、3-4...）
• 循环执行n/2轮奇偶交换的迭代次数，也满足奇偶排序收敛的要求（n个元素最多需要n轮全奇偶扫描即可完成排序）

关键修正示例 #

// 修正主机内存分配
const int Size = sizeof(int) * n;
a = (int*)malloc(Size);
// ... 中间数组初始化逻辑保持不变 ...
cudaMemcpy(ptr, a, Size, cudaMemcpyHostToDevice);

// 修正核函数传参与启动配置
for (int i = 0; i < n / 2; i++) {
    // 偶数阶段启动(n+1)/2个线程，传设备指针ptr
    Even<<<1, (n + 1) / 2 >>>(ptr, n);
    // 奇数阶段启动n/2个线程，传设备指针ptr
    Odd<<<1, n / 2 >>>(ptr, n);
    // 可选：添加设备同步与错误捕获
    cudaError_t err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        printf("核函数执行错误: %s\n", cudaGetErrorString(err));
        exit(-1);
    }
}
// 设备数据拷回主机的逻辑保持不变
cudaMemcpy(a, ptr, Size, cudaMemcpyDeviceToHost);

容易被忽略的细节 #

• 时间统计逻辑目前没有排除核函数异步执行的影响，当前统计的时间其实是“发起所有核函数调用的耗时”，并非真正的GPU排序耗时。如果要准确统计性能，需要把cudaDeviceSynchronize()放到计时开始之后、计时结束之前。
• 单块线程的启动方式在n值较大时（超过1024）会直接触发启动失败，实际工程中可以调整为多块网格配置，配合blockIdx.x * blockDim.x + threadIdx.x的全局索引计算方式适配更大的数组规模。

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