OpenCL ( англ. Open C omputing L anguage — открытый язык вычислений) — фреймворк для написания компьютерных программ , связанных с параллельными вычислениями на различных графических и центральных процессорах , а также FPGA . В OpenCL входят язык программирования , который основан на стандарте языка программирования Си C99 , и API . OpenCL обеспечивает параллелизм на уровне инструкций и на уровне данных и является осуществлением техники GPGPU . 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-ядерной системы для отрисовки видеоэффектов, рекомендуется посмотреть следующий демо-ролик: от 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 .

Применение

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 .

Ссылки

• от 9 августа 2011 на Wayback Machine Официальная страница стандарта (eng.)
• от 13 сентября 2008 на Wayback Machine , MacWorld , август 2008
• от 14 апреля 2010 на Wayback Machine
• от 6 августа 2020 на Wayback Machine
• Редакция THG. (рус.) . Tom's Hardware (18 апреля 2012). Дата обращения: 24 апреля 2012. 26 мая 2012 года.
• от 24 декабря 2013 на Wayback Machine
• от 2 сентября 2014 на Wayback Machine
• от 15 июля 2017 на Wayback Machine