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

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

15 голосов
/ 17 января 2012

Я использую ядро ​​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 массива с плавающей точкой, 1 int и 1 boolean, которые являются общими для всей рабочей группы (то есть для всех 100 рабочих элементов). Но я не могу найти метод, который делает это правильно ...

Ответы [ 3 ]

25 голосов
/ 17 января 2012

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

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

Затем вы устанавливаете аргумент ядра с value из NULL и 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, мне это кажется намного более уродливым.

1 голос
/ 01 марта 2018

Вы также можете объявить свои массивы следующим образом: 

__local float LP[LENGTH];

И передать ДЛИНУ как определение в вашей компиляции ядра. 

int lp_size = 128; // this is an example; could be dynamically calculated
char compileArgs[64];
sprintf(compileArgs, "-DLENGTH=%d", lp_size);
clBuildProgram(program, 0, NULL, compileArgs, NULL, NULL);
1 голос
/ 04 августа 2015

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

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

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

...