CUDA归并排序GPU函数输入超65536时失效，求配置指导

我正在编写代码对比CPU与GPU基于CUDA执行归并排序的性能，CPU端函数运行正常，但GPU端函数在输入数组大小超过65536时失效。作为CUDA新手，我知道需要调整线程块与线程数的配置，但不清楚具体操作方法。相关代码如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda.h>
#include <iostream>
#include <cstdlib>
#include <ctime>
#include <algorithm>

__global__ void mergeSortGPU(int* arr, int* temp, int size, int mergeSize) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    int start = tid * mergeSize * 2;
    int mid = start + mergeSize;
    int end = start + 2 * mergeSize;

    if (start >= size || mid >= size)
        return;

    if (end > size)
        end = size;

    int i = start;
    int j = mid;
    int k = start;

    while (i < mid && j < end) {
        if (arr[i] <= arr[j])
            temp[k++] = arr[i++];
        else
            temp[k++] = arr[j++];
    }

    while (i < mid)
        temp[k++] = arr[i++];
    while (j < end)
        temp[k++] = arr[j++];

    for (int idx = start; idx < end; ++idx)
        arr[idx] = temp[idx];
}

int main() {
    int size;
    std::cout << "Enter the size of the array: ";
    std::cin >> size;

    // Allocate memory for the array
    int* arr = new int[size];
    int* carr = new int[size];
    int* temp = new int[size];

  srand(static_cast<unsigned int>(time(nullptr)));
    for (int i = 0; i < size; ++i) {
        arr[i] = rand() % 100;
        carr[i] = arr[i];
    }


    // GPU variables
    int* gpuArr;
    int* gpuTemp;
    int maxThreadsPerBlock = 1024;
    int threadsPerBlock = std::min(1024, size / 2);
    int blocksPerGrid = (size + (2 * threadsPerBlock) - 1) / (2 * threadsPerBlock);
    blocksPerGrid = std::max(blocksPerGrid, 1);

    // Allocate memory on GPU
    cudaMalloc((void**)&gpuArr, size * sizeof(int));
    cudaMalloc((void**)&gpuTemp, size * sizeof(int));

    // Copy the input array to GPU memory
    cudaMemcpy(gpuArr, arr, size * sizeof(int), cudaMemcpyHostToDevice);


    for (int mergeSize = 1; mergeSize < size; mergeSize *= 2) {
        mergeSortGPU <<<blocksPerGrid, threadsPerBlock>>> (gpuArr, gpuTemp, size, mergeSize);
        cudaDeviceSynchronize();
    }


    // Free allocated memory
    delete[] arr;
    delete[] carr;
    delete[] temp;
    cudaFree(gpuArr);
    cudaFree(gpuTemp);

    return 0;
}

我曾尝试将blocksPerGrid和threadsPerBlock设为最大值，但感觉这并非正确解决方法。

问题根源 #

你的代码中线程块和线程数仅在初始化时计算一次，但归并排序的每一轮需要处理的合并段数量是动态变化的：

• 第一轮mergeSize=1，需要处理size/2个合并段对
• 第二轮mergeSize=2，需要处理size/4个合并段对
• 以此类推，每一轮需要的线程数为ceil(size/(2*mergeSize))

初始的线程配置只适配第一轮，后续轮次要么线程数量不足（无法处理所有合并段），要么线程过多（导致大量线程提前返回，浪费资源），当数组超过65536时，这种不匹配会直接导致排序失效。

正确的线程配置方法 #

每一轮归并前，都需要根据当前mergeSize重新计算所需的线程数，再分配线程块和线程：

1. 计算当前轮次需要的总线程数：int numThreads = (size + 2*mergeSize -1) / (2*mergeSize);（向上取整）
2. 线程数上限设为CUDA支持的最大线程块大小（通常为1024）：int threadsPerBlock = std::min(maxThreadsPerBlock, numThreads);
3. 计算所需线程块数量：int blocksPerGrid = (numThreads + threadsPerBlock -1) / threadsPerBlock;（向上取整）

