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

StackOverflow https://stackoverflow.com/questions/2541929

  •  23-09-2019
  •  | 
  •  

Вопрос

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

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

Это было полезно?

Решение

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

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

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

Другие советы

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

__local float localBuffer[1024];

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

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

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

Лицензировано под: CC-BY-SA с атрибуция
Не связан с StackOverflow
scroll top