Среди новых возможностей OpenCL 2.0 появилось несколько новых полезных встроенных функций, так называемых функций рабочих групп. Эти встроенные функции предоставляют широко используемые параллельные примитивы, работающие на уровне рабочих групп. В этой статье кратко описываются функции рабочих групп, приводятся данные производительности для устройства OpenCL Intel HD Graphics, а также рассматривается пример использования неоднородных рабочих групп.
Функции рабочих групп включают три классических алгоритма уровня рабочих групп (value broadcast, reduce и scan), а также две встроенные функции, проверяющие логический результат операции, проведенной для всей рабочей группы. Алгоритмы reduce и scan поддерживают операции add, min и max.
Функциональность встроенных функций рабочих групп очевидна из названий.
Важное ограничение, касающееся перечисленных встроенных функций: они действуют только для скалярных типов данных (например, популярные типы int4 и float4 не поддерживаются). Кроме того, не поддерживаются 8-разрядные типы данных, такие как char или uchar.
Функции рабочих групп, что следует из их названия, всегда работают параллельно для целой рабочей группы. Из этого проистекает неявное следствие: любой вызов функции рабочей группы действует в качестве барьера.
Использование функций рабочих групп связано с двумя основными идеями. Во-первых, функции рабочих групп удобны. Гораздо проще использовать одну встроенную функцию вместо написания достаточно крупного фрагмента кода, который бы потребовался для реализации такой же функциональности в OpenCL 1.2. Во-вторых, функции рабочих групп эффективнее с точки зрения производительности, поскольку они используют оптимизацию под оборудование.
Для примера рассмотрим следующую задачу (которая может являться частью какого-либо алгоритма): вычисление сумм префиксов для подчиненных массивов равного размера какого-либо более крупного массива. Итак, нам нужно вычислить сумму префикса для каждого элемента каждого подчиненного массива и сохранить ее в целевой области памяти с той же разметкой. Исходная и целевая компоновки данных показаны на следующей схеме.
Простое ядро OpenCL для выполнения этой задачи может выглядеть так:
Соответствующий код показан ниже.
Такой примитивный подход крайне неэффективен в большинстве случаев (кроме очень маленьких рабочих групп). Очевидно, что внутренний цикл for() выполняет слишком много избыточных операций загрузки и сложения; эту процедуру явно можно оптимизировать. Причем с увеличением размера рабочей группы возрастает и избыточность. Для более эффективного использования аппаратных ресурсов Intel HD Graphics требуется более эффективный алгоритм, например Blelloch. Мы не будем подробно его рассматривать: он замечательно описан в классической статье GPU Gems.
Код OpenCL 1.2 с параллельным сканированием будет выглядеть так.
Как правило, такой код работает эффективнее и образует не столь высокую нагрузку на аппаратные ресурсы, но с некоторыми оговорками.
В этом коде появляются издержки на перемещение данных между локальной и глобальной памятью, а также некоторые запреты. Для достижения действительно высокой эффективности алгоритму требуется достаточно большой размер рабочей группы. При небольших рабочих группах (<16) производительность вряд ли окажется выше, чем у простого цикла.
Кроме того, обратите внимание на усложнение кода и дополнительную логику, предназначенную для исключения конфликтов в общей локальной памяти (например, макрос BANK_OFFSET).
Использование рабочих групп позволяет обойти все упомянутые проблемы. Соответствующий вариант оптимизированного кода OpenCL приведен ниже.
Результаты производительности обоих оптимизированных алгоритмов измерены для достаточно большого объема входных данных (каждая рабочая группа сканирует 65 536 элементов, что, в зависимости от локального размера, соответствует 8192 … 2048 итерациям внешнего цикла).
Как и ожидалось, простой цикл работает гораздо медленнее при увеличении локального размера, а производительность обоих оптимизированных вариантов возрастает.
Если задать оптимальный размер рабочей группы для заданного алгоритма, то сравнение ядер будет таким.
Обратите внимание, что применение work_group_scan_exclusive_add() значительно повышает производительность рабочей группы любого размера и одновременно упрощает код.
Модель выполнения OpenCL включает понятие рабочих групп, которые являются группами отдельных рабочих элементов в NDRange. Если приложение использует OpenCL 1.x, то размеры NDRange должны нацело (без остатка) делиться на размеры рабочих групп. Если вызов clEnqueueNDRangeKernel включает параметры global_size и local_size, которые не делятся нацело, вызов возвратит код ошибки CL_INVALID_WORK_GROUP_SIZE. Если же вызов clEnqueueNDRangeKernel указывает значение NULL для параметра local_size, разрешая выполняемому модулю выбрать размер рабочей группы, то выполняемому модулю потребуется выбрать размер, на который можно нацело разделить глобальные размеры NDRange.
Необходимость выбора такого размера рабочих групп, чтобы на него нацело делился размер NDRange, может вызвать затруднение у разработчиков. Рассмотрим простой алгоритм размытия изображения 3х3. В этом алгоритме каждый выходной пиксель вычисляется как среднее значение для значений входных пикселей в соседней области размером 3х3. Проблема возникает при обработке выходных пикселей, расположенных на рамке изображения, поскольку эти пиксели зависят от пикселей вне границ входного изображения.
В некоторых приложениях входные значения рамок не имеют значения, их можно пропустить. В этом случае размер NDRange совпадает с размером выходного изображения за вычетом области рамки. При этом зачастую получается размер NDRange, который трудно нацело разделить. Например, для применения фильтра 3x3 к изображению 1920x1080 требуется рамка толщиной в один пиксель с каждой стороны. Проще всего это сделать с помощью ядра 1918х1078. Но ни 1918, ни 1078 не делятся нацело на значения, дающие рабочие группы оптимального размера.
В OpenCL 2.0 появилась новая возможность, в которой устранены проблемы, описанные в предыдущем разделе. Речь идет о так называемых неоднородных рабочих группах: выполняемый модуль OpenCL 2.0 может разделить NDRange на рабочие группы неоднородного размера по любому измерению. Если разработчик укажет размер рабочей группы, на который размер NDRange не делится нацело, выполняемый модуль разделит NDRange таким образом, чтобы создать как можно больше рабочих групп с указанным размером, а остальные рабочие группы будут иметь другой размер.
Благодаря этому OpenCL может использовать рабочие группы любого размера для любого размера NDRange, когда разработчик передает значение NULL параметра local_size в clEnqueueNDRangeKernel. В целом использование значения NULL в параметре local_size остается предпочитаемым методом выполнения ядер, если логика вашего приложения не требует какого-либо определенного размера рабочей группы.
Внутри кода ядра встроенная функция get_local_size () возвращает фактический размер рабочей группы, из которой она была вызвана. Если ядру требуется точный размер, указанный для параметра local_size в clEnqueueNDRangeKernel, то встроенная функция get_get_enqueued_local_size () возвращает эти значения.
Чтобы включить использование неоднородных рабочих групп, необходимо скомпилировать ядро с флагом -cl-std=CL2.0, включающим эту и другие возможности OpenCL 2.0. Без использования этого флага компилятор будет использовать версию OpenCL 1.2, даже если устройство поддерживает OpenCL 2.0. Кроме того, неоднородные рабочие группы можно отключить для ядер, скомпилированных для флага -cl-std=CL2.0, с помощью флага -cl-uniform-work-group-size. Это может быть полезно для устаревшего кода ядра до полного перехода на OpenCL 2.0.
Функция неоднородных рабочих групп в OpenCL 2.0 повышает простоту использования OpenCL и может повысить производительность некоторых ядер. Разработчики больше не добавляют код системы и ядра для работы с размерами NDRange, которые не делятся нацело. Код, созданный для использования этой возможности, может эффективно использовать SIMD и выравнивание доступа к памяти: такие преимущества обеспечиваются правильным выбором размера рабочих групп.
В коде учебной программы реализован алгоритм размытия 3х3, описанный выше. Самая интересная часть кода находится в файле main.cpp.
Для каждого варианта в пп. 5—8 результаты вызова get_local_size () и get_get_enqueued_local_size () в каждом из четырех углов NDRange отображаются на экране. Таким образом, мы видим, как происходит разделение NDRange на рабочие группы. Ядро, реализующее алгоритм размытия, хранится в BoxBlur.cl. Оно содержит очень простую реализацию, но не является наиболее эффективным способом применения размытия.
Для сборки и запуска этой учебной программы нужен ПК, соответствующий следующим требованиям:
Учебная программа представляет собой консольное приложение, которое прочитывает входное растровое изображение и записывает выходные растровые изображения для каждой разновидности NDRange, описанной в приведенном выше разделе. Эта учебная программа поддерживает несколько параметров командной строки: -h, -? (отображение текста справки и выход), -i <входной префикс> (префикс входного растрового изображения), -o <выходной префикс> (префикс выходного растрового изображения).
После запуска учебной программы для предоставленного рисунка результат будет таким.
Входное изображение имеет размер 640x480, поэтому размер NDRange в каждом случае составит 638x478. Приведенный выше результат показывает, что запуск ядра OpenCL 1.2 со значением NULL параметра local_size вынуждает использовать нечетные размеры для каждой рабочей группы (1, 239). Размеры рабочих групп, не являющиеся степенями двойки, могут работать очень медленно в некоторых ядрах. Конвейеры SIMD могут простаивать, синхронный доступ к памяти может нарушиться.
Запуск ядра OpenCL 1.2 с указанным размером рабочей группы (16x16) выдает ошибку, поскольку ни 648, ни 478 не делятся нацело на 16.
Запуск ядра OpenCL 2.0 со значением NULL параметра local_size позволяет выполняемому модулю OpenCL разделить NDRange на рабочие группы любого размера. Выше показан результат: видно, что выполняемый модуль продолжает использовать однородный размер рабочих групп точно так же, как для ядра OpenCL 1.2.
Запуск ядра OpenCL 2.0 с заданным размером рабочей группы (16x16) приведет к тому, что размер NDRange будет разделен на неоднородные рабочие группы. Мы видим, что левая верхняя рабочая группа имеет размер 16х16, правая верхняя — 14х16, левая нижняя — 16х14 и правая нижняя — 14х14. Поскольку в большинстве случаев размер рабочей группы составляет 16х16, это ядро будет очень эффективно использовать конвейеры SIMD и доступ к памяти будет очень быстрым.
Полные версии статей на сайте IDZ:
Оригиналы статей на английском:
Описание функций рабочих групп
Функции рабочих групп включают три классических алгоритма уровня рабочих групп (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.
Функции рабочих групп, что следует из их названия, всегда работают параллельно для целой рабочей группы. Из этого проистекает неявное следствие: любой вызов функции рабочей группы действует в качестве барьера.
Использование функций рабочих групп связано с двумя основными идеями. Во-первых, функции рабочих групп удобны. Гораздо проще использовать одну встроенную функцию вместо написания достаточно крупного фрагмента кода, который бы потребовался для реализации такой же функциональности в OpenCL 1.2. Во-вторых, функции рабочих групп эффективнее с точки зрения производительности, поскольку они используют оптимизацию под оборудование.
Для примера рассмотрим следующую задачу (которая может являться частью какого-либо алгоритма): вычисление сумм префиксов для подчиненных массивов равного размера какого-либо более крупного массива. Итак, нам нужно вычислить сумму префикса для каждого элемента каждого подчиненного массива и сохранить ее в целевой области памяти с той же разметкой. Исходная и целевая компоновки данных показаны на следующей схеме.
Простое ядро OpenCL для выполнения этой задачи может выглядеть так:
- каждый массив (строка на иллюстрации) будет обрабатываться одной рабочей группой;
- для каждого рабочего элемента сканирование выполняется с помощью простого цикла for() для предшествующих элементов, затем добавляется совокупное значение префикса, а затем результат сохраняется в месте назначения;
- если размер рабочей группы меньше входного массива, то исходный и конечный индексы сдвигаются на размер рабочей группы, совокупный префикс обновляется и этот процесс повторяется до конца исходной строки.
Соответствующий код показан ниже.
Код
__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.
Код 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 включает понятие рабочих групп, которые являются группами отдельных рабочих элементов в NDRange. Если приложение использует OpenCL 1.x, то размеры NDRange должны нацело (без остатка) делиться на размеры рабочих групп. Если вызов clEnqueueNDRangeKernel включает параметры global_size и local_size, которые не делятся нацело, вызов возвратит код ошибки CL_INVALID_WORK_GROUP_SIZE. Если же вызов clEnqueueNDRangeKernel указывает значение NULL для параметра local_size, разрешая выполняемому модулю выбрать размер рабочей группы, то выполняемому модулю потребуется выбрать размер, на который можно нацело разделить глобальные размеры NDRange.
Необходимость выбора такого размера рабочих групп, чтобы на него нацело делился размер NDRange, может вызвать затруднение у разработчиков. Рассмотрим простой алгоритм размытия изображения 3х3. В этом алгоритме каждый выходной пиксель вычисляется как среднее значение для значений входных пикселей в соседней области размером 3х3. Проблема возникает при обработке выходных пикселей, расположенных на рамке изображения, поскольку эти пиксели зависят от пикселей вне границ входного изображения.
В некоторых приложениях входные значения рамок не имеют значения, их можно пропустить. В этом случае размер NDRange совпадает с размером выходного изображения за вычетом области рамки. При этом зачастую получается размер NDRange, который трудно нацело разделить. Например, для применения фильтра 3x3 к изображению 1920x1080 требуется рамка толщиной в один пиксель с каждой стороны. Проще всего это сделать с помощью ядра 1918х1078. Но ни 1918, ни 1078 не делятся нацело на значения, дающие рабочие группы оптимального размера.
В OpenCL 2.0 появилась новая возможность, в которой устранены проблемы, описанные в предыдущем разделе. Речь идет о так называемых неоднородных рабочих группах: выполняемый модуль OpenCL 2.0 может разделить NDRange на рабочие группы неоднородного размера по любому измерению. Если разработчик укажет размер рабочей группы, на который размер NDRange не делится нацело, выполняемый модуль разделит NDRange таким образом, чтобы создать как можно больше рабочих групп с указанным размером, а остальные рабочие группы будут иметь другой размер.
Благодаря этому OpenCL может использовать рабочие группы любого размера для любого размера NDRange, когда разработчик передает значение NULL параметра local_size в clEnqueueNDRangeKernel. В целом использование значения NULL в параметре local_size остается предпочитаемым методом выполнения ядер, если логика вашего приложения не требует какого-либо определенного размера рабочей группы.
Внутри кода ядра встроенная функция get_local_size () возвращает фактический размер рабочей группы, из которой она была вызвана. Если ядру требуется точный размер, указанный для параметра local_size в clEnqueueNDRangeKernel, то встроенная функция get_get_enqueued_local_size () возвращает эти значения.
Чтобы включить использование неоднородных рабочих групп, необходимо скомпилировать ядро с флагом -cl-std=CL2.0, включающим эту и другие возможности OpenCL 2.0. Без использования этого флага компилятор будет использовать версию OpenCL 1.2, даже если устройство поддерживает OpenCL 2.0. Кроме того, неоднородные рабочие группы можно отключить для ядер, скомпилированных для флага -cl-std=CL2.0, с помощью флага -cl-uniform-work-group-size. Это может быть полезно для устаревшего кода ядра до полного перехода на OpenCL 2.0.
Функция неоднородных рабочих групп в OpenCL 2.0 повышает простоту использования OpenCL и может повысить производительность некоторых ядер. Разработчики больше не добавляют код системы и ядра для работы с размерами NDRange, которые не делятся нацело. Код, созданный для использования этой возможности, может эффективно использовать SIMD и выравнивание доступа к памяти: такие преимущества обеспечиваются правильным выбором размера рабочих групп.
В коде учебной программы реализован алгоритм размытия 3х3, описанный выше. Самая интересная часть кода находится в файле main.cpp.
Код
//1. Загрузка входного растрового файла.
//2. Сборка ядра OpenCL C с помощью параметров OpenCL 1.2.
// Get the box blur kernel compiled using OpenCL 1.2 (which is the
// default compilation, even on an OpenCL 2.0 device). This allows
// the code to show the pre-OpenCL 2.0 behavior.
cl::Kernel kernel_1_2 = GetKernel(device, context);
//3. Сборка ядра OpenCL C с помощью параметров OpenCL 2.0 (обратите внимание на передачу параметров сборки с флагом OpenCL 2.0).
// Get the box blur kernel compiled using OpenCL 2.0. OpenCL 2.0
// is required in order to use the non-uniform work-groups feature.
kernel_2_0 = GetKernel(device, context, "-cl-std=CL2.0");
//4. Задание глобального размера, который используется для всех разновидностей запущенных ядер.
// Set the size of the global NDRange, to be used in all NDRange cases.
// Since this is a box blur, we use a global size that is two elements
// smaller in each dimension. This creates a range which often doesn't
// divide nicely by local work sizes we might commonly pick for running
// kernels.
cl::NDRange global_size = cl::NDRange(input.get_width() - 2,
input.get_height() - 2);
//5. Размытие изображения с помощью версии ядра, скомпилированной для OpenCL 1.2, при значении параметра local_size имеет значение NULL.
// Blur the image with a NULL local range using the OpenCL 1.2 compiled
// kernel.
cout << "Compiled with OpenCL 1.2 and using a NULL local size:"
<< end1 << end1;
output = RunBlurKernel(context, queue, kernel_1_2, global_size,
cl::NullRange, input, true);
//6. Размытие изображения с помощью версии ядра, скомпилированной для OpenCL 1.2, при значении параметра local_size 16x16.
// Blur the image with an even local range using the OpenCL 1.2
// compiled kernel. This won't work, even if we are running on an
// OpenCL 2.0 implementation. The kernel has to be explicitly compiled
// with OpenCL 2.0 compilation enabled in the compiler switches.
try
{
cout << "Compiled with OpenCL 1.2 and using an even local size:"
<< end1 << end1;
output = RunBlurKernel(context, queue, kernel_1_2,
global_size, cl::NDRange(16, 16), input,
true);
cout << end1;
output.Write(output_files[1]);
}
catch (...)
{
cout << "Trying to launch a non-uniform workgroup with a kernel "
"compiled using" << end1 <<
"OpenCL 1.2 failed (as expected.)" << end1 << end1;
}
//7. Размытие изображения с помощью версии ядра, скомпилированной для OpenCL 2.0, при значении параметра local_size NULL.
// Blur the image with a NULL local range using the OpenCL 2.0
// compiled kernel.
cout << "Compiled with OpenCL 2.0 and using a NULL local size:"
<< end1 << end1;
output = RunBlurKernel(context, queue, kernel_2_0, global_size,
cl::NullRange, input, true);
//8. Размытие изображения с помощью версии ядра, скомпилированной для OpenCL 2.0, при значении параметра local_size 16x16.
// Blur the image with an even local range using the OpenCL 2.0
// compiled kernel. This will only work on an OpenCL 2.0 device
// and compiler.
cout << "Compiled with OpenCL 2.0 and using an even local size:"
<< end1 << end1;
output = RunBlurKernel(context, queue, kernel_2_0,
global_size, cl::NDRange(16, 16), input,
true);
//9. Запись выходных файлов, созданных в пп. 2—5.
Для каждого варианта в пп. 5—8 результаты вызова get_local_size () и get_get_enqueued_local_size () в каждом из четырех углов NDRange отображаются на экране. Таким образом, мы видим, как происходит разделение NDRange на рабочие группы. Ядро, реализующее алгоритм размытия, хранится в BoxBlur.cl. Оно содержит очень простую реализацию, но не является наиболее эффективным способом применения размытия.
Для сборки и запуска этой учебной программы нужен ПК, соответствующий следующим требованиям:
- Процессор серии Intel® Core™ с кодовым названием Broadwell.
- Microsoft Windows* 8 или 8.1.
- Intel® SDK для приложений OpenCL™ версии 2014 R2 или более поздней.
- Microsoft Visual Studio* 2012 или более поздней версии.
Учебная программа представляет собой консольное приложение, которое прочитывает входное растровое изображение и записывает выходные растровые изображения для каждой разновидности NDRange, описанной в приведенном выше разделе. Эта учебная программа поддерживает несколько параметров командной строки: -h, -? (отображение текста справки и выход), -i <входной префикс> (префикс входного растрового изображения), -o <выходной префикс> (префикс выходного растрового изображения).
После запуска учебной программы для предоставленного рисунка результат будет таким.
Скрытый текст
Input file: input.bmp
Output files: output_0.bmp, output_1.bmp, output_2.bmp, output_3.bmp
Device: Intel(R) HD Graphics 5500
Vendor: Intel(R) Corporation
Compiled with OpenCL 1.2 and using a NULL local size:
Work Item get_global_id() get_local_size() get_enqueued_local_size()
-------------------------------------------------------------------------
Top left (0,0) (1,239) undefined
Top right (637,0) (1,239) undefined
Bottom left (0,477) (1,239) undefined
Bottom right (637,477) (1,239) undefined
Compiled with OpenCL 1.2 and using an even local size:
Trying to launch a non-uniform workgroup with a kernel compiled using OpenCL 1.2 failed (as expected.)
Compiled with OpenCL 2.0 and using a NULL local size:
Work Item get_global_id() get_local_size() get_enqueued_local_size()
Top left (0,0) (1,239) (1,239)
Top right (637, 0) (1,239) (1,239)
Bottom left (0,477) (1,239) (1,239)
Bottom right (637,477) (1,239) (1,239)
Compiled with OpenCL 2.0 and using an even local size:
Work Item get_global_id() get_local_size() get_enqueued_local_size()
Top left (0,0) (16,16) (16,16)
Top right (637,0) (14,16) (16,16)
Bottom left (0,477) (16,14) (16,16)
Bottom right (637,477) (14,14) (16,16)
Done!
Входное изображение имеет размер 640x480, поэтому размер NDRange в каждом случае составит 638x478. Приведенный выше результат показывает, что запуск ядра OpenCL 1.2 со значением NULL параметра local_size вынуждает использовать нечетные размеры для каждой рабочей группы (1, 239). Размеры рабочих групп, не являющиеся степенями двойки, могут работать очень медленно в некоторых ядрах. Конвейеры SIMD могут простаивать, синхронный доступ к памяти может нарушиться.
Запуск ядра OpenCL 1.2 с указанным размером рабочей группы (16x16) выдает ошибку, поскольку ни 648, ни 478 не делятся нацело на 16.
Запуск ядра OpenCL 2.0 со значением NULL параметра local_size позволяет выполняемому модулю OpenCL разделить NDRange на рабочие группы любого размера. Выше показан результат: видно, что выполняемый модуль продолжает использовать однородный размер рабочих групп точно так же, как для ядра OpenCL 1.2.
Запуск ядра OpenCL 2.0 с заданным размером рабочей группы (16x16) приведет к тому, что размер NDRange будет разделен на неоднородные рабочие группы. Мы видим, что левая верхняя рабочая группа имеет размер 16х16, правая верхняя — 14х16, левая нижняя — 16х14 и правая нижняя — 14х14. Поскольку в большинстве случаев размер рабочей группы составляет 16х16, это ядро будет очень эффективно использовать конвейеры SIMD и доступ к памяти будет очень быстрым.
Полные версии статей на сайте IDZ:
Оригиналы статей на английском:
