создание временных глобальных переменных памяти внутри ядер - PullRequest
Новая
       44

создание временных глобальных переменных памяти внутри ядер

0 голосов
/ 08 мая 2019

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

Есть ли способ выделить временные глобальные переменные изнутри ядра?

В настоящее время я выделяю временный глобальный и передаю указательк моему ядру.Это не очень удобно для пользователя.

TL; DR: требуется временная переменная для атомарного сложения между различными блоками без необходимости явного выделения глобала и передачи указателяк ядру

1 Ответ

0 голосов
/ 08 мая 2019

Вы можете использовать malloc() внутри кода ядра.Тем не менее, это редко хорошая идея сделать это.Как правило, гораздо лучше предварительно выделить пространство для создания нуля перед запуском ядра, передать его в качестве аргумента и дать каждому потоку или группе потоков некоторую формулу для определения местоположения, которое они будут использовать для своей общей атомики в этой области для очистки..

Итак, вы написали, что это не очень "удобно для пользователя";Я предполагаю, что вы имеете в виду дружественные к разработчикам.Ну, это можно сделать более дружелюбным!Например, моя библиотека CUDA Modern C ++ API предлагает эквивалент std::unique_ptr - но для памяти устройства: 

#include <cuda/api_wrappers.hpp>

//... etc. etc. ...

{
    auto scratch = cuda::memory::device::make_unique<float[]>(1024, my_cude_device_id);
    my_kernel<<<blah,blah,blah>>>(output, input, scratch.get();
} // the device memory is released here!

(это, конечно, для синхронных запусков.)

Что-то еще, что вы можете сделать более дружественным для разработчика, - это использовать какую-то прокси-функцию, чтобы получить место в этой временной памяти, относящееся к конкретному потоку / деформации / группе потоков / что угодно, которое использует тот же адрес для атомарного элемента,Это должно, по крайней мере, скрыть некоторую повторяющуюся, раздражающую адресную арифметику, которую может использовать ваше ядро.

Существует также возможность использования глобальных __device__ переменных (как упомянуто @RobertCrovella), но яне поощряет это: размер должен быть фиксирован во время компиляции, и вы не сможете использовать его из двух ядер одновременно, безболезненно и т. д.

...