При передаче параметра по значению в функцию ядра, куда копируются параметры?

Question

Я новичок в программировании CUDA и у меня есть вопрос.

Когда я передаю параметры по значению, например:

__global__ void add(int a, int b, int *c) {
    // some operations
}

Поскольку переменные a и b передаются в функцию ядра add как скопированное значение в стеке вызовов функций, я предположил, что для копирования потребуется некоторое пространство памяти.

Если я прав, это дополнительное пространство памяти, куда эти параметры копируются в GPU или в основную память хоста?

Причина, по которой я задаюсь вопросом, состоит в том, что я должен передать большую структуру функции ядра.

Я также подумал передать указатель структуры, но кажется, что этот способ необходим для вызова cudamalloc для структуры и каждой переменной-члена.

Answer 1

Очень короткий ответ заключается в том, что все аргументы в ядра CUDA передаются по значению, и эти аргументы копируются хостом через API в выделенный буфер аргументов памяти на графическом процессоре. В настоящее время этот буфер хранится в постоянной памяти, и существует ограничение в 4 КБ аргументов на запуск ядра - см. Здесь.

---

Более подробно, стандарт PTX (технически, так как появились вычислительные возможности 2.0 и CUDA ABI) определяет выделенный логический вызов пространства состояний .param которые содержат аргументы параметров ядра и устройства. Смотрите здесь. Цитирую из этой документации:

Каждое определение функции ядра включает в себя необязательный список параметров. Эти параметры являются адресуемыми переменными только для чтения, объявленными в пространстве состояний.param. К значениям, передаваемым от хоста к ядру, обращаются через эти переменные параметра, используя ld.param инструкции. Переменные параметров ядра являются общими для всех CTA в сетке.

Он также отмечает, что:

Примечание. Расположение пространства параметров зависит от конкретной реализации. Например, в некоторых реализациях параметры ядра находятся в глобальной памяти. В этом случае не обеспечивается защита доступа между параметром и глобальным пространством. Аналогично, параметры функции отображаются на регистры передачи параметров и / или расположения в стеке на основе соглашений о вызовах функций двоичного интерфейса приложений (ABI).

Таким образом, точное расположение пространства состояний параметров зависит от конкретной реализации. На первой итерации оборудования CUDA оно фактически отображалось в разделяемую память для аргументов ядра и регистров для аргументов функций устройства. Тем не менее, поскольку аппаратное обеспечение вычисляет 2.0 и стандарт PTX 2.2, в большинстве случаев оно отображается на постоянную память для ядер. Документация говорит следующее по этому вопросу:

Константа (.const) пространство состояний - это постоянная память, инициализированная хостом. Постоянная память доступна с ld.const инструкция. Размер постоянной памяти ограничен размером, в настоящее время ограниченным 64 КБ, который может использоваться для хранения постоянных переменных статического размера. Существует еще 640 КБ постоянной памяти, организованной в виде десяти независимых областей по 64 КБ. Драйвер может выделять и инициализировать постоянные буферы в этих областях и передавать указатели на буферы в качестве параметров функции ядра. Поскольку десять областей не являются смежными, драйвер должен гарантировать, что постоянные буферы распределены так, чтобы каждый буфер полностью помещался в область размером 64 КБ и не выходил за границы области.

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

[Акцент мой]

Таким образом, хотя аргументы ядра хранятся в постоянной памяти, это не та же самая постоянная память, которая отображается на .const пространство состояний, доступное путем определения переменной как __constant__ в CUDA C или эквивалент в Fortran или Python. Скорее, это внутренний пул памяти устройства, управляемый драйвером и не доступный напрямую для программиста.