用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 个线程中处理一行，并将排序的行保留在共享内存中。冒泡排序实际上很容易并行实现：

__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低于你的块的尺寸，对于 100 个元素来说这是很容易实现的。

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

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

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