NVIDIA __constant memory: как заполнить постоянную память с хоста как в OpenCL, так и в CUDA?

У меня есть буфер (массив) на хосте, который должен находиться в области постоянной памяти устройства (в данном случае это графический процессор NVIDIA).

Итак, у меня есть два вопроса:

  1. Как я могу выделить кусок постоянной памяти? Учитывая тот факт, что я отслеживаю доступную постоянную память на устройстве и точно знаю, что у нас есть доступный нам объем памяти (на данный момент)

  2. Как я могу инициализировать (заполнить) эти массивы значениями, которые вычисляются во время выполнения на хосте?

Я искал это в Интернете, но нет четкого документа, документирующего это. Я был бы признателен, если бы предоставленные примеры были как в OpenCL, так и в CUDA. Пример для OpenCL для меня важнее, чем CUDA.

Для CUDA: stackoverflow.com/questions/28987495/…

Sebastian 07.05.2022 00:11

Cuda может обрабатывать 704 КБ постоянной памяти на ядро ​​с непрерывными массивами до 64 КБ, если я правильно прочитал этот раздел: docs.nvidia.com/cuda/parallel-thread-execution/…

Sebastian 07.05.2022 07:23

@Sebastian: Ваше прочтение не совсем верно. Для каждого ядра доступно только 64 КБ постоянной памяти программиста. Другие банки зарезервированы драйвером для внутренних статических распределений, таких как аргументы ядра.

talonmies 07.05.2022 07:32

@talonmies По умолчанию (только) банк 2 используется для определяемой пользователем постоянной памяти в C. Ассемблер (ISA) может получить доступ ко всем банкам (сам никогда не пробовал). См. docs.nvidia.com/cuda/parallel-thread-execution/…, как это делалось до версии 2.1, и docs.nvidia.com/cuda/parallel-thread-execution/…, как это делалось, начиная с версии 2.2. Вы можете передать ядрам несколько указателей на разные константные пространства.

Sebastian 07.05.2022 07:50
Стоит ли изучать PHP в 2023-2024 годах?
Стоит ли изучать PHP в 2023-2024 годах?
Привет всем, сегодня я хочу высказать свои соображения по поводу вопроса, который я уже много раз получал в своем сообществе: "Стоит ли изучать PHP в...
Поведение ключевого слова "this" в стрелочной функции в сравнении с нормальной функцией
Поведение ключевого слова "this" в стрелочной функции в сравнении с нормальной функцией
В JavaScript одним из самых запутанных понятий является поведение ключевого слова "this" в стрелочной и обычной функциях.
Приемы CSS-макетирования - floats и Flexbox
Приемы CSS-макетирования - floats и Flexbox
Здравствуйте, друзья-студенты! Готовы совершенствовать свои навыки веб-дизайна? Сегодня в нашем путешествии мы рассмотрим приемы CSS-верстки - в...
Тестирование функциональных ngrx-эффектов в Angular 16 с помощью Jest
В системе управления состояниями ngrx, совместимой с Angular 16, появились функциональные эффекты. Это здорово и делает код определенно легче для...
Концепция локализации и ее применение в приложениях React ⚡️
Концепция локализации и ее применение в приложениях React ⚡️
Локализация - это процесс адаптации приложения к различным языкам и культурным требованиям. Это позволяет пользователям получить опыт, соответствующий...
Пользовательский скаляр GraphQL
Пользовательский скаляр GraphQL
Листовые узлы системы типов GraphQL называются скалярами. Достигнув скалярного типа, невозможно спуститься дальше по иерархии типов. Скалярный тип...
2
5
67
3
Перейти к ответу Данный вопрос помечен как решенный

Ответы 3

  1. How can I allocate a chunk of constant memory? Given the fact that I am tracing the available constant memory on the device and I know, for a fact, that we have that amount of memory available to us (at this time)

В CUDA нельзя. Нет выделения постоянной памяти во время выполнения, только статическое определение памяти с помощью спецификатора __constant__, которое сопоставляется со страницами постоянной памяти при сборке. Вы мог генерируете некоторый код, содержащий такое статическое объявление во время выполнения, и компилируете его с помощью nvrtc, но это кажется большим усилием для чего-то, что, как вы знаете, может иметь размер только до 64 КБ. Мне кажется гораздо проще (по крайней мере, мне) просто статически объявить постоянный буфер размером 64 КБ и использовать его во время выполнения по своему усмотрению.

  1. How can I initialize (populate) those arrays from values that are computed at the runtime on the host?

Как отмечено в комментариях, см. здесь. Для этой цели был создан API cudaMemcpyToSymbol, который работает так же, как стандартный memcpy.

Функционально нет разницы между __constant в OpenCL и __constant__ в CUDA. Применяются те же ограничения: статическое определение во время компиляции (которое является временем выполнения в стандартной модели выполнения OpenCL), ограничение в 64 КБ.

Отказ от ответственности: я не могу помочь вам с CUDA.

