Где хранятся данные параметра ядра?

Question

Как в заголовке, в программах cuda, где параметр ядра находится после запуска ядра, в локальной памяти или глобальной памяти GPU?

Например, в LLVM IR программы Cuda:

__global__ kernel(int param1):

%0 = alloca int

store param1, %0

Таким образом, в данном случае, где делает% 0 пункт? локальной памяти или глобальной памяти?

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

Answer 1

Как заметил в своем комментарии Роберт Корвелла: параметры хранятся в памяти памяти GPU.

Однако делает alloca, а хранилище param1 в выделенное пространство перемещает, копирует параметр из постоянной памяти в локальную память. alloca инструкции младшего разряда для распределения стека в коде PTX. В clang это канонический способ обработки параметров функции во время генерации кода. Однако на графических процессорах это может (поскольку PTX оптимизируется при понижении до SASS, просто говоря: может) приведет к снижению производительности, поскольку локальная память проходит через все уровни кэша до глобальной памяти и намного медленнее, чем постоянная память.

В LLVM у вас есть пропуск для оптимизатора mem2reg. Этот проход способствует распределению всех папок памяти в стек для регистров. В случае параметров ядра вы, скорее всего, хотите эту оптимизацию. Команды alloca и store исчезнут из вашего ИК-диапазона, и этот параметр будет использоваться непосредственно вместо ненужной копии.