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回報位置的前值被原子更新。這是加載到共享內存時需要使用的值。