Как использовать локальную память в 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];
}

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


opencl
person splicer    schedule 29.03.2010    source источник

Ответы (3)


arrow_upward
32
arrow_downward

Ознакомьтесь с примерами в 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];
    }
}
person Tom    schedule 02.04.2010
comment
Я прочитал вводный материал NVIDIA, но все еще считаю примеры слишком сложными. Я ищу сверхпростой одномерный пример использования локальной памяти, чтобы намочить ноги. - person splicer; 02.04.2010
comment
Спасибо за добавление кода в ваше последнее изменение! Я не могу заставить ваше ядро ​​работать… Как мне использовать clSetKernelArg () для временного хранения? Нужно ли мне использовать clCreateBuffer () для временного хранения? Кроме того, в вашем ядре есть несколько опечаток: temp * temp должно быть temp [ltid] * temp [ltid], а закрывающая скобка должна быть вставлена ​​перед последней строкой. - person splicer; 04.04.2010
comment
Запустив CPU под Snow Leopard, я попробовал clSetKernelArg (kernel, 2, sizeof (cl_float), NULL); но он вылетает. Любые идеи? - person splicer; 04.04.2010
comment
Исправил опечатки - служит мне для набора текста на айподе. Ваш clSetKernelArg не выделяет достаточно памяти, вам нужно место для одного cl_float на поток (вы выделили только одно число с плавающей запятой). Попробуйте: clSetKernelArg(kernel, 2, sizeof(cl_float) * local_work_size[0], NULL);, где local_work_size[0] - размер рабочей группы в измерении 0. - person Tom; 04.04.2010
comment
Спасибо! Похоже, вам не хватает точки с запятой в строке 11. На процессоре get_local_size (0) возвращает мне 1, так что не должно ли сработать мое использование clSetKernelArg? Это ошибка в реализации Apple? - person splicer; 06.04.2010
comment
Обратите внимание, что вы можете объявить переменные как локальные с помощью квалификатора __local. Например, вы можете сделать __local float values[GROUP_SIZE];, а затем заставить каждый поток писать values[get_local_id(0)] = .... Доступ к локальной памяти не требуется через указатель, переданный в ядро. - person ; 26.05.2010
comment
@EdwardLuong: на самом деле, ваше предложение не работает, по крайней мере, на моем Macbook AMD Radeon HD 6750M. Хотелось бы, чтобы это могло сделать мой класс Objective-C OpenCL намного проще. К вашему сведению, при реализации этого подхода я получаю код ошибки = -11; clBuildProgram (CL_BUILD_PROGRAM_FAILURE). Спасибо, в любом случае. - person Bruce Dean; 05.06.2012
comment
Помните, что локальное хранилище видно только той же рабочей группе. Таким образом, размер вашей рабочей группы должен быть ›1, чтобы ее можно было использовать в нескольких потоках. - person Tim Child; 02.08.2012
comment
Я думал, нужно использовать барьер локальной памяти после записи в локальную? - person mike; 19.04.2015
comment
@mike: в приведенном выше примере каждый поток использует уникальную часть локальной памяти в качестве своего личного блокнота, поэтому барьер не требуется. Однако, если потоки собираются обмениваться данными через локальную память, то есть читать данные, записанные другим потоком, тогда, да, потребуется барьер. Учитывая, что это ответ новичку, я должен был упомянуть об этом, поскольку это очевидная ошибка в будущем. - person Tom; 02.05.2015
comment
Будет ли копирование данных в локальный буфер, как указано выше, улучшить скорость ядра? Я думаю, что оборудование могло загрузить данные в кеш L1. - person fyquah95; 04.10.2015

arrow_upward
29
arrow_downward

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

__local float localBuffer[1024];

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

person Rick-Rainer Ludwig    schedule 13.06.2011
comment
Это правда, но было бы гораздо полезнее, если бы вам не нужно было знать размер массива во время выполнения. Это желательно при инкапсуляции функциональности OpenCL внутри класса объекта. Например, см. Комментарий Эдварда Луонга выше; было бы здорово, если бы его предложение могло сработать (похоже, не работает для моего оборудования). Спасибо. - person Bruce Dean; 05.06.2012

arrow_upward
5
arrow_downward

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

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

person Hunter Wang    schedule 03.07.2013