Использование функций рабочих групп в OpenCL™ 2.0

Download PDF      Download Code

Для новичков в области программирования на OpenCL раздел в спецификации OpenCL 2.0, касающийся этой темы, может показаться слишком академичным и малопонятным. Коротко говоря, функции рабочих групп включают три классических алгоритма уровня рабочих групп (value broadcast, reduce и scan), а также две встроенные функции, проверяющие логический результат операции, проведенной для всей рабочей группы. Алгоритмы reduce и scan поддерживают операции add, min и max.

Описание функций рабочих групп

Для новичков в области программирования на OpenCL раздел в спецификации OpenCL 2.0, касающийся этой темы, может показаться слишком академичным и малопонятным. Коротко говоря, функции рабочих групп включают три классических алгоритма уровня рабочих групп (value broadcast, reduce и scan), а также две встроенные функции, проверяющие логический результат операции, проведенной для всей рабочей группы. Алгоритмы reduce и scan поддерживают операции add, min и max.

Функциональность встроенных функций рабочих групп очевидна из названий:

  • work_group_broadcast() распространяет значение выбранного рабочего элемента на все элементы рабочей группы
  • work_group_reduce() вычисляет значения sum, min или maxдля всех элементов рабочей группы, а затем распространяет полученное значение на все элементы рабочей группы
  • work_group_scan()  вычисляет значения sum, min или max для всех предшествующих рабочих элементов (с возможным включением текущих)
  • work_group_all() возвращает логическое И для одинакового логического выражения, вычисленного для каждого рабочего элемента
  • work_group_any() действует аналогично work_group_all(), но использует логическое ИЛИ

Важное ограничение, касающееся перечисленных встроенных функций: они действуют только для скалярных типов данных (например, популярные типы int4 и float4 не поддерживаются). Кроме того, не поддерживаются 8-разрядные типы данных, такие как char или uchar.

Полное описание см. в главе 6.13.15 спецификации OpenCL 2.0 C.

Функции рабочих групп, что следует из их названия, всегда работают параллельно для целой рабочей группы. Из этого проистекает неявное следствие: любой вызов функции рабочей группы действует в качестве барьера.

Использование функций рабочих групп связано с двумя основными идеями. Во-первых, функции рабочих групп удобны. Гораздо проще использовать одну встроенную функцию вместо написания достаточно крупного фрагмента кода, который бы потребовался для реализации такой же функциональности в OpenCL 1.2. Во-вторых, функции рабочих групп эффективнее с точки зрения производительности, поскольку они используют оптимизацию под оборудование.

Пример

Для примера рассмотрим следующую задачу (которая может являться частью какого-либо алгоритма): вычисление сумм префиксов для подчиненных массивов равного размера какого-либо более крупного массива. Итак, нам нужно вычислить сумму префикса для каждого элемента каждого подчиненного массива и сохранить ее в целевой области памяти с той же разметкой. Исходная и целевая компоновки данных показаны на следующей схеме:

Простое ядро OpenCL для выполнения этой задачи может выглядеть так:

  • каждый массив (строка на иллюстрации) будет обрабатываться одной рабочей группой
  • для каждого рабочего элемента сканирование выполняется с помощью простого цикла for() для предшествующих элементов, затем добавляется совокупное значение префикса, а затем результат сохраняется в месте назначения
  • если размер рабочей группы меньше входного массива, то исходный и конечный индексы сдвигаются на размер рабочей группы, совокупный префикс обновляется и этот процесс повторяется до конца исходной строки

Соответствующий код показан ниже:

