CUDA Динамически выделяемая постоянная или текстура памяти для массива структур

Question

Мне нужно использовать массив структур в постоянной памяти для моего ядра, где фактический размер массива неизвестен до времени выполнения. Как ответили в Правильный способ использования памяти __constant__ на CUDA? , я понял, что постоянная память выделяется во время компиляции, поэтому массив должен быть объявлен как:

__constant__ SKY_GRID_TYPE const_patch_grid_lat[5];

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

Этот ответ выше предлагает вместо этого использовать текстурную память, которая говорит, что «может быть установлена ​​динамически и кэширована». Однако тип данных, который мне нужен в моей памяти, является массивом struct и в соответствии с Структура в памяти текстур на CUDA , кажется, что память текстур поддерживает только встроенные типы CUDA.

Так есть ли обходной путь? Постоянная память была бы идеальной для моего массива структуры, но размер определяется динамически, поэтому он не работает. Память текстур работала бы, но она не допускает ничего, кроме встроенных типов CUDA. Могу ли я использовать что-нибудь еще или какой-нибудь умный способ обойти это?

Answer 1

Постоянная память - максимум 64 КБ. Насколько я знаю, нет недостатков в распределении всех 64 КБ.

Просто выделите максимальный размер (64 КБ / размер вашей структуры) для вашего массива. Используйте все, что вам нужно. Это также предполагает, что вы будете делать единый доступ через основу для каждого доступа.

Если вам нужно больше 64 кбайт, то, конечно, это не будет работать, но это ставит под сомнение основная предпосылка вашего вопроса.

Для больших постоянных областей и / или для ситуаций, когда у вас нет единообразного доступа, моя рекомендация для cc3.5 и более новых графических процессоров - использовать read- только механизм кэширования (const __restrict__ или __ldg()).