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