__kernel void Calc_wg_offsets_naive(

__kernel void Calc_wg_offsets_naive(
                            __global const uint* gHistArray,
                            __global uint* gPrefixsumArray,
                            uint bin_size
                            )
{
    uint lid = get_local_id(0);
    uint binId = get_group_id(0);

    //calculate source/destination offset for workgroup
    uint group_offset = binId * bin_size;
    local uint maxval;

    //initialize cumulative prefix
    if( lid == 0 )  maxval = 0;
    barrier(CLK_LOCAL_MEM_FENCE);

    do
    {
        //perform a scan for every workitem
        uint prefix_sum=0;
        for(int i=0; i<lid; i++)
            prefix_sum += gHistArray[group_offset + i];
        
        //store result
        gPrefixsumArray[group_offset + lid] = prefix_sum + maxval;
        prefix_sum += gHistArray[group_offset + lid];

        //update group offset and cumulative prefix
        if( lid == get_local_size(0)-1 )  maxval += prefix_sum;
        barrier(CLK_LOCAL_MEM_FENCE);

        group_offset += get_local_size(0);
    }
    while(group_offset < (binId+1) * bin_size);
}

Такой примитивный подход крайне неэффективен в большинстве случаев (кроме очень маленьких рабочих групп). Очевидно, что внутренний цикл for() выполняет слишком много избыточных операций загрузки и сложения; эту процедуру явно можно оптимизировать. Причем с увеличением размера рабочей группы возрастает и избыточность. Для более эффективного использования аппаратных ресурсов Intel HD Graphics требуется более эффективный алгоритм, например Blelloch. Мы не будем подробно его рассматривать: он замечательно описан в классической статье GPU Gems - http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html.

Код OpenCL 1.2 с параллельным сканированием будет выглядеть так:

#define WARP_SHIFT 4
#define GRP_SHIFT 8
#define BANK_OFFSET(n)     ((n) >> WARP_SHIFT + (n) >> GRP_SHIFT)

__kernel void Calc_wg_offsets_Blelloch(__global const uint* gHistArray,
                              __global uint* gPrefixsumArray,
                              uint bin_size
                              ,__local uint* temp
                            )
{
    int lid = get_local_id(0);
    uint binId = get_group_id(0);
    int n = get_local_size(0) * 2;

    uint group_offset = binId * bin_size;
    uint maxval = 0;
    do
    {
        // calculate array indices and offsets to avoid SLM bank conflicts
        int ai = lid;
        int bi = lid + (n>>1);  
        int bankOffsetA = BANK_OFFSET(ai);
        int bankOffsetB = BANK_OFFSET(bi);

        // load input into local memory
        temp[ai + bankOffsetA] = gHistArray[group_offset + ai];
        temp[bi + bankOffsetB] = gHistArray[group_offset + bi];

        // parallel prefix sum up sweep phase
        int offset = 1;
        for (int d = n>>1; d > 0; d >>= 1)
        {   
            barrier(CLK_LOCAL_MEM_FENCE);
            if (lid < d)
            {  
                int ai = offset * (2*lid + 1)-1;
                int bi = offset * (2*lid + 2)-1;
                ai += BANK_OFFSET(ai);
                bi += BANK_OFFSET(bi);
                temp[bi] += temp[ai];  
            }  
            offset <<= 1; 
        }

        // clear the last element
        if (lid == 0)
        {
            temp[n - 1 + BANK_OFFSET(n - 1)] = 0;
        }

        // down sweep phase
        for (int d = 1; d < n; d <<= 1)
        {  
            offset >>= 1;  
            barrier(CLK_LOCAL_MEM_FENCE);

            if (lid < d)
            {
                int ai = offset * (2*lid + 1)-1;
                int bi = offset * (2*lid + 2)-1;
                ai += BANK_OFFSET(ai);
                bi += BANK_OFFSET(bi);
         
                uint t = temp[ai];  
                temp[ai] = temp[bi];  
                temp[bi] += t;   
            }
        }  
        barrier(CLK_LOCAL_MEM_FENCE);

        //output scan result to global memory
        gPrefixsumArray[group_offset + ai] = temp[ai + bankOffsetA] + maxval;
        gPrefixsumArray[group_offset + bi] = temp[bi + bankOffsetB] + maxval;

        //update cumulative prefix sum and shift offset for next iteration
        maxval += temp[n - 1 + BANK_OFFSET(n - 1)] + gHistArray[group_offset + n - 1];
        group_offset += n;
    }
    while(group_offset < (binId+1) * bin_size);
}

Как правило, такой код работает эффективнее и образует не столь высокую нагрузку на аппаратные ресурсы, но с некоторыми оговорками.

В этом коде появляются издержки на перемещение данных между локальной и глобальной памятью, а также некоторые запреты. Для достижения действительно высокой эффективности алгоритму требуется достаточно большой размер рабочей группы. При небольших рабочих группах (<16) производительность вряд ли окажется выше, чем у простого цикла.

Кроме того, обратите внимание на усложнение кода и дополнительную логику, предназначенную для исключения конфликтов в общей локальной памяти (например, макрос BANK_OFFSET).

Использование рабочих групп позволяет обойти все упомянутые проблемы. Соответствующий вариант оптимизированного кода OpenCL приведен ниже:

__kernel void Calc_wg_offsets_wgf(
                            __global const uint* gHistArray,
                            __global uint* gPrefixsumArray,
                            uint bin_size
                            )
{
    uint lid = get_local_id(0);
    uint binId = get_group_id(0);

    uint group_offset = binId * bin_size;
    uint maxval = 0;

    do
    {
        uint binValue = gHistArray[group_offset + lid];
        uint prefix_sum = work_group_scan_exclusive_add( binValue );
        gPrefixsumArray[group_offset + lid] = prefix_sum + maxval;

        maxval += work_group_broadcast( prefix_sum + binValue, get_local_size(0)-1 );

        group_offset += get_local_size(0);
    }
    while(group_offset < (binId+1) * bin_size);
}

Результаты производительности обоих оптимизированных алгоритмов измерены для достаточно большого объема входных данных (каждая рабочая группа сканирует 65 536 элементов, что, в зависимости от локального размера, соответствует 8192 … 2048 итерациям внешнего цикла).

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

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

Обратите внимание, что применение work_group_scan_exclusive_add() значительно повышает производительность рабочей группы любого размера и одновременно упрощает код.

Заключение

Добавление функций рабочих групп — важное усовершенствование в OpenCL 2.0, значительно упрощающее работу разработчиков, применяющих OpenCL. Правильное использование этой новой функции поможет снизить трудоемкость разработки сложных и высокопроизводительных приложений OpenCL.

Для получения подробной информации о возможностях оптимизации компилятора обратитесь к нашему Уведомлению об оптимизации.
Возможность комментирования русскоязычного контента была отключена. Узнать подробнее.