Введение в OpenCL
Эта статья посвящена основам программирования на OpenCl. OpenCl -это язык программирования на GPU/CPU, по своей структуре близкий к стандарту c99. Его развитием занимается Khronos Group, где на их сайте доступна полная документация. Во избежание полемики на тему «ну это же всё тривиально, достаточно покопаться в инете» сразу оговорюсь: в рунете информация на эту тематику практически полностью отсутствует, а в западном инете доступна весьма в разрозненном состоянии на десятке сайтов. Здесь будет приведена некоторая компиляция базовых принципов, максимально упрощающая начинающему программисту жизнь, а так же позволяющая с самого первого проекта максимально задействовать вычислительные мощности видеокарты. Людям написавшим 2-3 серьёзных программы на OpenCl это будет уже неинтересно. Статья в некотором смысле является продолжением моей прошлой статьи.
Компилятор
В первую очередь вопрос: где писать сам код. Насколько мне известно под .NET пока что нет никакой свистелки, позволяющей обрабатывать код ядра непосредственно в студии. Поэтому приходиться использовать сторонние редакторы. AMD, nVidia и Intel прилагают их к своим пакетам SDK. Мне почему-то больше нравиться именно Интеловский. Так же, как вариант, есть несколько редакторов, написанных фанатами. Из них мне больше всего нравиться редактор, прилагающийся к OpenCLTemplate. Стоит отметить, что это именно редакторы. Компиляция кода происходит непосредственно перед запуском на GPU/CPU.
Модель памяти устройства
Прежде чем описывать сам язык я дам краткое описание физической модели устройства с которой он взаимодействует. Исполнение команд языка идёт на объектах, называемых «work-item». Каждый «work-item» не зависим от другого и может исполнять код параллельно с остальными. Если же процесс из одного work-item хочет получить данные, используемые или уже обработанные любым другим work-item он может это сделать через общую память. Общая память весьма медленная, зато имеет большой объём. Чтобы ускорить вычисления имеется локальная память. Если вы знакомы с CUDA, то там она называется «разделяемая память». Она значительно быстрее общей, но не любой процесс может получить к ней доступ. К локальной памяти могут обращаться только work-item одной группы. Эти группы называются «Compute Unit» или «Workgroup» (первое название относится к физическому разбиению на уровне железа, а второе к логическому на уровне программы). В зависимости от устройства в каждой из этих групп различное количество work-item (например 240 для NVIDIA GT200 или 256 для Radeon 5700 Series). Количество этих юнитов ограниченно достаточно маленьким числом (30 для NVIDIA GT200 или 9-10 для Radeon 5700 Series). Так же существует сверхбыстрая «private memory» к которой work-item может обращаться единолично.
Драйвера OpenCL устройств автоматизируют старт и работу work-item и workgroup. Например если нам нужно выполнить миллион процессов, а у нас в распоряжении всего тысяча work-item, то драйвера будут автоматически запускать каждый процесс со следующей задачей после его завершения. Понимание физического уровня требуется только для того, чтобы иметь представление о возможностях взаимодействия между процессами и доступа процессов в память.
Базовые особенности
Так как в основе языка лежит практически стандартный с++ я буду рассматривать только те особенности, которые отличают OpenCL от него. Рассмотрим код простейшей программы ядра, который я приводил в прошлой статье. Этот код складывает два вектора, v1 и v2, кладя результат в первый из них.
__kernel void
floatVectorSum(__global float * v1,
__global float * v2)
int i = get_global_id(0);
v1[i] = v1[i] + v2[i];
>* This source code was highlighted with Source Code Highlighter .
Объявление процедур
В первую очередь в глаза бросается загадочный «__kernel «. Этой директивой должна быть помечена любая процедура, которую мы хотим вызвать извне. Если процедура не нужна при работе извне, её можно не отмечать.
Типы памяти
Тип данных «__global » обозначает память, которая выделяется из глобального адресного пространства работающего устройства. Она достаточна медленная, зато вместительная. Для современных видеокарт измеряется гигобайтами. Если вы работаете на процессоре — под global подразумевается оперативная память.
Кроме global есть «__local «. К ней может обращаться только рабочая группа(workgroup). На каждую такую группу выделяется примерно 8 килобайт.
Так же быстрой памятью является «__privat «. Это память к которой имеет доступ только отдельный поток (work-item). Всего на поток выделяется 32 регистра этой памяти.
Остальные типы памяти, которые можно объявлять при создании ядра основаны на типе «__global «. Во-первых, это «__constant «, который может использоваться только для чтения. Во-вторых, это «__read_only», «__write_only» и «__read_write» — структуры, использование которых разрешено только для изображений.
Идентификаторы процессов
После запуска на видеокарте все процессы равнозначны и исполняют равнозначный код. Но, очевидно, нам не нужно многократное повторение одного и того же действия — каждый процесс должен делать свой кусок задачи. Для осознания своего места в окружающем мире служат идентификаторы процессов. Самый простой идентификатор — » get_global_id(0)». В случае приведённого примера он указывает на i номер вектора, который должен сложить этот процесс. Если же мы обрабатываем не одномерный вектор, а двухмерное изображение — нам нужно знать положение процесса по двум осям. Конечно, это значение можно вычислить. Но это лишние операции. Поэтому для удобства при запуске можно указать, что нам нужно пространство двухмерной размерности. Тогда в процессе можно получить оба идентификатора положения: «get_global_id(0)», «get_global_id(1)». Так же можно сделать и для трёхмерного пространства. Часто может потребоваться и размерности пространства в которой мы работаем. Например для изображения практически при любой его обработке нам нужны его ширина и высота. Для получения размерности пространства используется идентификатор «get_global_size(i)». Кроме этого есть идентификаторы процессов внутри рабочей группы — «get_local_id(i)», «get_local_size(i)» и идентификатор самой группы -» get_group_id(i)», «get_num_groups(i)». Большая часть этих соотношений связанна друг с другом: num_groups * local_size = global_size, local_id + group_id * local_size = global_id, global_size % local_size = 0.
Оптимизация расчётов
Разработчики OpenCL и видеокарт понимали, что основная цель их детища — ускорить сложные расчёты. Для этого в язык был добавлен ряд специализированных особенностей, позволяющих при их использовании получить прирост в скорости на математических задачах.
Встроенные вектора
Самой первой особенностью являются вектора и векторная математика. В OpenCl можно объявлять в качестве переменных 2, 4, 8 и 16-мерные вектора. Это делается соответственно: int2, int4, int8, int16. Так же можно объявлять double, byte и все остальные типы. Вектора соответствующей размерности можно складывать/вычитать/делить/умножать, а так же любой вектор можно делить/умножать на число:
uint4 sumall = (uint4)(1,1,1,1);
small += (uint4)(1,1,1,1);
sumall = sumall/2;* This source code was highlighted with Source Code Highlighter .
Кроме того есть ряд функций, оптимизированных под вектора и позволяющих работать непосредственно с ними. К таким функциям относятся функции вычисления расстояния, функции векторного произведения. Например:
float4 dir1 = (float4)(1, 1, 1, 0);
float4 dir2 = (float4)(1, 2, 3, 0);
float4 normal = cross(dir1, dir2);* This source code was highlighted with Source Code Highlighter .
Ещё можно сливать вектора, беря части от одного и другого, а так же склеивать в более большие:
int4 vi0 = (int4) -7 ;
int4 vi1 = (int4) ( 0, 1, 2, 3 ) ;
vi0.lo = vi1.hi; // слияние
int8 v8 = (int8)(vi0.s0123, vi1.s0123); //склейка* This source code was highlighted with Source Code Highlighter .
Простые функции
Следующей особенностью OpenCl является встроенная библиотека функций. Кроме стандартного набора math.lib в OpenCl имеются так называемые native функции. Это функции, основаны непосредственно на использовании некоторых функций видеокарт и на загрублённой математике. Не советуется применять их при сверхточных расчётах, но в случае фильтрации изображений разницу невозможно заметить. К таким функциям, например, относятся: «native_sin», «native_cos», «native_powr». Я не буду приводить более подробное объяснение этих функций, их очень много, да и принципы разные. Если они вам понадобятся — смотрите документацию.
Часто встречающиеся функции
Кроме «простых функций» разработчики создали целый ряд называемый common function. Это функции, часто встречающиеся при обработке изображений. Например: mad(a,b,c) = a*b + c, mix(a,b,c) = a + (b-a)*c. Эти функции выполняются быстрее, чем соответствующие им математические действия.
Пример
На сайте www.cmsoft.com.br есть замечательный пример, показывающий возможности при оптимизации кода средствами native и common функций:
kernel void regularFuncs()
for ( int i=0; i float a=1, b=2, c=3, d=4;
float e = a*b+c;
e = a*b+c*d;
e = sin(a);
e = cos(b);
e = a*b+c*d;
e = sin(a);
e = cos(b);
e = a*b+c*d;
e = sin(a);
e = cos(b);
float4 vec1 = (float4)(1, 2, 3, 0);
float4 vec2 = (float4)(-1, 3, 1, 0);
float4 vec = distance(vec1, vec2);
double x=1, y=2, z=3;
double resp = x*y+z;
>
>
kernel void nativeFuncs()
for ( int i=0; i float a=1, b=2, c=3, d=4;
float e = mad(a,b,c);
e = mad(a,b,c*d);
e = native_sin(a);
e = native_cos(b);
e = mad(a,b,c*d);
e = native_sin(a);
e = native_cos(b);
e = mad(a,b,c*d);
e = native_sin(a);
e = native_cos(b);
float4 vec1 = (float4)(1, 2, 3, 0);
float4 vec2 = (float4)(-1, 3, 1, 0);
float4 vec = fast_distance(vec1, vec2);
double x=1, y=2, z=3;
double resp = mad(x,y,z);
>
>* This source code was highlighted with Source Code Highlighter .
Вторая процедура (использующая оптимизацию) выполняется в 35 раз быстрее.
Разрешения
Стоит отметить, что в OpenCl есть ряд директив, позволяющих включать различные дополнительные функционалы. Причин к этому две. Первая — исторически не все из этих функционалов поддерживались. Вторая — эти функционалы могут влиять на производительность. Обычно функционал включается следующей командой:
#pragma OPENCL EXTENSION extension name : behavior
* This source code was highlighted with Source Code Highlighter .
Для примера. Слудющие команды включают: возможность использования типа byte, двойную точность рассчётов и все математические функции
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#pragma OPENCL EXTENSION cl_khr_fp64 : enable* This source code was highlighted with Source Code Highlighter .
Синхронизация
Барьеры
Часто в вычислениях нужна синхронизация. Это достигается несколькими способами. Первый — барьеры. Барьер это такая команда, на которой остановится процесс пока все прочие процессы или процессы его рабочей группы не достигнут. Приведём два примера:
kernel void localVarExample()
int i = get_global_id(0);
__local int x[10];
x[i] = i;
barrier(CLK_LOCAL_MEM_FENCE);
if (i>0) int y = x[i-1];
>
kernel void globalVarExample()
int i = get_global_id(0);
__global int x[10];
x[i] = i;
barrier(CLK_GLOBAL_MEM_FENCE);
if (i>0) int y = x[i-1];
>* This source code was highlighted with Source Code Highlighter .
В первом примере на команде barrier ожидают все процессы рабочей группы, во втором — все процессы OpenCL устройства.
Стоит отметить особенность этого примера, команды «__local int x[10];» и «__global int x[10];». Они позволяют выделить глобальную переменную в группе процессов и во всех процессах уже во время их исполнения.
Единичные операции
Второй вариант синхронизации между потоками — atomic. Это функции, предотвращающие одновременное обращение к памяти. Перед их использованием нужно включить следующие директивы:
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable* This source code was highlighted with Source Code Highlighter .
Простейший пример работы этой функции:
__kernel void test(global int * num)
atom_inc(&num[0]);
>* This source code was highlighted with Source Code Highlighter .
Если бы вместо «atom_inc(&num[0]);» было написано num++, то результат исполнения программы был непредсказуем, так как все процессы бы одновременно обратились к памяти и считали там одно и то же значение. Всего есть 11 функций еденичных операций: «add, sub, xchg, inc, dec, cmp_xchg, min, max, and, or, xor».
При помощи этих функций не сложно создать семафор.
void GetSemaphor(__global int * semaphor) int occupied = atom_xchg(semaphor, 1);
while (occupied > 0)
occupied = atom_xchg(semaphor, 1);
>
>
void ReleaseSemaphor(__global int * semaphor)
int prevVal = atom_xchg(semaphor, 0);
>
* This source code was highlighted with Source Code Highlighter .
Работа с изображениями
Последней вещью, которую я хочу включить в этот guide является работа с изображениями через OpenCL. Создатели попробовали сделать так, чтобы работа с изображениями требовала минимума мозга пользователя. Это очень приятно. Загрузка изображений возможна в типы image2d_t и image3d_t. Первые — это обычные изображения, вторые — трёхмерные. Так же загружаемое изображение должно быть одного из форматов: » __read_only», » __write_only», «__read_write». Чтение и запись данных из изображения возможны только специальными процедурами: значение = read_imageui(изображение, сэмплер, положение), write_imageui(изображение, положение, значение).
На мой взгляд здесь всё понятно кроме понятия «сэмплер». Сэмплер — это штука, которая будет оптимизировать вашу работу с изображением. У него есть три параметра: «normalized coords», «address mode», «filter mode». Первый имеет два значения: «CLK_NORMALIZED_COORDS_TRUE, CLK_NORMALIZED_COORDS_FALSE». В соответствии с названием он должен показывать, нормализованы ли входные координаты или нет. Второй показывает, что делать в случае, если вы пробуете прочитать координаты из-за пределов границ изображения. Возможные варианты: зеркально продолжить изображение(CLK_ADDRESS_MIRRORED_REPEAT), взять ближайшее граничное значение (CLK_ADDRESS_CLAMP_TO_EDGE), взять базовый цвет (CLK_ADDRESS_CLAMP), ничего не делать (пользователь гарантирует что такого не произойдёт CLK_ADDRESS_NONE). Третий показывает, что делать, если на входе не целые координаты. Возможные варианты: приблизить ближайшим значением (CLK_FILTER_NEAREST), линейно проинтерполировать (CLK_FILTER_LINEAR).
Краткий пример. Замыливаем изображение по среднему значению в области:
__kernel void ImageDiff(__read_only image2d_t bmp1, __write_only image2d_t bmpOut)
<
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST;
int2 coords = (int2)(get_global_id(0), get_global_id(1));
uint4 sumall = (uint4)(0,0,0,0);
int sum = 0;
for ( int i=-10;i <11;i++)
for ( int j=-10;j <11;j++)
int2 newpol = (int2)(i,j)+coords;
sumall+= read_imageui(bmp1, smp, newpol);
sum++;
>
sumall = sumall/sum;
write_imageui(bmpOut, coords, sumall);
>
* This source code was highlighted with Source Code Highlighter .
Полезности
Ну, думаю, с кратким описанием я справился. Теперь несколько ссылок для более подробного изучения, ежели кому понадобиться.
Официальный сайт с документацией.
Сайт с примерами и понятными описаниями.
Неплохой пдфничек, там структура OpenCl устройств неплохо прорисованна.
Так же есть 2 презентации про OpenCL на русском. В них довольно мало инфы и нет связующего текста. Правда есть примеры неплохие. Первая. Вторая.
OpenCL. Что это такое и зачем он нужен? (если есть CUDA)
Многие, наверное, слышали или читали на хабре об OpenCL – новом стандарте для разработки приложений для гетерогенных систем. Именно так, это не стандарт для разработки приложений для GPU, как многие считают, OpenCL изначально задумывался как нечто большее: единый стандарт для написания приложений, которые должны исполняться в системе, где установлены различные по архитектуре процессоры, ускорители и платы расширения.
Предпосылки появления OpenCL
Основным местом, где можно встретить гетерогенные системы, являются высокопроизводительные вычисления: от моделирования физических процессов в пограничном слое до кодирования видео и рендеринга трехмерных сцен. Раньше подобные задачи решали применяя суперкомпьютеры либо очень мощные настольные системы. С появлением технологий NVidia CUDA/AMD Stream стало возможным относительно просто писать программы, использующие вычислительные возможности GPU.
Стоит отметить, что подобные программы создавались и раньше, но именно NVidiaа CUDA обеспечила рост популярности GPGPU за счет облегчения процесса создания GPGPU приложений. Первые GPGPU приложения в качестве ядер (kernel в CUDA и OpenCL) использовали шейдеры, а данные запаковывались в текстуры. Таким образом необходимо было быть хорошо знакомым OpenGL или DirectX. Чуть позже появился язык Brook, который немного упрощал жизнь программиста (на основе этого языка создавалась AMD Stream (в ней используется Brook+) ).
CUDA стала набирать обороты, а между тем (а точнее несколько ранее) в кузнице, расположенной глубоко под землей, у подножия горы Фуджи (Fuji), японскими инженерами был выкован процессор всевластия Cell (родился он в сотрудничестве IBM, Sony и Toshiba). В настоящее время Cell используется во всех суперкомпьютерах, поставляемых IBM, на его основе постоены самые производительные в мире суперкомпьютеры (по данным top500). Чуть менее года назад компания Toshiba объявила о выпуске платы расширения SpursEngine для PC для ускорения декодирования видео и прочих ресурсоемких операций, используя вычислительные блоки (SPE), разработанные для Cell. В википедии есть статья, в кратце описывающая SpursEngine и его отличия от Cell.
Примерно в то же время (около года назад) оживилась и S3 Graphics (на самом деле VIA), представив на суд общественности свой новый графический адаптер S3 Graphics Chrome 500. По заявлениям самой компании этот адаптер так же умеет ускорять всяческие вычисления. В комплекте с ним поставляется программный продукт (графический редактор), который использует все прелести такого ускорения. Описание технологии на сайте производителя.
Итак, что мы имеем: машина, на которой проводятся вычисления может содержать процессоры x86, x86-64, Itanium, SpursEngine (Cell), NVidia GPU, AMD GPU, VIA (S3 Graphics) GPU. Для каждого из этих типов процессов существует свой SDK (ну кроме разве что VIA), свой язык программирования и программная модель. То есть если Вы захотите чтобы ваш движок рендеринга или программа расчета нагрузок на крыло боинга 787 работала на простой рабочей станции, суперкомпьютере BlueGene, или компьютере оборудованном двумя ускорителями NVidia Tesla – Вам будет необходимо переписывать достаточно большую часть программы, так как каждая из платформ в силу своей архитектуры имеет набор жестких ограничений.
Так как программисты – народ ленивый, и не хотят писать одно и то же для 5 различных платформ с учетом всех особенностей и учиться использовать разные программные средства и модели, а заказчики – народ жадный и не хотят платить за программу для каждой платформы как за отдельный продукт и оплачивать курсы обучения для программистов, было решено создать некий единый стандарт для программ, исполняющихся в гетерогенной среде. Это означает, что программа, вообще говоря, должна быть способна исполняться на компьютере, в котором установлены одновременно GPU NVidia и AMD, Toshiba SpursEngine итд.
Решение проблемы
Для разработки открытого стандарта решили привлечь людей, у которых уже есть опыт (весьма успешный) в разработке подобного стандарта: Khronos Group, на чьей совести уже OpenGL и OpenML и еще много всего. OpenCL является торговой маркой Apple Inc., как сказано на сайте Khronos Group: «OpenCL is a trademark of Apple Inc., and is used under license by Khronos. The OpenCL logo and guidelines for its usage in association with Conformant products can be found here:
http://developer.apple.com/softwarelicensing/agreements/opencl.html». В разработке (и финансировании, конечно же), кроме Apple, участвовали такие воротилы IT как AMD, IBM, Activision Blizzard, Intel, NVidia итд. (полный список тут).
Компания NVidia особо не афишировала свое участие в проекте, и быстрыми темпами наращивала функциональность и производительность CUDA. Тем временем несколько ведущих инженеров NVidia участвовали в создании OpenCL. Вероятно, именно участие NVidia в большой мере определило синтаксическую и идеологическую схожесть OpenCL и CUDA. Впрочем программисты от этого только выиграли – проще будет перейти от CUDA к OpenCL при необходимости.
Первая версия стандарта была опубликована в конце 2008 года и с тех пор уже успела претерпеть несколько ревизий.
Почти сразу после того как стандарт был опубликован, компания NVidia заявила что поддержка OpenCL не составит никакой сложности для нее и в скором времени будет реализована в рамках GPU Computing SDK поверх CUDA Driver API. Ничего подобного от главного конкурента NVidia – AMD слышно не было.
Драйвер для OpenCL был выпущен NVidia и прошел проверку на совместимость со стандартом, но все еще доступен только для ограниченного круга людей – зарегистрированных разработчиков (заявку на регистрацию подать может любой желающий, в моем случае рассмотрение заняло 2 недели, после чего по почте пришло приглашение). Ограничения доступа к SDK и драйверам заставляют задуматься о том, что на данный момент существуют какие-то проблемы или ошибки, которые пока не удается исправить, то есть продукт все еще находится в стадии бета-тестирования.
Реализация OpenCL для NVidia была достаточно легкой задачей, так как основные идеи сходны: и CUDA и OpenCL – некоторые расширения языка С, со сходным синтаксисом, использующие одинаковую программную модель в качестве основной: Data Parallel (SIMD), так же OpenCL поддерживает Task Parallel programming model – модель, когда одновременно могут выполняться различные kernel (work-group содержит один элемент). О схожести двух технологий говорит даже то что NVidia выпустила специальный документ о том как писать для CUDA так, чтобы потом легко перейти на OpenCL.
Как обстоят дела на настоящий момент
Основной проблемой реализации OpenCL от NVidia является низкая производительность по сравнению с CUDA, но с каждым новым релизом драйверов производительность OpenCL под управлением CUDA все ближе подбирается к производительности CUDA приложений. По заявлениям разработчиков такой же путь проделала и производительность самих CUDA приложений – от сравнительно невысокой на ранний версиях драйверов до впечатляющей в настоящее время.
А что же делала в этот момент AMD? Ведь именно AMD (как сторонник открытых стандартов – закрытый PhysX vs. открытый Havoc; дорогой Intel Thread Profiler vs. бесплатный AMD CodeAnalyst) делала большие ставки на новую технологию, учитывая что AMD Stream не удавалось хоть сколь-нибудь соревноваться в популярности с NVidia CUDA – виною тому отставание Stream от CUDA в техническом плане.
Летом 2009 года компания AMD сделала заявление о поддержке и соответствии стандарту OpenCL в новой версии Stream SDK. На деле же оказалось, что поддержка была реализована только для CPU. Да, именно так, это ничему не противоречит – OpenCL стандарт для гетерогенных систем и ничего не мешает Вам запустить kernel на CPU, более того – это очень удобно в случае если в системе нет другого OpenCL устройства. В таком случае программа будет продолжать работать, только медленнее. Или же вы можете задействовать все вычислительные мощности, которые есть в компьютере – как GPU так и CPU, хотя на практике это не имеет особого смысла, так как время исполнения kernel’ов которые исполняются на CPU будет намного больше тех что исполняются на GPU – скорость процессора станет узким местом. Зато для отладки приложений это более чем удобно.
Поддержка OpenCL для графических адаптеров AMD так же не заставила себя долго ждать – по последним сообщениям компании версия для графических чипов сейчас находится на стадии подтверждения соответствия спецификациям стандарта. После чего она станет доступна всем желающим.
Так как OpenCL должен работать поверх некоторой специфической для железа оболочки, а значит для того чтобы можно этот стандарт действительно стал единым для различных гетерогенных систем – надо чтобы соответствующие оболочки (драйверы) были выпущены и для IBM Cell и для Intel Larrabie. Пока от этих гигантов IT ничего не слышно, таким образом OpenCL остается еще одним средством разработки для GPU на ряду с CUDA, Stream и DirectX Compute.
- OpenTK — библиотека-обертка над OpenGL, OpenAL и OpenCL для .Net.
- PyOpenCL – обертка над OpenCL для Pyton.
- Java обертка для OpenCL.
Заключение
Технология OpenCL представляет интерес для различных компания IT сферы – от разработчиков игр до производителей чипов, а это означает что у нее большие шансы стать фактическим стандартом для разработки высокопроизводительных вычислений, отобрав этот титул у главенствующей в этом секторе CUDA.
В будущем я планирую более подробную статью о самом OpenCL, описывающую что из себя представляет эта технология, ее особенности, достоинства и недостатки.
Спасибо за внимание.
Opencl что это amd
В статье рассматриваются основные принципы дизайна OpenCL согласно стандарту версии 1.1. Не вдаваясь в излишние на данном уровне изложения подробности описаны 4 модели, на которых держится стандарт: модель платформы, модель исполнения, модель памяти и модель программирования. В статье не приведено ни единой строчки программного кода, так как цель — лишь ввести читателя в мир разработки на OpenCL, осветив различные стороны его дизайна.
Microsoft представила C++ AMP
Опубликовано shapovalovts в Чт, 06/16/2011 — 10:07
На конференции AMD Fusion 11 Developer Summit Герб Саттер (Herb Sutter) анонсировал новую технологию для разработки гетерогенных приложений на языке С++, получившую название C++ Accelerated Massive Parallelism (AMP). Данная технлогия, по заявлению разработчиков Microsoft, позволит использовать возможности параллельного выполнения кода на CPU, так и выполнения кода на GPU.
IBM OpenCL Development Kit для Linux на платформе Power v0.3
Опубликовано shapovalovts в Пнд, 04/11/2011 — 19:11
31 марта IBM представила средства разработки OpenCL-приложений для операционной системы Linux, работающей на платформе Power. По заявлению разработчиков была реализована спецификация OpenCL 1.1 в полном объеме.
NVIDIA выпускает CUDA 4.0 Toolkit
Опубликовано shapovalovts в Втр, 03/01/2011 — 16:07
Анонсирована новая версия инструментария CUDA. В 4-й версии разработчики сделали упор на упрощение портирования CPU-кода на платформу GPGPU-вычислений. Этого они попытались достичь с помощью следующих основных нововведений: NVIDIA GPUDirect 2.0, Unified Virtual Addressing, а также библиотеки Thrust C++ Template Performance Primitives.
AMD представляет OpenCL University Kit
Опубликовано shapovalovts в Вс, 02/13/2011 — 23:05
OpenCL University Kit — это набор материалов от компании AMD для обучения студентов технологии OpenCL в течение одного семестра. Материалы логически поделены на 13 лекций, каждая из которых включает информацию для преподавателя и план проведения лекций. Также данный набор дополнен примерами кода для трех лекций (2,3 и 13), кодом полноценного OpenCL-приложения (лекция 9) и набором упражнений для закрепления материала.
Представлена альфа-версия Intel OpenCL SDK
Опубликовано shapovalovts в Пт, 11/19/2010 — 11:59
Intel представила новый SDK для написания программ с использованием OpenCL 1.1, оптимизированный для процессоров семейства Intel Core. Первый вариант SDK требует операционную систему Microsoft Windows 7 или Vista. Сейчас этот SDK поддерживает компиляцию лишь 32 битных приложений.
Представлен Aparapi — API для написание OpenCL-ядер на Java
Опубликовано shapovalovts в Сб, 09/25/2010 — 14:30
Aparapi — новое API языка программирования JAVA для написания и выполнения OpenCL-ядер на графических акселераторах компании AMD. Если целевая платформа поддерживает OpenCL, то байткод будет проанализирован «на лету» и в случае отсутствия неконвертируемого в OpenCL кода, он транслируется в OpenCL, компилируется и запускается на GPU (если подходящий GPU присутствует). В противном случае, код будет запущен как обычный байткод на CPU.
Особенности следования стандартам
Опубликовано taleks в Пт, 09/24/2010 — 16:06
Недавно возникла необходимость сравнить производительность работы кернела (kernel) на CPU и GPU под Linux. Однако, выяснилось, что устройства типа CPU не оказалось в списке доступных устройств для платформы, в связи с чем было произведено обновление драйверов nVidia под Linux со старых версий до последних. Результат был неожиданным: получилась неработающая программа. Почему и как побороть — далее.
В MATLAB добавлена поддержка GPGPU
Опубликовано shapovalovts в Пт, 09/24/2010 — 14:12
В новой версии MATLAB R2010b в компонент Parallel Computing Toolbox была добавлена официальная поддержка вычислений с применением GPU. В настоящий момент объявлено о поддержке графических процессоров NVidia. Об этом заявила глава направления Parallel Computing в компании MathWorks Silvina Grad-Freilich на конференции NVIDIA GPU Technology Conference.
PGI представляет CUDA C компилятор для платформы x86
На конференции SC10 в Новом Орлеане компания PGI (The Portland Group) продемонстрирует новый CUDA компилятор для 32 и 64 битных платформ x86. Об этом заявил CEO компании PGI Jen-Hsun Huang на конференции GPU Technology Conference 2010, проходящей в San Jose.
OpenCL
OpenCL (от англ. Open Computing Language — открытый язык вычислений) — фреймворк для написания компьютерных программ, связанных с параллельными вычислениями на различных графических (англ. GPU ) и центральных процессорах (англ. CPU ). В фреймворк OpenCL входят язык программирования, который базируется на стандарте C99, и интерфейс программирования приложений (англ. API ). OpenCL обеспечивает параллелизм на уровне инструкций и на уровне данных и является реализацией техники GPGPU. OpenCL является полностью открытым стандартом, его использование не облагается лицензионными отчислениями.
Цель OpenCL состоит в том, чтобы дополнить OpenGL и OpenAL, которые являются открытыми отраслевыми стандартами для трёхмерной компьютерной графики и звука, пользуясь возможностями GPU. OpenCL разрабатывается и поддерживается некоммерческим консорциумом Khronos Group, в который входят много крупных компаний, включая Apple, AMD, Intel, Nvidia, ARM, Sun Microsystems, Sony Computer Entertainment и другие.
История
OpenCL первоначально был разработан в компании Apple Inc. Apple внесла предложения по разработке спецификации в комитет Khronos. Вскоре компания AMD решила поддержать разработку OpenCL (и DirectX 11), который должен заменить фреймворк Close to Metal. [1] [2]
16 июня 2008 года была сформирована рабочая группа Khronos Compute для разработки спецификаций OpenCL. В нее вошли Apple, nVidia, AMD, IBM, Intel, ARM, Motorola и другие компании, в том числе специализирующиеся на создании компьютерных игр. Работа велась в течение пяти месяцев, по истечении которых 9 декабря 2008 года организация Khronos Group представила первую версию стандарта.
OpenCL 1.0 была выпущена вместе с Mac OS X 10.6. [3]
5 апреля 2009 года компания AMD анонсировала доступность для загрузки beta-версии комплекта разработчика ATI Stream SDK v2.0 в который входит язык мультипроцессорного программирования OpenCL.
20 апреля 2009 года nVidia представила бета-драйвер и комплект для разработки программного обеспечения (SDK) с поддержкой открытого GPGPU-стандарта OpenCL. Данный бета-драйвер предназначен для разработчиков, участвующих в программе «OpenCL Early Access», которые уже с 20 апреля могут принять участие в тестировании бета-версии. Для участников программы «GPU Computing Registered Developers» бета-версия драйвера OpenCL будет доступна позже. [4] [5] [6]
26 ноября 2009 года компания nVidia выпустила драйвер с поддержкой OpenCL 1.0 (rev 48).
Для получения наглядного представления, как технология OpenCL использует возможности 24-ядерной системы для отрисовки видеоэффектов, рекомендуется посмотреть следующий демо-ролик: [1].
OpenCL 1.1 был представлен организацией Khronos Group 14 июня 2010 года. В новой версии значительно расширены функциональные возможности для параллельного программирования, гибкость и производительность, а также:
- Новые типы данных, включая 3-компонентные векторы и дополнительные форматы изображений;
- Обработка команд из нескольких потоков хоста и обработки буфера между несколькими устройствами;
- Операции по регионам буфера включая чтение, запись и копирование 1D, 2D или 3D прямоугольных областей;
- Расширенное использование события для управления и контроля выполнения команд;
- Улучшенное взаимодействие с OpenGL за счет эффективного обмена изображениями.
OpenCL 1.2 увидел свет 15 ноября 2011 года. В новой версии отмечено множество небольших улучшений, связанных с увеличением гибкости языка и оптимизацией производительности. Из добавленных в OpenCL 1.2 значительных новшеств отмечается:
- Партицирование устройств — возможность разбиения на уровне OpenCL-приложения устройства на несколько подустройств для непосредственной привязки работ к конкретным вычислительным блокам, резервирования ресурсов для более приоритетных задач или более эффективного совместного использования аппаратных ресурсов, таких как кэш;
- Раздельная компиляция и связывание объектов — появилась возможность создания динамических библиотек, позволяющих использовать в сторонних программах, ранее реализованные подпрограммы с OpenCL-вычислениями;
- Расширенная поддержка изображений, включая возможность работы с одномерными изображениями и массивами одномерных или двухмерных изображений. Кроме того, в расширении для организации совместного доступа (sharing) добавлена возможность создания OpenCL-изображения на основе отдельных текстур OpenGL или массивов текстур;
- Встроенные OpenCL-ядра теперь позволяют использовать возможности специализированного или непрограммируемого аппаратного обеспечения и связанных с ним прошивок. Например, появилась возможность использования возможностей и более тесной интеграции с фреймворком OpenCL таких устройств, как DSP-процессоры или видео кодировщики/декодировщики;
- Возможность бесшовного совместного использования поверхностей (Media Surface Sharing) между OpenCL и API DirectX 9/11.
События
- 3 марта 2011 — Khronos Group объявляет о создании рабочей группы WebCL для разработки JavaScript-интерфейса к стандарту OpenCL. Это создает потенциал для того, чтобы использовать GPU и многоядерных процессоры для параллельной обработки вычислений в веб-браузере. [7]
- 4 мая 2011 — подразделение Nokia Research представило открытое расширение WebCL для браузера Firefox. [7]
- 1 июля 2011 — Samsung Electronics представила открытый прототип WebCL для движка WebKit. [7]
- 8 августа 2011 — AMD выпустила OpenCL-драйвер AMD Accelerated Parallel Processing (APP) Software Development Kit (SDK) v2.5, заменив ATI Stream SDK.
- 15 ноября 2011 — комитет Khronos представил обновлённую спецификацию OpenCL 1.2. В новой версии отмечено множество небольших улучшений, связанных с увеличением гибкости языка и оптимизацией производительности.
- 1 декабря 2012 — комитет Khronos представил очередное обновление спецификации OpenCL 1.2. В новой версии улучшено взаимодействие с OpenGL, улучшена безопасность в WebGL, добавлена поддержка загрузки OpenCL программ из промежуточного представления SPIR.
Особенности языка
Ключевыми отличиями используемого языка от Си (стандарт ISO 1999 года) являются:
- Отсутствие поддержки указателей на функции, рекурсии, битовых полей, массивов переменной длины (VLA), стандартных заголовочных файлов [8]
- Расширения языка для параллелизма: векторные типы, синхронизация, функции для Work-items/Work-Groups [8]
- Квалификаторы типов памяти: __global, __local, __constant, __private
- Иной набор встроенных функций
Примеры
// создание вычислительного контекста для GPU (видеокарты) context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL); // создание очереди команд queue = clCreateCommandQueue(context, NULL, 0, NULL); // выделение памяти в виде буферов memobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float)*2*num_entries, srcA, NULL); memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*2*num_entries, NULL, NULL); // создание программы из исходных текстов program = clCreateProgramWithSource(context, 1, &fft1D_1024_kernel_src, NULL, NULL); // компиляция программы clBuildProgram(program, 0, NULL, NULL, NULL, NULL); // создание объекта kernel из скомпилированной программы kernel = clCreateKernel(program, "fft1D_1024", NULL); // подготовка аргументов clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobjs[0]); clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobjs[1]); clSetKernelArg(kernel, 2, sizeof(float)*(local_work_size[0]+1)*16, NULL); clSetKernelArg(kernel, 3, sizeof(float)*(local_work_size[0]+1)*16, NULL); // задание N-D диапазона с размерностями work-item и отправка в очередь исполнения global_work_size[0] = num_entries; local_work_size[0] = 64; clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
Непосредственные вычисления (Основаны на отчете «Fitting FFT onto the G80 Architecture») [10]
// Данный код вычисляет FFT длины 1024, путем разбиения на 16, 16 и 4 __kernel void fft1D_1024 (__global float2 *in, __global float2 *out, __local float *sMemx, __local float *sMemy) int tid = get_local_id(0); int blockIdx = get_group_id(0) * 1024 + tid; float2 data[16]; // адрес начала обрабатываемых данных в глобальной памяти in = in + blockIdx; out = out + blockIdx; globalLoads(data, in, 64); // coalesced global reads fftRadix16Pass(data); // in-place radix-16 pass twiddleFactorMul(data, tid, 1024, 0); // локальная перестановка с использованием локальной памяти localShuffle(data, sMemx, sMemy, tid, (((tid & 15) * 65) + (tid >> 4))); fftRadix16Pass(data); // in-place radix-16 pass twiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplication localShuffle(data, sMemx, sMemy, tid, (((tid >> 4) * 64) + (tid & 15))); // 4 вызова БПФ порядка 4 fftRadix4Pass(data); // radix-4 function number 1 fftRadix4Pass(data + 4); // radix-4 function number 2 fftRadix4Pass(data + 8); // radix-4 function number 3 fftRadix4Pass(data + 12); // radix-4 function number 4 // coalesced global writes globalStores(data, out, 64); >
Полноценная реализация БПФ на OpenCL доступна на сайте Apple [11]
Применение
OpenCL находит применение, как одна из реализаций концепции GPU общего назначения, в различном ПО.
- WinZip v16.5 (2012) от Corel — помимо обновлённого движка для улучшения оптимизации его для многоядерных процессоров, добавлена поддержку OpenCL для GPU AMD (однако, не для Intel и Nvidia) — при этом прирост производительности в этом приложении на APU Trinity и Llano составил до 45 %. [12]
Вы поможете проекту, исправив и дополнив его.
См. также
- GPGPU
- AMD FireStream
- CUDA
- Close to Metal
- APU (гибридные процессоры)
Примечания
- ↑AMD Drives Adoption of Industry Standards in GPGPU Software Development. AMD. Архивировано из первоисточника 19 марта 2012.
- ↑AMD Backs OpenCL, Microsoft DirectX 11. eWeek. Архивировано из первоисточника 19 марта 2012.
- ↑Apple Previews Mac OS X Snow Leopard to Developers. Apple. Архивировано из первоисточника 19 марта 2012.
- ↑Andrew HumberNVIDIA Releases OpenCL Driver To Developers (англ.) . NVIDIA (20 апреля 2009 года). — Оригинальная новость на официальном сайте NVIDIA Corporation. Архивировано из первоисточника 19 марта 2012.Проверено 21 апреля 2009.
- ↑Павел ШубскийNVIDIA открыла GPGPU для разработчиков под OpenCL. Игромания (журнал) (21 апреля 2009 года). (недоступная ссылка — история) Проверено 21 апреля 2009.
- ↑Сергей и Марина БондаренкоДрайвер OpenCL для разработчиков от NVIDIA. 3DNews (21 апреля 2009 года). Проверено 21 апреля 2009.
- ↑ 123Для WebKit представлена реализация технологии WebCL (рус.) . opennet.ru (4 июля 2011 года). Архивировано из первоисточника 19 марта 2012.Проверено 31 октября 2011.
- ↑ 12 AMD. Introduction to OpenCL Programming 201005, page 89-90
- ↑OpenCL. SIGGRAPH2008 (14 августа 2008). Архивировано из первоисточника 19 марта 2012.Проверено 14 августа 2008.
- ↑Fitting FFT onto G80 Architecture (PDF). Vasily Volkov and Brian Kazian, UC Berkeley CS258 project report (May 2008). Архивировано из первоисточника 19 марта 2012.Проверено 14 ноября 2008.
- ↑ . OpenCL on FFT. Apple (16 Nov 2009). Проверено 7 декабря 2009.
- ↑AMD Trinity: тесты — OpenCL // THG
Ссылки
- www.khronos.org/opencl Официальная страница стандарта (eng.)
- «OpenCL: What you need to know», MacWorld, август2008
- OpenCL Русскоязычный сайт
- Редакция THGOpenCl: приложения с ускоренной постобработкой (рус.) . Tom’s Hardware (18 апреля 2012 года). Архивировано из первоисточника 15 мая 2012.Проверено 24 апреля 2012.
| Стандарты The Khronos Group | |
|---|---|
| COLLADA · EGL · OpenCL · OpenGL · OpenGL ES · OpenGL SC · OpenKODE · OpenMAX · OpenML · OpenSL ES · OpenVG · OpenWF · WebGL |
- Программное обеспечение по алфавиту
- GPGPU
- Кроссплатформенные реализации языков программирования
Wikimedia Foundation . 2010 .
Полезное
Смотреть что такое «OpenCL» в других словарях:
- OpenCL — Entwickler Khronos Group Aktuelle Version 1.2 (16. November 2011) Betriebssystem plattformunabhängig Kategorie Programmierschnittstelle … Deutsch Wikipedia
- OpenCL C — OpenCL Entwickler: Khronos Group Aktuelle Version: 1.0 (8. Dezember 2008) Betriebssystem: plattformunabhängig Kateg … Deutsch Wikipedia
- OpenCL — Desarrollador Grupo Khronos http://www.khronos.org/opencl Información general Diseñador Apple … Wikipedia Español
- OpenCL — This article is about the parallel computing library. For the cryptographic library formerly named OpenCL, see Botan (programming library). OpenCL Original author(s) Apple Inc. Developer(s) Khronos Group … Wikipedia
- OpenCL — Ne doit pas être confondu avec OpenGL. OpenCL (Open Computing Language) est la combinaison d une API et d un langage de programmation dérivé du C, proposé comme un standard ouvert par le Khronos Group. OpenCL est conçu pour programmer des… … Wikipédia en Français
- Open CL — OpenCL Ne doit pas être confondu avec OpenGL. OpenCL (Open Computing Language) est la combinaison d une API et d un langage de programmation dérivé du C, proposé comme un standard ouvert par le Khronos Group. L objectif d OpenCL est de… … Wikipédia en Français
- Bullet Physics Library — Физический движок … Википедия
- Mac OS X Snow Leopard — Mac OS X v10.6 Snow Leopard Part of the Mac OS X family … Wikipedia
- Compute Unified Device Architecture — Entwickler Nvidia Aktuelle Version 4.0 (Mai 2011) Betriebssystem Windows, Linux, MacOS X Kategorie … Deutsch Wikipedia
- AMD Fusion — Codename(s) Fusion Desna Ontario Zacate Llano Hondo (cancelled) Wichita (cancelled) Krishna (cancelled) Trinity Weatherford Richland IGP Wrestler WinterPark BeaverCreek ATI/Radeon Driver related BTC[1] [2] … Wikipedia
- Обратная связь: Техподдержка, Реклама на сайте
- �� Путешествия
Экспорт словарей на сайты, сделанные на PHP,
WordPress, MODx.
- Пометить текст и поделитьсяИскать в этом же словареИскать синонимы
- Искать во всех словарях
- Искать в переводах
- Искать в ИнтернетеИскать в этой же категории