Использование переменной устройства несколькими потоками в CUDA

Question

Я играю с Cuda.

На данный момент у меня проблема. Я тестирую большой массив для конкретных ответов, и когда я получаю ответ, мне нужно скопировать данные в другой массив.

Например, мой тестовый массив из 5 элементов выглядит следующим образом:
[] [] [v1] [] [] [v2]

Результат должен выглядеть так:
[V1] [v2]

Проблема в том, как рассчитать адрес второго массива для сохранения результата? Все элементы первого массива проверяются параллельно.

Я думаю объявить переменную устройства int addr = 0. Каждый раз, когда я нахожу ответ, я буду увеличивать addr . Но я не уверен в этом, поскольку это означает, что к addr могут обращаться несколько потоков одновременно. Это вызовет проблемы? Или поток будет ждать, пока другой поток не завершит использование этой переменной?

Answer 1

Не так тривиально, как кажется. Я только что закончил, чтобы реализовать один, и я могу сказать, что вам нужно прочитайте статью scan Gpu Gems 3 , в частности главу 39.3.1 Сжатие потока .

Для реализации собственного запуска из примера LargeArrayScan в SDK, который даст вам только предварительное сканирование. Предполагая, что у вас есть массив выбора в памяти устройства (массив 1 и 0, означающий 1 - выберите 0 - отбросьте), dev_selection_array a dev_elements_array элементов для выбора dev_prescan_array и dev_result_array все по размеру N , тогда вы делаете

prescan(dev_prescan_array,dev_selection_array, N);
scatter(dev_result_array, dev_prescan_array,
         dev_selection_array, dev_elements_array, N);

где разброс

 __global__ void scatter_kernel( T*dev_result_array, 
                   const T* dev_prescan_array, 
                   const T* dev_selection_array,
                   const T* dev_elements_array, std::size_t size){

unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= size) return;
if (dev_selection_array[idx] == 1){
    dev_result_array[dev_prescan_array[idx]] = dev_elements_array[idx];
}
}

для другого хорошего применения предварительного сканирования см. Бумагу Ble93

Веселись!

Answer 2

Вы говорите о классическом сжатии потока. Обычно я бы рекомендовал посмотреть Thrust или CUDPP (эти ссылки приведены в документации по уплотнению). Оба из них с открытым исходным кодом, если вы хотите накатить свой собственный, то я бы также посоветовал взглянуть на образец «отсканированного» SDK.