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

Question

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

__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 раз).

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

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

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

Answer 1

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

Answer 2

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

Answer 3

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

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