Как использовать локальную память в 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 локальная память предназначена для обмена данными между всеми рабочими элементами в рабочей группе. И обычно требуется выполнить барьерный вызов перед использованием данных локальной памяти (например, один рабочий элемент хочет прочитать данные локальной памяти, записанные другими рабочими элементами). Барьер дорогостоящий в оборудовании. Имейте в виду, локальная память должна использоваться для повторного чтения / записи данных. Банковский конфликт следует избегать в максимально возможной степени.

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