修改后的代码 #

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda.h>
#include <iostream>
#include <cstdlib>
#include <ctime>
#include <algorithm>

__global__ void mergeSortGPU(int* arr, int* temp, int size, int mergeSize) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    int start = tid * mergeSize * 2;
    int mid = start + mergeSize;
    int end = start + 2 * mergeSize;

    if (start >= size)
        return;

    if (mid >= size) {
        // 只剩单个段，无需合并
        return;
    }

    if (end > size)
        end = size;

    int i = start;
    int j = mid;
    int k = start;

    while (i < mid && j < end) {
        if (arr[i] <= arr[j])
            temp[k++] = arr[i++];
        else
            temp[k++] = arr[j++];
    }

    while (i < mid)
        temp[k++] = arr[i++];
    while (j < end)
        temp[k++] = arr[j++];

    for (int idx = start; idx < end; ++idx)
        arr[idx] = temp[idx];
}

int main() {
    int size;
    std::cout << "Enter the size of the array: ";
    std::cin >> size;

    // Allocate memory for the array
    int* arr = new int[size];
    int* carr = new int[size];
    int* temp = new int[size];

    srand(static_cast<unsigned int>(time(nullptr)));
    for (int i = 0; i < size; ++i) {
        arr[i] = rand() % 100;
        carr[i] = arr[i];
    }

    // CPU排序对比
    std::sort(carr, carr + size);

    // GPU variables
    int* gpuArr;
    int* gpuTemp;
    const int maxThreadsPerBlock = 1024;

    // Allocate memory on GPU
    cudaMalloc((void**)&gpuArr, size * sizeof(int));
    cudaMalloc((void**)&gpuTemp, size * sizeof(int));

    // Copy the input array to GPU memory
    cudaMemcpy(gpuArr, arr, size * sizeof(int), cudaMemcpyHostToDevice);

    // 每轮归并前动态计算线程配置
    for (int mergeSize = 1; mergeSize < size; mergeSize *= 2) {
        int numThreads = (size + 2 * mergeSize - 1) / (2 * mergeSize);
        int threadsPerBlock = std::min(maxThreadsPerBlock, numThreads);
        int blocksPerGrid = (numThreads + threadsPerBlock - 1) / threadsPerBlock;

        mergeSortGPU <<<blocksPerGrid, threadsPerBlock>>> (gpuArr, gpuTemp, size, mergeSize);
        cudaDeviceSynchronize();

        // 检查CUDA错误，调试必备
        cudaError_t err = cudaGetLastError();
        if (err != cudaSuccess) {
            std::cerr << "CUDA Error: " << cudaGetErrorString(err) << std::endl;
            return -1;
        }
    }

    // 将GPU结果拷回CPU验证正确性
    int* gpuResult = new int[size];
    cudaMemcpy(gpuResult, gpuArr, size * sizeof(int), cudaMemcpyDeviceToHost);

    // 对比CPU和GPU结果
    bool isCorrect = true;
    for (int i = 0; i < size; ++i) {
        if (gpuResult[i] != carr[i]) {
            isCorrect = false;
            std::cerr << "Mismatch at index " << i << ": CPU=" << carr[i] << ", GPU=" << gpuResult[i] << std::endl;
            break;
        }
    }
    if (isCorrect) {
        std::cout << "Sorting is correct!" << std::endl;
    }

    // Free allocated memory
    delete[] arr;
    delete[] carr;
    delete[] temp;
    delete[] gpuResult;
    cudaFree(gpuArr);
    cudaFree(gpuTemp);

    return 0;
}

额外注意事项 #

• CUDA错误检查：一定要在核函数调用后检查错误（cudaGetLastError()），这能帮你快速定位线程配置错误、内存访问越界等问题。
• 数组边界处理：原代码中if (start >= size || mid >= size)的判断可以简化为if (start >= size)，因为当start < size但mid >= size时，单个段无需合并，直接返回即可。
• 性能优化：当前实现是基础的归并排序，还可以进一步优化（比如使用共享内存减少全局内存访问），但先确保正确性再优化。

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