Как вы храните данные в быстрой памяти GPU (l1 / shared) при вызовах ядра?

Question

Как вы храните данные в быстрой памяти GPU при вызовах ядра?

Предположим, мне нужно ответить на 1 миллион запросов, каждый из которых содержит около 1,5 МБ данных, которые можно повторно использовать при вызовах, и около 8 КБ данных, уникальных для каждого запроса.

Один из подходов - запускать ядро ​​для каждого запроса, каждый раз копируя 1,5 МБ + 8 КБ данных в общую память. Однако затем я трачу много времени на копирование 1,5 МБ данных, которые действительно могут сохраняться в запросах.

Другой подход состоит в том, чтобы «перезапускать» потоки графического процессора (см. https://stackoverflow.com/a/49957384/3738356).). Это включает в себя запуск одного ядра, которое немедленно копирует 1,5 МБ данных в общую память. Затем ядро ​​ожидает поступления запросов, ожидая для отображения 8 КБ данных перед продолжением каждой итерации. Действительно, кажется, что CUDA не должен использоваться таким образом. Если кто-то просто использует управляемую память и монотонно увеличивающиеся счетчики volatile + для синхронизации, все еще нет гарантии, что данные, необходимые для вычисления ответа, будут находиться на графическом процессоре, когда вы начнете его читать. Вы можете заполнить значения в памяти фиктивными значениями, такими как -42, которые указывают, что значение еще не попало в графический процессор (через механизмы кэширования / управляемой памяти), а затем заняты ожиданием, пока значения не станут действительными. Теоретически это должно сработать. Однако у меня было достаточно ошибок памяти, которые я на данный момент оставил, и преследовал ... .

Другой подход все еще использует переработанные потоки, но вместо этого синхронизирует данные через cudaMemcpyAsync, потоки cuda, события cuda и еще пару изменчивых + монотонно увеличивающихся счетчиков. Я слышал, что для правильной работы cudaMemcpyAsync мне нужно прикрепить 8 КБ данных, которые свежие с каждым запросом. Но асинхронная копия не блокируется - ее эффекты просто не наблюдаемы. Я подозреваю, что с достаточным количеством песка я тоже могу заставить это работать.

Однако все вышеперечисленное заставляет меня думать: «Я делаю это неправильно». Как хранить чрезвычайно многократно используемые данные в кэшах графического процессора, чтобы к ним можно было обращаться от одного запроса к другому?

Answer 1

• Прежде всего наблюдать за эффектами потоков и асинхронного копирования вам определенно нужно закрепить память хоста. Тогда вы можете наблюдать параллельные вызовы ядра "почти" происходят одновременно. Я бы предпочел использовать асинхронное копирование, так как это заставляет меня чувствовать ситуация.
• Во-вторых, вы можете просто сохранить данные в глобальной памяти и загрузить это в общей памяти, когда вам это нужно. Насколько мне известно память известна только самому ядру и уничтожается после прекращение. Попробуйте использовать асинхронные копии во время работы ядра и синхронизировать потоки соответственно. Не забудьте __syncthreads () после загрузки в общую память. Надеюсь, это поможет.