c++ - 使用 Cuda 的排序算法。内仁还是外仁？

Question

我有一个大小为 50000x100 的矩阵，我需要使用 C++ 中的 Cuda 对每一行进行排序。我的架构是 K80 NVidia 卡。

由于列数较少，我目前正在内核中运行排序算法。我正在使用在矩阵的所有行上运行的修改后的气泡算法。

我想知道是否有更有效的方法来进行。我尝试在我的内核中使用 thrust::sort 但它要慢得多。我还尝试了合并排序算法，但该算法的递归部分在我的内核中不起作用。

==编辑==

这是我的内核:

__global__ void computeQuantilesKernel(float *matIn, int nRows, int nCols, int nQuantiles, float *outsideValues, float *quantilesAve, int param2)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float values[100];//big enough for 100 columns
    int keys[100];
    int nQuant[100];//big enough for 100 quantiles (percentiles)
    float thisQuantile[100];
    int quant;

    if (idx >= nRows) return;

    //read matIn from global memory
    for (int i = 0; i < nCols; i++)
    {
        values[i] = matIn[idx * nCols + i + param2 * nCols * nRows];
        keys[i] = i;
    }

    //bubble Sort:
    int i, j;
    int temp;
    float tempVal;

    for (i = 0; i < nCols - 1; i++)
    {
        for (j = 0; j < nCols - i - 1; j++)
        {
            if (values[j + 1] < values[j])      // ascending order simply changes to <
            {
                tempVal = values[j];             // swap elements
                temp = keys[j];             // swap elements
                values[j] = values[j + 1];
                keys[j] = keys[j + 1];
                values[j + 1] = tempVal;
                keys[j + 1] = temp;
            }
        }
    }
    //end of bubble sort

    //reset nQuant and thisQuantile
    for (int iQuant = 0; iQuant < nQuantiles; iQuant++)
    {
        nQuant[iQuant] = 0;
        thisQuantile[iQuant] = 0;
    }

    //Compute sum of outsideValues for each quantile
    for (int i = 0; i < nCols; i++)
    {
        quant = (int)(((float)i + 0.5) / ((float)nCols / (float)nQuantiles));//quantile like Matlab
        nQuant[quant]++;
        thisQuantile[quant] += outsideValues[idx * nCols + keys[i]];
    }

    //Divide by the size of each quantile to get averages
    for (int iQuant = 0; iQuant < nQuantiles; iQuant++)
    {
        quantilesAve[idx + nRows * iQuant + param2 * nQuantiles * nRows] = thisQuantile[iQuant] / (float)nQuant[iQuant];
    }
}

Answer 1

您的代码目前使用单个线程来分别处理每一行。因此，您急需快速暂存内存(寄存器、L1 缓存、共享内存)。您为每个线程至少分配了 1600 个字节——很多!您希望每个线程保持在 128 字节左右(32 个寄存器，每个寄存器 32 位)。其次，您使用的是在运行时可寻址的本地数组——这些数组将溢出到本地内存中，丢弃您的 L1 缓存并再次进入全局内存(1600B x 32 线程提供 51KB，这已经达到或超过限制shmem/L1)。

出于这个原因，我建议改为在 64 或 128 个线程的每个 block 中处理一行，并将您排序的行保留在共享内存中。冒泡排序其实很容易并行实现:

__shared__ float values[nCols];
... load the data ...
__syncthreads();
for (int i = 0; i < nCols/2; i++)
{
    int j = threadIdx.x;
    if (j % 2 == 0 && j<nCols-1)
        if (values[j+1] < values[j])
            swap(values[j+1], values[j]);
    __syncthreads();
    if (j % 2 == 1 && j<nCols-1)
        if (values[j+1] < values[j])
            swap(values[j+1], values[j]);
    __syncthreads();
}

请注意您的内部 for j = ... 循环如何被 threadIdx 替换，但算法的核心思想保持不变。在每次迭代中，我首先仅对偶数对执行气泡交换，然后仅对奇数对执行气泡交换以避免并行冲突。

我假设 nCols 低于您的 block 的维度，这对于 100 个元素来说很容易实现。

上面的代码还有很多方法可以进一步改进，例如

• 将线程数减半，假设循环的前半部分为 j=threadIdx.x*2，j=threadIdx.x*2+1下半场。这样就没有线程处于空闲状态。
• 仅使用 32 个线程，每个线程依次处理 j 的两个值。这样你的问题将适合单个扭曲，允许你完全放弃 __syncthreads() 。使用 32 个线程，您可以使用 warp shuffle 内在函数。
• 尝试使用 #pragma unroll，尽管生成的代码量可能不可行。分析会有所帮助。

还可以考虑尝试使用硬编码合并排序而不是冒泡排序。如果我没记错的话，当我实现一个扭曲大小的冒泡排序和所有循环展开的合并排序时，合并排序的执行速度几乎是冒泡排序的两倍。请注意，这是几年前的第一代支持 CUDA 的卡。