Cuda 共享内存错误答案

【问题标题】：Cuda shared memory bugCuda 共享内存错误
【发布时间】：2015-03-24 13:04:40
【问题描述】：

我有这个简单的基数排序（它只按位排序，只有一个块）。我的第一个版本有效，但我尝试首先对共享内存上的键进行排序，以实现对 DRAM 的合并内存写入。然而这个版本产生了不好的结果，它没有排序。

第一个简单的工作版本：

__global__ void dev_radix(unsigned int *in_keys, const unsigned int *histo, unsigned int desp, unsigned int *out_keys){
int tid=threadIdx.x;

//Get offset by using prefix sum scan.
__shared__ unsigned int s_sum[1024];
unsigned int first=((in_keys[tid]>>desp)&1)==0;
s_sum[tid]=first;
__syncthreads();
int pos=tid-1;
for (int off=1; pos>=0; off=off*2, pos=tid-off){
    int a=s_sum[pos];
    int b=s_sum[tid];
    __syncthreads();
    s_sum[tid]=a+b;
}
__syncthreads();

int offset=s_sum[tid]-first;
if (first==0){
    //Get offset for '1' bit keys
    offset=histo[0]+tid-offset;
}

out_keys[offset]=in_keys[tid];

}

第二版：

__global__ void dev_radix(unsigned int *in_keys, const unsigned int *histo, unsigned int desp, unsigned int *out_keys){
int tid=threadIdx.x;

//Get offset by using prefix sum scan.
__shared__ unsigned int s_sum[1024];
unsigned int first=((in_keys[tid]>>desp)&1)==0;
s_sum[tid]=first;
__syncthreads();
int pos=tid-1;
for (int off=1; pos>=0; off=off*2, pos=tid-off){
    int a=s_sum[pos];
    int b=s_sum[tid];
    __syncthreads();
    s_sum[tid]=a+b;
}
__syncthreads();

int offset=s_sum[tid]-first;
if (first==0){
    //Get offset for '1' bit keys
    offset=histo[0]+tid-offset;
}

__syncthreads();
s_sum[offset]=in_keys[tid];
__syncthreads();
out_keys[tid]=s_sum[tid];

}

【问题讨论】：

如果没有完整的重现案例和对问题的更具体描述，那么很难说您的代码可能有什么问题，然后“不排序”
我不是 Cuda 专家，但这是一个 4 行差异问题，我认为这不是太难。

标签： cuda parallel-processing

【解决方案1】：

问题是我在条件代码上调用了 __syncthreads()。只允许在对块中的所有线程具有相同执行路径的条件代码上调用 __syncthreads()。正确版本：

__global__ void dev_radix(unsigned int *in_keys, const unsigned int *histo, unsigned int desp, unsigned int *out_keys){
__shared__ unsigned int s_sum[1024];
int tid=threadIdx.x;

//Get offset by using prefix sum scan.
unsigned int v=in_keys[tid];
unsigned int first=((v>>desp)&1)==0;
s_sum[tid]=first;
__syncthreads();
int pos=tid-1;
for (int off=1; off<1024;){
    int a,b;
    if (pos>=0){
        a=s_sum[pos];
        b=s_sum[tid];
    }
    __syncthreads();
    if (pos>=0){
        s_sum[tid]=a+b;
    }
    __syncthreads();
    off=off*2;
    pos=tid-off;
}
__syncthreads();

int offset=s_sum[tid]-first;
if (first==0){
    //Get offset for '1' bit keys
    offset=histo[0]+tid-offset;
}
__syncthreads();
s_sum[offset]=v;
__syncthreads();
out_keys[tid]=s_sum[tid];

}

【讨论】：