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

Недавно я играл с OpenCL, и я могу писать простые ядра, которые используют только глобальную память. Теперь я хотел бы начать использовать локальную память, но я не могу понять, как использовать get_local_size() и get_local_id() для вычисления одного "куска" вывода за раз.

Например, скажем, я хотел преобразовать ядро Apple OpenCL Hello World в нечто, использующее локальную память. Как бы вы это сделали? Здесь исходный источник ядра:

__kernel square(
    __global float *input,
    __global float *output,
    const unsigned int count)
{
    int i = get_global_id(0);
    if (i < count)
        output[i] = input[i] * input[i];
}

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

Ответ 1

Проверьте образцы в NVIDIA или AMD SDK, они должны указать вам в правильном направлении. Например, матричная транспозиция будет использовать локальную память.

Используя ваше квадратичное ядро, вы можете создавать данные в промежуточном буфере. Не забудьте передать дополнительный параметр.

__kernel square(
    __global float *input,
    __global float *output,
    __local float *temp,
    const unsigned int count)
{
    int gtid = get_global_id(0);
    int ltid = get_local_id(0);
    if (gtid < count)
    {
        temp[ltid] = input[gtid];
        // if the threads were reading data from other threads, then we would
        // want a barrier here to ensure the write completes before the read
        output[gtid] =  temp[ltid] * temp[ltid];
    }
}

Ответ 2

Есть еще одна возможность сделать это, если размер локальной памяти постоянный. Без использования указателя в списке параметров ядер локальный буфер можно объявить в ядре, просто объявив его __local:

__local float localBuffer[1024];

Это удаляет код из-за меньших вызовов clSetKernelArg.

Ответ 3

В OpenCL локальная память предназначена для обмена данными по всем рабочим элементам рабочей группы. И обычно требуется сделать вызов барьера, прежде чем данные локальной памяти могут быть использованы (например, один рабочий элемент хочет прочитать данные локальной памяти, написанные другими рабочими элементами). Барьер является дорогостоящим аппаратным обеспечением. Имейте в виду, что локальная память должна использоваться для повторного чтения/записи данных. Банковский конфликт следует избегать как можно больше.

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