CUDA: атомарная операция на разделяемой памяти

Question

Мое ядро ​​cuda генерирует что-то, что подается на хост в конце выполнения блока.
Скелет выглядит следующим образом.
host_data, в который записываются данные, выделяется как отображенная память хоста.
host_data_count - также отображенная память, которая указывает количество произведенных данных.
Я использую графический процессор GTX 580 с архитектурой Fermi и CC 2.0.

__global__ void kernel(host_data, host_data_count)
{
    __shared__ int  shd_data[1024];
    __shared__ int  shd_cnt;
    int i;

    if (threadIdx.x == 0)
        shd_cnt = 0;
    __syncthreads();

    while ( ... )
    {
        if (something happens)
        {
            i = atomicAdd(&shd_cnt, 1);
            shd_data[i] = d;
        }
    }

    __syncthreads();
    if (threadIdx.x == 0)
    {
        i = atomicAdd(host_data_count, shd_cnt);
        memcpy(&host_data[i], shd_data, shd_cnt * 4);
    }
}

Чего мне не хватает в этом коде ядра?
Кто-нибудь может помочь?

Answer 1

Трудно сказать, чего вам не хватает, потому что вы на самом деле не указали, с какой проблемой вы столкнулись. Я вижу мало возможностей, но это зависит от фактической реализации некоторых ваших высших концепций.

• Я предполагаю, что host_data_count имеет тип int* (или похожий?). Он либо указывает на глобальную память, либо размещает память через закрепленную на карте память. Я настоятельно рекомендую использовать глобальную память ради скорости.
• Если host_data_count является закрепленной памятью, имейте в виду, что атомарные операции являются атомарными только в графическом процессоре. Если, тем временем, CPU что-то делает с ним, это может сломать атомику. Скорее всего, вам потребуется синхронизировать поток хоста после вызова ядра и перед чтением / использованием значения. Вызовы ядра всегда асинхронны.
• Что такое memcpy в коде устройства? Я полагаю, вы реализовали это сами, верно? Вы копируете память, используя один поток или целый блок? Использование всего блока будет быстрее, но тогда вам нужно использовать эту функцию за пределами if (threadIdx.x==0), а переменная i должна быть общей.