GPU排序与CPU排序

Question

我对归并排序算法做了一个非常简单的实现，然后我将其用于 CUDA，只需很少的实现更改，算法代码如下：

//Merge for mergesort
__device__ void merge(int* aux,int* data,int l,int m,int r)
{
    int i,j,k;
    for(i=m+1;i>l;i--){
        aux[i-1]=data[i-1];
    }
    //Copy in reverse order the second subarray
    for(j=m;j<r;j++){
        aux[r+m-j]=data[j+1];
    }
    //Merge
    for(k=l;k<=r;k++){
        if(aux[j]<aux[i] || i==(m+1))
            data[k]=aux[j--];
        else
            data[k]=aux[i++];
    }
}

//What this code do is performing a local merge
//of the array
__global__
void basic_merge(int* aux, int* data,int n)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    int tn = n / (blockDim.x*gridDim.x);
    int l = i * tn;
    int r = l + tn;
    //printf("Thread %d:  %d,%d: \n",i,l,r);
    for(int i{1};i<=(tn/2)+1;i*=2)
        for(int j{l+i};j<(r+1);j+=2*i)
        {
            merge(aux,data,j-i,j-1,j+i-1);
        }
    __syncthreads();
    if(i==0){
        //Complete the merge
        do{
            for(int i{tn};i<(n+1);i+=2*tn)
                merge(aux,data,i-tn,i-1,i+tn-1);
            tn*=2;
        }while(tn<(n/2)+1);
    }
}

问题是，无论我在 GTX 760 上启动多少个线程，排序性能总是比在 CPU 上运行 8 个线程的相同代码差得多（我的 CPU 硬件支持最多 8 个并发线程）。

例如，在 CPU 上对 1.5 亿个元素进行排序需要几百毫秒，在 GPU 上则需要长达 10 分钟（即使每个块有 1024 个线程）！显然我在这里遗漏了一些重要的观点，你能给我一些评论吗？我强烈怀疑问题出在第一个线程执行的最终合并操作中，此时我们有一定数量的子数组（确切的数量取决于线程数），这些子数组已排序并需要我合并，这是仅由一个线程（一个微小的 GPU 线程）完成。

我认为我应该在这里使用某种减少，因此每个线程并行执行更多合并，并且“完成合并”步骤只是合并最后两个排序的子数组。

我对 CUDA 很陌生。

编辑（附录）：

感谢您的链接，我必须承认，在充分利用该材料之前，我仍然需要一些时间来学习更好的 CUDA。无论如何，我能够重写排序函数，以便尽可能长时间地利用多线程，我的第一个实现在合并过程的最后阶段存在瓶颈，该过程仅由一个多处理器执行。

现在，在第一次合并之后，我每次最多使用 (1/2)*(n/b) 个线程，其中 n 是要排序的数据量，b 是每个线程排序的数据块的大小。

性能的提升是令人惊讶的，仅使用 1024 个线程，对 3000 万个元素进行排序大约需要 10 秒。不幸的是，这仍然是一个糟糕的结果！问题出在线程同步上，但首先，让我们看一下代码：

__global__
void basic_merge(int* aux, int* data,int n)
{
    int k = blockIdx.x*blockDim.x + threadIdx.x;
    int b = log2( ceil( (double)n / (blockDim.x*gridDim.x)) ) + 1;
    b = pow( (float)2, b);
    int l=k*b;
    int r=min(l+b-1,n-1);
    __syncthreads();
    for(int m{1};m<=(r-l);m=2*m)
    {
        for(int i{l};i<=r;i+=2*m)
        {
            merge(aux,data,i,min(r,i+m-1),min(r,i+2*m-1));
        }
    }
    __syncthreads();
    do{
        if(k<=(n/b)*.5)
        {
            l=2*k*b;
            r=min(l+2*b-1,n-1);
            merge(aux,data,l,min(r,l+b-1),r);
        }else break;
        __syncthreads();
        b*=2;
    }while((r+1)<n);
}

“合并”功能与以前相同。现在的问题是，我只使用 1024 个线程，而不是可以在 CUDA 设备上运行的 65000 个甚至更多线程，问题是 __syncthreads 不能在网格级别作为同步原语工作，而只能在块级别工作！

所以我最多可以同步 1024 个线程，这是每个块支持的线程数量。如果没有适当的同步，每个线程都会弄乱另一个线程的数据，并且合并过程将不起作用。

为了提高性能，我需要在网格中的所有线程之间进行某种同步，似乎不存在用于此目的的 API，并且我读到了一个涉及从主机代码启动多个内核的解决方案，使用主机作为屏障对于所有线程。

我对如何在我的合并排序函数中实现这项技术有一定的计划，我将在不久的将来向您提供代码。您自己有什么建议吗？

谢谢

Answer 1

看起来所有的工作都是在 __global __ 内存中完成的。每次写入都需要很长时间，每次读取也需要很长的时间，这使得函数变慢。我认为首先将数据复制到 __shared __ 内存，然后在那里完成工作，然后在排序完成时（对于该块）将结果复制回全局内存会有所帮助。

全局内存大约需要 400 个时钟周期（如果数据恰好位于 L2 缓存中，则大约需要 100 个时钟周期）。另一方面，共享内存只需要 1-3 个时钟周期即可写入和读取。

上述内容对性能有很大帮助。您可以尝试的其他一些超级小事情是.. (1) 删除第一个 __syncthreads(); 它实际上并没有做任何事情，因为此时扭曲之间没有数据经过。(2) 移动“int b = log2( ceil( (double)n / (blockDim.x*gridDim.x)) ) + 1; b = pow( (float)2, b);” 在内核之外，只需传入 b 即可。这是被一遍又一遍地计算，而实际上只需要计算一次。

我试图遵循你的算法，但没能做到。变量名称很难理解......或者......你的代码超出了我的理解范围，我无法理解。=) 希望以上内容有所帮助。