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

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 КБ и использовать его во время выполнения по своему усмотрению.

2. 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 КБ - некоторые реализации могут «разлиться» на глобальную память, если вы попытаетесь превысить это, что потеряет все преимущества производительности, которые вы получили бы от использования постоянной памяти.

Для 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__ стоит затраченных усилий? Есть ли у вас какие-то результаты тестов, чтобы показать, что действительно улучшает производительность? (преждевременная оптимизация - источник всех зол). Возможно, ваш алгоритм работает с регистровым тайлингом и источник данных значения не имеет?