CUDA : Copy non-zero values position of a vector to another

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 存储到共享内存的读取不是原子的，这将产生 undefined行为。

你的内核可能应该是这样的:

__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 返回被原子更新的位置的先前值。这是加载到共享内存时需要使用的值。