OpenCL

OpenCL (англ. Open Computing Language — открытый язык вычислений) — фреймворк для написания компьютерных программ, связанных с параллельными вычислениями на различных графических и центральных процессорах, а также FPGA. В OpenCL входят язык программирования, который основан на стандарте языка программирования Си C99, и API. OpenCL обеспечивает параллелизм на уровне инструкций и на уровне данных и является осуществлением техники GPGPU. OpenCL является полностью открытым стандартом, его использование не облагается лицензионными отчислениями.

OpenCL
Тип	API
Автор	Apple
Разработчик	Khronos Group
Написана на	C++ и Си
Операционная система	Кроссплатформенное программное обеспечение
Первый выпуск	28 августа 2009
Аппаратные платформы	x86-64, IA-32 и ARM
Последняя версия	3.0.11 (6 мая 2022)
Лицензия	EULA
Сайт	khronos.org/opencl/
Медиафайлы на Викискладе

Цель OpenCL состоит в том, чтобы дополнить открытые отраслевые стандарты для трёхмерной компьютерной графики и звука — OpenGL и OpenAL, соответственно, — возможностями GPU для высокопроизводительных вычислений. OpenCL разрабатывается и поддерживается некоммерческим консорциумом Khronos Group, в который входят много крупных компаний, включая AMD, Apple, ARM, Intel, Nvidia, Sony Computer Entertainment и другие.

История

OpenCL первоначально был разработан в компании Apple Inc. Apple внесла предложения по разработке спецификации в комитет Khronos. Вскоре компания AMD решила поддержать разработку OpenCL (и DirectX 11), который должен заменить фреймворк Close to Metal.

16 июня 2008 года была образована рабочая группа Khronos Compute для разработки спецификаций OpenCL. В неё вошли Apple, nVidia, AMD, IBM, Intel, ARM, Motorola и другие компании, в том числе специализирующиеся на создании компьютерных игр. Работа велась в течение пяти месяцев, по истечении которых 9 декабря 2008 года организация Khronos Group представила первую версию стандарта.

OpenCL 1.0 был впервые показан общественности 9 июня 2008, а выпущен вместе с Mac OS X 10.6, 28 августа 2009 года.

5 апреля 2009 года компания AMD анонсировала доступность для загрузки бета-версии набора разработчика ATI Stream SDK v2.0, в который входит язык мультипроцессорного программирования OpenCL.

20 апреля 2009 года nVidia представила бета-драйвер и набор для разработки программного обеспечения (SDK) с поддержкой открытого GPGPU-стандарта OpenCL. Этот бета-драйвер предназначен для разработчиков, участвующих в программе «OpenCL Early Access», которые уже с 20 апреля могут принять участие в испытании бета-версии. Для участников программы «GPU Computing Registered Developers» бета-версия драйвера OpenCL будет доступна позже.

26 ноября 2009 года компания nVidia выпустила драйвер с поддержкой OpenCL 1.0 (rev 48).

Для получения наглядного представления, как технология OpenCL использует возможности 24-ядерной системы для отрисовки видеоэффектов, рекомендуется посмотреть следующий демо-ролик:[1] Архивная копия от 9 марта 2017 на Wayback Machine.

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.

OpenCL 2.0 был представлен 22 июля 2013 года и стандартизирован 18 ноября того же года.

• Общая виртуальная память - Позволяет ядрам узла и устройств совместно использовать структуры данных, основанные на комплексных адресных ссылках, устраняя явные пересылки между узлом и устройствами, повышая при этом гибкость программирования.
• Вложенный параллелизм - Обновление улучшило возможности программирования и увеличило производительность приложений.
• Универсальное адресное пространство - Позволяет записать функции без наименования адресного пространства, что повышает гибкость и экономит время за счет устранения необходимости записи нескольких функций.
• Атомарные операции C11 со стороны устройства - Подмножество атомарных и синхронизирующих операций C11 обеспечивает параллельное выполнение потоков для безопасной работы над общими наборами данных.
• Каналы - Объекты памяти, организованные по принципу FIFO, что упрощает структуры данных общей очереди.

OpenCL 2.1 был представлен 3 марта 2015 года и стандартизирован 16 ноября того же года. В нём было переписано ядро с языка C на C++14.

OpenCL 3.0 был представлен 27 апреля 2020 года и стандартизирован 30 сентября того же года. Среди заметных изменений можно отметить то, что API OpenCL 3.0 теперь охватывает все версии OpenCL (1.2, 2.x), без предоставления отдельных спецификаций для каждой версии.

События

• 3 марта 2011 — Khronos Group объявляет о создании рабочей группы WebCL для разработки JavaScript-интерфейса к стандарту OpenCL. Это создаёт потенциал для того, чтобы использовать GPU и многоядерные процессоры для параллельной обработки вычислений в веб-браузере.
• 4 мая 2011 — подразделение Nokia Research представило открытое расширение WebCL для браузера Firefox.
• 1 июля 2011 — Samsung Electronics представила открытый прототип WebCL для движка WebKit.
• 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), стандартных заголовочных файлов;
• расширения языка для параллелизма: векторные типы, синхронизация, функции для Work-items/Work-Groups;
• квалификаторы типов памяти: __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»):

  // Данный код вычисляет 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.

Применение

См. также: en:List of OpenCL applications

OpenCL находит применение, как одна из реализаций концепции GPU общего назначения, в различном ПО.

• WinZip v16.5 (2012) от Corel — помимо обновлённого движка для улучшения оптимизации его для многоядерных процессоров, добавлена поддержка OpenCL для GPU AMD (однако, не для Intel и Nvidia) — при этом прирост производительности в этом приложении на APU Trinity и Llano составил до 45 %.

См. также

• GPGPU
• AMD FireStream
• CUDA
• Close to Metal
• APU (гибридные процессоры)

Литература

• Aaftab Munshi; Benedict R. Gaster; Timothy G. Mattson; James Fung; Dan Ginsburg. OpenCL Programming Guide. — Addison-Wesley Professional, 2011. — 648 p. — ISBN 978-0-321-74964-2.

Ссылки

• www.khronos.org/opencl Архивная копия от 9 августа 2011 на Wayback Machine Официальная страница стандарта (eng.)
• «OpenCL: What you need to know» Архивная копия от 13 сентября 2008 на Wayback Machine, MacWorld, август 2008
• OpenCL Русскоязычный сайт Архивная копия от 14 апреля 2010 на Wayback Machine
• OpenCL Англоязычный сайт Архивная копия от 6 августа 2020 на Wayback Machine
• Редакция THG. OpenCl: приложения с ускоренной постобработкой  (рус.). Tom's Hardware (18 апреля 2012). Дата обращения: 24 апреля 2012. Архивировано 26 мая 2012 года.
• AMD OpenCL SDK
• Nvidia OpenCL Samples Архивная копия от 24 декабря 2013 на Wayback Machine
• Intel OpenCL SDK Архивная копия от 2 сентября 2014 на Wayback Machine
• Международный семинар по OpenCL Архивная копия от 15 июля 2017 на Wayback Machine