CUDA：将一个向量的非零值位置复制到另一个

Question

在 GPGPU 上，使用 cuda 我的问题是： 我有一个256个元素的向量，我想做一个程序，可以提取非零值的位置并将它们复制到另一个向量中。

我的代码不起作用：

dev_Hist : 是数据源，初始向量；

dev_Xn : 是 dev_Hist 上非零值位置的向量；

nN : 是 dev_Hist 上非零值的数量

1.内核调用：

gpu_Xn<<<1, nN>>>(dev_Hist, nN, dev_Xn) ;

2。设备功能

__global__ void gpu_Xn(int *pHist, int pnN, int* pXn) 
{
    int Tid ;
    Tid = threadIdx.x ;

    __shared__ T tmpXn[256] ;

    tmpXn[Tid] = 0 ;

    __syncthreads() ;

    __shared__ int idx ;

    if(Tid == 0)
        idx = -1  ;

    syncthreads() ;

    if(pHist[Tid] !=0)
    {
        atomicAdd(&idx, 1) ; 
        tmpXn[idx] = Tid ;
    }

    __syncthreads() ;
    if(Tid < pnN)
        pXn[Tid] = tmpXn[Tid] ;
}

Answer 1

这里的问题是您没有正确使用atomicAdd。尽管您以原子方式递增 idx 的值，但 读取 idx 以存储到共享内存不是原子的，这将产生未定义的行为。

您的内核可能应该如下所示：

__global__ void gpu_Xn(int *pHist, int pnN, int* pXn) 
{
    int Tid ;
    Tid = threadIdx.x ;

    __shared__ int tmpXn[256] ;
    __shared__ int idx ;

    tmpXn[Tid] = -1 ;
    if(Tid == 0) idx = 0  ;

    __syncthreads() ;

    if(pHist[Tid] !=0)
    {
        int x = atomicAdd(&idx, 1) ; 
        tmpXn[x] = Tid ;
    }

    __syncthreads() ;
    if(Tid < pnN)
        pXn[Tid] = tmpXn[Tid] ;
}

[免责声明：在浏览器中编写，从未编译，使用风险自负]

请注意，atomicAdd 返回被原子更新的位置的先前值。这是加载到共享内存时需要使用的值。