CUDA: копирование ненулевых значений положения вектора в другое

На 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] ;
}

Заранее спасибо

1 2
gpgpu cuda

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 возвращает Предыдущее значение местоположения, которое атомарно обновляется. Это то, что ценят вас нужно использовать при загрузке в общую память.

2