Как объявить локальную память в OpenCL?

Я запускаю ядро OpenCL ниже с двумерным глобальным размером работы 1000000 x 100 и локальным размером работы 1 x 100.

__kernel void myKernel(
        const int length, 
        const int height, 
        and a bunch of other parameters) {

    //declare some local arrays to be shared by all 100 work item in this group
    __local float LP [length];
    __local float LT [height];
    __local int bitErrors = 0;
    __local bool failed = false;

    //here come my actual computations which utilize the space in LP and LT
}

Однако он отказывается компилировать, поскольку параметры length и height не известны во время компиляции. Но мне совершенно не ясно, как это сделать правильно. Должен ли я использовать указатели с memalloc? Как справиться с этим так, чтобы память выделялась только один раз для всей рабочей группы, а не один раз за рабочий элемент?

Все, что мне нужно, - это 2 массива float, 1 int и 1 boolean, которые разделены между всей рабочей группой (так что все 100 рабочих элементов). Но я не могу найти какой-либо метод, который сделает это правильно...

Ответ 1

Это относительно просто, вы можете передать локальные массивы в качестве аргументов вашему ядру:

kernel void myKernel(const int length, const int height, local float* LP, 
                     local float* LT, a bunch of other parameters)

Затем вы устанавливаете ядроаргумент с value of NULL и a size равным размеру, который вы хотите выделить для аргумента (в байте). Поэтому это должно быть:

clSetKernelArg(kernel, 2, length * sizeof(cl_float), NULL);
clSetKernelArg(kernel, 2, height* sizeof(cl_float), NULL);

локальная память всегда разделяется рабочей группой (в отличие от частной), поэтому я думаю, что bool и int должны быть точными, но если вы не всегда можете передавать их как аргументы.

На самом деле это не связано с вашей проблемой (и не обязательно релевантно, поскольку я не знаю, какое оборудование вы планируете запускать), но по крайней мере gpus не отличается от процессов, которые не кратно определенной мощности два (я думаю, это было 32 для nvidia, 64 для amd), что означает, что, вероятно, создаст рабочие группы с 128 элементами, из которых последние 28 в основном потрачены впустую. Поэтому, если вы используете opencl на gpu, это может помочь производительности, если вы напрямую используете рабочие группы размером 128 (и соответствующим образом изменяете глобальный размер работы)

В качестве побочного примечания: я никогда не понимал, почему каждый использует вариант подчеркивания для kernel, local and global, кажется мне намного уродливее.

Ответ 2

Вам не нужно выделять всю локальную память вне ядра, особенно если это простая переменная вместо массива.

Причина, по которой ваш код не может скомпилировать, заключается в том, что OpenCL не поддерживает инициализацию локальной памяти. Это указано в документе (https://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/local.html). Это также невозможно в CUDA (Есть ли способ установить значение по умолчанию для массива разделяемой памяти?)

ps: Ответ от Grizzly достаточно хорош, и было бы лучше, если бы я мог опубликовать его в качестве комментария, но меня ограничивает политика репутации. К сожалению.