Для OpenCL constant память фактически рассматривается как global память только для чтения с точки зрения программиста/API или определена встроенно в исходный код ядра.

  1. Определите константные переменные, массивы и т. д. в коде ядра, например constant float DCT_C4 = 0.707106781f;. Обратите внимание, что вы можете динамически генерировать код ядра на хосте во время выполнения, чтобы генерировать производные постоянные данные, если хотите.
  2. Передайте постоянную память от хоста к ядру через объект буфера, как и для global памяти. Просто укажите параметр указателя в области памяти constant в прототипе функции ядра и установите буфер на стороне хоста с помощью clSetKernelArg(), например:
kernel void mykernel(
    constant float* fixed_parameters,
    global const uint* dynamic_input_data,
    global uint* restrict output_data)
{
    cl_mem fixed_parameter_buffer = clCreateBuffer(
        cl_context,
        CL_MEM_READ_ONLY | CL_MEM_HOST_NO_ACCESS | CL_MEM_COPY_HOST_PTR,
        sizeof(cl_float) * num_fixed_parameters, fixed_parameter_data,
        NULL);
    clSetKernelArg(mykernel, 0, sizeof(cl_mem), &fixed_parameter_buffer);

Обязательно учитывайте значение, указанное для CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, для используемого контекста! Обычно не помогает использование constant буферов памяти для потоковой передачи входных данных, это лучше хранить в global буферах, даже если они помечены для ядра как доступные только для чтения. constant память наиболее полезна для данных, которые используются большей частью рабочих элементов. Обычно для него существует довольно жесткое ограничение по размеру, например 64 КБ - некоторые реализации могут «разлиться» на глобальную память, если вы попытаетесь превысить это, что потеряет все преимущества производительности, которые вы получили бы от использования постоянной памяти.

Спасибо за Ваш ответ. Пытается ли API clCreateBuffer вернуть указатель на constant float* fixed_parameters? Если да, то как он узнает, какой символ мы ищем?

mgNobody 09.05.2022 08:30
clCreateBuffer возвращает дескриптор объекта буфера (cl_mem), который необходимо освободить, когда вы закончите с ним работать. В примере я копирую массив хостов fixed_parameter_data в буфер. Однако есть много других способов заполнения буферов OpenCL — это точно так же, как и для global памяти.
pmdj 09.05.2022 08:45

В OpenCL нет разрешения символов между кодом хоста и ядра, за исключением обращения к функциям ядра по имени в API хоста. Второй аргумент clSetKernelArg — это индекс аргумента ядра, здесь 0, потому что я случайно указал буфер констант в качестве первого параметра функции ядра.

pmdj 09.05.2022 08:48

Я понимаю. Итак, по сути, вы выступаете за создание буфера READ_ONLY в глобальной памяти. Я искал способ справиться с __constant памятью на устройстве (используя память текстур).

mgNobody 09.05.2022 19:48

Нет, содержимое буфера будет скопировано в constant память (кэш с прямой адресацией/регистры/все, что использует аппаратное обеспечение) при постановке ядра в очередь. Это вызывает объявление аргумента constant в сигнатуре функции ядра вместо global. На стороне хоста вы управляете этими данными идентично глобальной памяти.

pmdj 09.05.2022 19:53
Ответ принят как подходящий

Для cuda я использую API драйвера и NVRTC и создаю строку ядра с глобальным массивом констант следующим образом:

auto kernel = R"(
..
__constant__ @@Type@@ buffer[@@SIZE@@] = {
   @@elm@@
};
..
__global__ void test(int * input)
{   }

)";   

затем замените слова @@-шаблона информацией о размере и значении элемента во время выполнения и компилировать следующим образом:

__constant__ int buffer[16384] = { 1,2,3,4, ....., 16384 };

Итак, это время выполнения для хоста, время компиляции для устройства. Недостатком является то, что строка ядра становится слишком большой, менее удобочитаемой, а соединяющие классы требуют явного связывания (как если бы вы компилировали сторонний проект C++) другие единицы компиляции. Но для простых вычислений только с вашими собственными реализациями (без непосредственного использования хост-определений) это то же самое, что и API времени выполнения.

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

Вы уверены, что просто __constant__ стоит затраченных усилий? Есть ли у вас какие-то результаты тестов, чтобы показать, что действительно улучшает производительность? (преждевременная оптимизация - источник всех зол). Возможно, ваш алгоритм работает с регистровым тайлингом и источник данных значения не имеет?

Спасибо @huseyin. Это имеет смысл. Думаю, я искал более формальный способ (например, enqueueMemCopy) справиться с этим.

mgNobody 09.05.2022 18:34

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

huseyin tugrul buyukisik 09.05.2022 18:45

Кроме того, использование определений класса хоста требует связывания через API-интерфейс драйвера. Это похоже на написание dll для проекта C++ и использование хоста в качестве dll с точки зрения ядра cuda.

huseyin tugrul buyukisik 09.05.2022 18:48

Я думаю, что talonmies дал лучший ответ здесь.

huseyin tugrul buyukisik 09.05.2022 18:57

Спасибо за вашу помощь. На самом деле, ваш подход хорош для OpenCL. Компиляция и сборка исходного файла CL в OpenCL — очень распространенная вещь (как вы знаете). Единственным недостатком является то, что нам приходится выполнять некоторые манипуляции с текстом и компилировать/собирать ядра с нуля.

mgNobody 09.05.2022 19:45

Другие вопросы по теме