OpenCL: какой тип памяти использовать?

У меня есть какое-то фильтрующее ядро, что-то вроде этого:

__kernel void filterKernel (__global float4 *filter, __global float4* in_array, __global float4* out_array)
{
...
out_array[tid] = in_array[tid] * filter[fid];
...
}

  • kernel filterKernel вызывается несколько раз (около 1000 раз).

  • Переменная filter представляет собой массив чисел с плавающей запятой, которые никогда не меняют своих значений (остаются одинаковыми для всех рабочих групп и для всех вызовов ядра).

  • in_array содержит 32768 чисел с плавающей запятой.

Как лучше всего объявить эту переменную filter? __постоянный? __местный? Может быть, разместить «const» здесь и там? Что больше всего помогает компилятору? Что делает код самым быстрым?


memory local constants opencl
person Frizz    schedule 07.01.2011    source источник
comment
На ваш вопрос ответили? Если да, то, пожалуйста, примите ответ.   -  person koan    schedule 06.08.2011

Ответы (3)


arrow_upward
3
arrow_downward

Вы должны использовать постоянное адресное пространство (__constant), так как большинство графических процессоров имеют специальные кеши для постоянной памяти. Единственная проблема заключается в том, что постоянная память имеет небольшой размер (порядка 16-64 КБ).

person Dr. Snoopy    schedule 07.01.2011
comment
В случае, если я использую постоянную память: следует ли добавить определение (например, ___constant float4* filter = {1.0, 2.0, 3.0, ...}; ) в мой файл .cl или сохранить filter в моем интерфейс функции и заполнить массив на ЦП, прежде чем я вызову функцию? Есть ли вообще разница? - person Frizz; 07.01.2011
comment
@Frizz: заполните его из ЦП, используя clSetKernelArg. - person Dr. Snoopy; 07.01.2011
comment
Это намного медленнее, чем заполнение массива __constant в файле .cl. Смотрите мои комментарии ниже. - person Frizz; 10.01.2011

arrow_upward
1
arrow_downward

__local будет неправильным, так как вы не можете ничего инициализировать. Вы, вероятно, захотите использовать __constant, при условии, что он подходит.

person arsenm    schedule 07.01.2011

arrow_upward
0
arrow_downward

Если он не слишком велик, попробуйте определить свой фильтр глобально внутри файла .cl.
Там вы можете попробовать разместить его либо в __constant, либо в __local и сравните, какой из них быстрее. Но не все SDK поддерживают глобальные переменные в адресном пространстве __local (я смотрю на вас ATI).

Если вы все же хотите передать фильтр в качестве аргумента ядра, рассмотрите возможность вызова его SetKernelArg(0, ...) только один раз. Нет необходимости вызывать SetKernelArg() 1000 раз, пока значение или индекс аргумента ядра не меняется. Хотя это может и не оказать заметного влияния на производительность, оно все же чище.

person vobject    schedule 07.01.2011
comment
Из всех различных методов, которые я пробовал до сих пор (постоянная память, LDS, кеш текстур), определение его в файле .cl как __constant на сегодняшний день (!) является самым быстрым. Удивительно, но есть огромная (!) разница между A) передачей массива __constant в ядро ​​через SetkernelArg() перед его вызовом и B) определением массива __constant в самом файле .cl. К сожалению, метод B) работает только для ATI/AMD. Реализация OpenCL от Nvidia содержит ошибки (я уже сообщал об этой ошибке, она будет исправлена ​​в CUDA 3.3). - person Frizz; 10.01.2011
comment
Да, я второй. CUDA SDK очень глючный. Я попытался реализовать функцию byte_swap() внутри ядра, используя операции побитового сдвига. Компилятор OpenCL распознает назначение функции и пытается сопоставить ее с bswap32 соответственно. Инструкция bswap64 - которую невозможно найти или решить... - person vobject; 10.01.2011
comment
Frizz: Вы проверяли влияние установки аргумента ядра для буфера __constant только один раз по сравнению с каждой постановкой в ​​очередь ядра? - person dietr; 11.01.2011
comment
Нет, не проверял, есть ли разница. Сделаю. ... Обычно я ТОЛЬКО проверяю, сколько времени требуется для завершения вызова clEnqueueNDRange(). - person Frizz; 12.01.2011
comment
Проверено: это не имеет никакого значения (установка аргумента ядра __constant buffer только один раз, а не перед каждым вызовом ядра). По крайней мере, не на моем оборудовании (GT240). - person Frizz; 12.01.2011