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

Question

Я недавно играл с 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];
}

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

Answer 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];
    }
}

Answer 2

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

__local float localBuffer[1024];

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

Answer 3

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

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