OpenCL

OpenCL
Тип	API
Разработчик	Apple Inc., Khronos Group
Операционная система	Кроссплатформенное программное обеспечение
Последняя версия	1.2 (15 ноября 2011)
Лицензия	Свободная
Сайт	http://www.khronos.org/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.

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

OpenCL 1.0 была выпущена вместе с Mac OS X 10.6.

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

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

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 и многоядерных процессоры для параллельной обработки вычислений в веб-браузере.
• 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. В новой версии отмечено множество небольших улучшений, связанных с увеличением гибкости языка и оптимизацией производительности.

Особенности языка

Ключевыми отличиями используемого языка от C99 являются:

• Отсутствие поддержки указателей на функции, рекурсии, битовых полей, массивов переменной длины (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

См. также

• GPGPU
• AMD FireStream
• CUDA
• Close to Metal
• Intel Larrabee

Ссылки

• www.khronos.org/opencl Официальная страница стандарта (eng.)
• «OpenCL: What you need to know», MacWorld, август 2008
• OpenCL-Z: OpenCL Information utility
• OpenCL Русскоязычный сайт
• Заголовочные файлы для Delphi и FreePascal
• Статья: «Введение в OpenCL»  (рус.)