在 CUDA 中使用共享内存实现并行合并不会提高性能答案

【问题标题】：Implementing parallel merging in CUDA with shared memory does not enhance performance在 CUDA 中使用共享内存实现并行合并不会提高性能
【发布时间】：2013-01-24 20:44:05
【问题描述】：

我正在尝试在 CUDA 中实现并行合并算法。该算法设计为在一个线程块中执行。它的基本思想是计算两个输入序列中每个元素的全局等级。由于这两个输入序列是排序的，一个元素的全局排名等于它在其原始序列中的索引加上它在另一个序列中的排名，这是通过二进制搜索计算的。我认为实现这种算法的最佳策略是将两个序列加载到共享内存中以减少全局内存读取。但是，当我比较两个版本的实现时，一个使用共享内存，一个不使用共享内存，我看不到性能增强。我想知道我是不是做错了什么。

硬件：GeForce GTX 285、Linux x86_64。两种实现合并两个 1024 个元素序列的时间约为 0.068672 毫秒。

__global__ void localMerge(int * A, int numA,int * B,int numB,int * C){
extern __shared__ int  temp[]; // shared memory for A and B; 
int tx=threadIdx.x;
int size=blockDim.x;
int *tempA=temp;
int *tempB=temp+numA;

int i,j,k,mid;
    //read sequences into shared memory 
for(i=tx;i<numA;i+=size){
    tempA[i]=A[i];
}
for(i=tx;i<numB;i+=size){
    tempB[i]=B[i];
}
__syncthreads();
    //compute global rank for elements in sequence A
for(i=tx;i<numA;i+=size){
    j=0;
    k=numB-1;
    if(tempA[i]<=tempB[0]){
        C[i]=tempA[i];
    }
    else if(tempA[i]>tempB[numB-1]){
        C[i+numB]=tempA[i];
    }
    else{
        while(j<k-1){
            mid=(j+k)/2;
            if(tempB[mid]<tempA[i]){
                j=mid;
            }
            else{
                k=mid;
            }
        }
        //printf("i=%d,j=%d,C=%d\n",i,j,tempA[i]);
        C[i+j+1]=tempA[i];
    }
}   
    //compute global rank for elements in sequence B
for(i=tx;i<numB;i+=size){
    j=0;
    k=numA-1;
    if(tempB[i]<tempA[0]){
        C[i]=tempB[i];
    }
    else if(tempB[i]>=tempA[numA-1]){
        C[i+numA]=tempB[i];
    }
    else{
        while(j<k-1){
            mid=(j+k)/2;
            if(tempA[mid]<=tempB[i]){
                j=mid;
            }
            else{
                k=mid;
            }
        }
        //printf("i=%d,j=%d,C=%d\n",i,j,tempB[i]);
        C[i+j+1]=tempB[i];
    }
}    
}

【问题讨论】：

所以我没有误会，你是用 1 个块运行这个内核？
有这么多循环和分支语句，我很惊讶 cuda 版本并不慢。同样，在处理并行架构时，即使串行代码的实现速度较慢，双音合并也是前进的方法。
@talonmies 是的。我用一个线程块执行内核。
@xhe8：运行单个块不会让您隐藏架构中的大部分延迟，而且这些延迟可能会比共享内存等优化带来的任何性能改进带来更大的降低。
@talonmies 实际上我确实实现了一种并行合并算法，它将长输入序列划分为许多较小的子序列，并且一个线程块执行上述算法来合并每对子序列。我发布这个问题的原因是我想知道为什么使用共享内存的实现不如没有共享内存的实现。

标签： cuda merge shared-memory

【解决方案1】：

应用“merge path”算法可能比通过__shared__ 内存中的两个输入列表依赖并行细粒度二进制搜索集合更幸运。使用__shared__ 内存来解决这个问题不太重要，因为缓存可以很好地捕获该算法中存在的重用。

使用此合并算法，其想法是 CTA 的每个线程负责在合并结果中生成 k 输出。这有一个很好的特性，即每个线程的工作大致一致，并且涉及的二进制搜索是相当粗粒度的。

线程i同时搜索两个输入列表，以找到k*ith 输出元素在每个列表中的位置。然后工作就很简单了：每个线程依次合并输入列表中的k 项并将它们复制到输出中的位置k*i。

详情可参考Thrust's implementation。

【讨论】：