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


Мое ядро ​​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);
    }
}

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


gpu-shared-memory cuda gpu-atomics
person superscalar    schedule 03.02.2012    source источник
comment
Не могли бы вы отредактировать свой вопрос, чтобы объяснить, в чем ваша проблема, потому что это очень неясно. Кроме того, если вам нужна помощь в исправлении определенного фрагмента кода, опубликуйте фактический код, который демонстрирует проблему и который кто-то другой может скомпилировать и запустить. Отладка псевдокода очень сложна.   -  person talonmies    schedule 03.02.2012
comment
Не могли бы вы также подтвердить, что вы компилируете свой код с параметром -arch=sm_20?   -  person pQB    schedule 03.02.2012
comment
@pQB Да, целевая архитектура установлена ​​​​на sm_20   -  person superscalar    schedule 02.03.2012

Ответы (1)


arrow_upward
1
arrow_downward

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

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