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 должна быть общей.