Обновлено: 22.01.2025

OpenCL (англ. Open Computing Language — открытый язык вычислений) — фреймворк для написания компьютерных программ, связанных с параллельными вычислениями на различных графических и центральных процессорах, а также FPGA. В OpenCL входят язык программирования, который основан на стандарте языка программирования Си C99, и интерфейс программирования приложений. 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. [1] [2]

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 года. [3]

5 апреля 2009 года компания AMD анонсировала доступность для загрузки бета-версии набора разработчика 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.

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

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

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

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

События

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

Непосредственные вычисления (основаны на отчете «Fitting FFT onto the G80 Architecture») [14] :

Полноценная реализация БПФ на OpenCL доступна на сайте Apple [15] .

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

Intel® FPGA SDK for OpenCL™ software technology 1 is a world class development environment that enables software developers to accelerate their applications by targeting heterogeneous platforms with Intel CPUs and FPGAs. This environment combines Intel’s state-of-the-art software development frameworks and compiler technology with the revolutionary, new Intel® Quartus® Prime Software to deliver next generation development environment that abstracts FPGA details while delivering optimized results. Intel® FPGA SDK for OpenCL™ software technology enables you to fully leverage the unique capabilities of FPGAs to deliver acceleration performance with power efficiency and low latency.

Features

• Microsoft* Visual Studio or Eclipse*-based Intel® Code Builder for OpenCL™ API now with FPGA support
• Fast FPGA emulation based on Intel's compiler technology
• Create OpenCL™ project jump-start wizard

• Common Development Environment for both host (CPU) and accelerator (FPGA)
• Syntax highlighting and code auto-completion features
• Availability of Board Support Packages (BSP) for FPGA development kits

• Quick static FPGA resource and performance analysis
• Support for fast and incremental FPGA compile

What's New in 20.1

Intel® FPGA SDK for OpenCL­™ software technology v20.1 includes key enhancements that deliver improved performance and productivity.

An Intuitive Design Environment

• Dynamic Profiling via Intel® VTune analysis.
  
• Improved utilization and padding to reduce local memory usage for depths that are not a power of two.
  

Meeting Performance Requirements

• Improved reporting and pragmas for fine-tuned control over loop fusion to merge sequential loops at the same level to improve throughput and reduce area.
  
• Enabled stall enable clusters pragma to reduce local memory usage by removing FIFO instantiations and latency between functions.
  

Getting Started

to target the latest Intel® Stratix® 10, Intel® Arria® 10, and Intel® Cyclone® 10 GX devices to target Stratix® V and Cyclone® V SoC devices
• The software installation file includes the OpenCL™ software and Intel® Quartus® Prime Pro Edition Software or Intel® Quartus® Prime Design Software
• The Intel® Quartus® Prime Software does not require a license purchase when used as part of the Intel® FPGA SDK for OpenCL™ software technology
  
• A Board Support Package (BSP) is needed to run your OpenCL™ application. Download an Intel BSP (development kit BSP) or purchase a partner provided BSP, or create a custom BSP
• Read the Intel® FPGA SDK for OpenCL™ software technology getting started guide

Frequently Asked Questions

What is OpenCL™ application?

The OpenCL™ standard is the first open, royalty-free, unified programming model for accelerating algorithms on heterogeneous systems. OpenCL™ application allows the use of a C-based programming language for developing code across different platforms, such as CPUs, GPUs, and FPGAs. A key benefit of OpenCL™ application is that it is a portable, open, royalty-free standard, which is a key differentiator versus proprietary programming models. OpenCL™ standard is a programming model for software engineers and a methodology for system architects. It is based on standard ANSI C (C99) with extensions to extract parallelism. OpenCL™ platform also includes an application programming interface (API) for the host to communicate with the hardware accelerator traditionally over PCI Express* or one kernel to communicate with another without host interaction. Intel® FPGA SDK for OpenCL™ software technology provides a vendor extension, an I/O, and a Host Channel API to stream data into a kernel directly from a streaming I/O interface such as 10 GB Ethernet.

What is a Board Support Package (BSP) and where can I find reference BSPs?

In order to create and run applications in the Intel® FPGA SDK for OpenCL™ software technology, you need to either use an Intel or partner provided BSP or create your own BSP (see links in ‘Getting Started’ section). A board support package (BSP) is a collection of libraries and drivers that will form the lowest layer of your application software stack. Your Software Applications must link against or run on top of a given software platform.

To find reference BSPs, see the “Getting Started” section.

Which Intel® FPGA development kits have supported BSPs?

The following development kits have an available BSP:

• Intel® Stratix® 10 GX
• Stratix® V
• Intel® Arria® 10 SoC
• Intel® Arria® 10 GX

Where can I find OpenCL™ design examples?

OpenCL™ design examples are part of the Intel® FPGA SDK for OpenCL TM software technology. You can find the design example folder at <OpenCL_Installation_Directory>/examples_aoc.

Стандартный подход к написанию программ является линейным – операция b выполняется после завершения операции a. Но что делать в случае если таких операций десятки тысяч, а задача требует быстрого произведения данных операций?

Допустим, вы разрабатываете приложение с искусственным интеллектом для распознавания предметов основываясь на данных получаемых с сенсора, или пишете графическое приложение, которое требует множество вычислений на каждом из пикселей. Как сделать так, чтобы обработка информации происходила в реальном времени, а задержка между получением информации, её обработкой и выдачей результата была минимальной?

Те кто сталкивался с данной проблемой, наверняка знают про многопоточность (multithreading) — способ параллельного запуска определённых фрагментов кода для ускорения обработки этого кода. Но, к сожалению, этот метод подходит не всегда, и когда требуется произведение одинаковых операций на большом количестве данных, гораздо эффективнее будет прибегнуть к массивному параллелизму и использовать фреймворки OpenCL или CUDA.

В этой статье мы познакомимся с массивным параллелизмом и напишем программу для параллельных вычислений, используя фреймворк OpenCL.

OpenCL (Open Computing Language) — фреймворк разработанный Apple в 2008 году и поддерживаемый Khronos Group с 2009 года. Он позволяет создавать программы для параллельного исполнения на различных вычислительных девайсах (CPU и GPU), упакованные в "кернели"-ядра (kernels) — части кода, которые будут отправлены на вычислительный девайс для произведения каких-то операций.

OpenCL — замечательный инструмент, который может ускорить вашу программу в десятки, если не сотни\тысячи раз. К сожалению, из-за отсутствия простых и доступных гайдов по данной спецификации, в особенности на русском языке, а так-же особенностям работы с ней (о чём мы поговорим чуть позже), начинающему разработчику может быть очень трудно разобраться в теме и попытки изучить спецификацию остаются безуспешными. Это большое упущение, и цель этой статьи — не только доступно и подробно рассказать о спецификации и продемонстрировать как работать с ней, но и написать гайд, который мне хотелось бы прочитать самому когда я начал изучать OpenCL.

Массивный параллелизм — топик для целого цикла лекций по информатике и инженерии, по этому давайте ограничимся небольшим обобщением, чтобы понимать с чем мы имеем дело.

В линейном программировании мы имеем чёткую последовательность действий: a, b, c, d; действие b не будет выполнено до того как завершиться a и c не будет выполнено пока не завершиться b. Но что делать, если нам, например, требуется найти суммы элементов из двух массивов (листов), и в каждом массиве по 100,000 элементов? Последовательное вычисление заняло бы достаточно долгое время, так-как нам пришлось бы совершить минимум 100,000 операций. А что если такая процедура требует многочисленного повторения и результат нужен в реальном времени с минимальной задержкой? Тут нам и приходит на помощь массивный параллелизм!

Допустим, мы хотим вычислить суммы 0 + 3, 1 + 2, 2 + 1, 3 + 0 и записать результаты в массив. В линейном программировании, мы воспользуемся циклом for или while, где операции будут выполняться последовательно, и схема вычислений будет выглядеть примерно так:

Через последовательную итерацию, мы будем выполнять первое действие, записывать результат в массив и переходить к следующему действию. Поскольку для выполнения данных вычислений нам требуется 4 итерации, для наглядности, давайте обозначим что время (t) потраченное на выполнение данных операций равняется 4.

Имея достаточно место в массиве, нам не обязательно ждать пока заполнится предыдущий элемент массива перед тем как записывать следующий, мы можем присвоить задачам индексы, соответствующие их "месту" в массиве, и тем самым, записывать ответ каждого вычисления в правильное место, не дожидаясь пока предыдущая задача будет выполнена.

Другими словами, если 0 + 3 — вычисление номер один, а 1 + 2 — вычисление номер два, мы можем посчитать 1 + 2 и записать ответ во второе место в массиве (res[1]) не зависимо от того записан ли ответ 0 + 3 в первое место (res[0]). При этом, если мы попытаемся одновременно записать информацию в ячейки res[0] (первое место) и res[1] (второе место), у нас не возникнет ошибки.

Таким образом, с помощью массивного параллелизма мы можем одновременно и независимо друг от друга выполнить все нужные нам операции, и записать все ответы в массив, содержащий результаты, всего за одно действие, тем самым, сократив t до 1. Мы только что сократили временную сложность алгоритма (время работы алгоритма) до константного значения 1 (одно действие).

Перед тем как мы перейдём к практике, важно прояснить специфику работы с OpenCL и разобраться как переписывать линейный код в параллельном формате. Не волнуйтесь — это совсем не сложно!

Кернель (kernel, вектор — функция, отправляемая на вычислительный девайс в контексте работы OpenCL) и хост (host — код, вызывающий OpenCL; ваш код) существуют, по большему счёту, изолированно друг от друга.

Компиляция и запуск кернеля OpenCL происходят внутри вашего кода, во время его исполнения (онлайн, или runtime execution). Так-же, важно понимать что кернель и хост не имеют общего буфера памяти и мы не можем динамически выделять память (с помощью malloc() и подобных) внутри кернеля. Обмен информацией между двумя системами происходит посредством отправления между ними заранее выделенных регионов памяти.

Модель памяти OpenCL выглядит так:

При создании контекста OpenCL (среда в которой существует данный инстанс OpenCL), мы обозначаем нашу задачу как NDRange — общий размер вычислительной сетки — количество вычислений которые будут выполнены (в примере с суммами использованном выше, NDRange = 4). Информацию записанную в глобальную память (global memory) мы можем получить из любого элемента NDRange.

При отправке задачи на девайс OpenCL (например, видеокарту компьютера), наш NDRange разбивается на рабочие группы (work-groups) — локальная память (local memory), которые содержат в себе рабочие единицы (work-items, изолированные инстансы кернеля) — приватная память (private memory).

Мы не можем считать объект из локальной памяти в глобальную. Это значит что значение существующее внутри одной рабочей группы не будет доступно из другой группы.

Локальная память имеет доступ к объектам из глобальной памяти, но не может получить информацию из приватной памяти. Наконец, приватная память может считывать из глобальной и локальной памяти, но всё что существует внутри её, остаётся в ней.

Стоит добавить, что от выбора памяти зависит скорость выполнения программы. Так, запрос к глобальной памяти является самым времязатратным, и работать с глобальной памятью стоит только в случае необходимости.

В большинстве случаев, будет грамотнее объявить переменную в локальной или приватной памяти (нет смысла создавать переменную в глобальной памяти, если она не используется за пределами одной рабочей единицы). Так-же, зачастую будет грамотным решением перевести глобальный объект в локальную или приватную память, после получения его внутри кернеля, если требуется произвести с ним множество операций внутри одной рабочей единицы, так-как это может значительно ускорить выполнение программы.

Для обмена данными между хостом и кернелем, внутри хоста мы создаём специальный буфер, содержащий информацию, нужную для вычислений на девайсе OpenCL, а так-же буфер с выделенным местом (памятью), в который будут записаны результаты вычислений кернеля. Эти элементы отправляются в кернель, он производит свои вычисления, записывает результат в буфер который мы специально для этого подготовили и отправляет его обратно в хост.

Технология параллельного программирования OpenCL

Эту статью можно считать продолжением статьи [1] о технологии CUDA, здесь мы поговорим о технологии параллельного программирования OpenCL.

1. Введение

OpenCL (Open Computing Language ) это спецификация, описывающая технологию параллельного программирования, которая в первую очередь ориентирована на GPGPU. Изначально она была разработана компанией Apple, в последствии для развития спецификаций OpenCL был образована группа разработчиков Khronos Compute [2], в неё вошли Apple, nVidia, AMD, IBM, Intel, ARM, Motorola и др. Первая версия стандарта была опубликована в конце 2008 года.

В отличии от nVidia CUDA, AMD Stream и т.п., в OpenCL изначально закладывалась мультиплатформенность, т.е. OpenCL программа должна без изменений в коде работать на GPU разных типов (разных производителей). Такая программа без изменений должна работать даже на CPU без GPU, хотя в этом случае она может выполняться существенно медленнее чем на GPU.

2. Схема работы с аппаратурой

Итак - хотим мультиплатформенность и желательно без существенных потерь в производительности. Достигается этот результат следующим образом [3,4].

OpenCL-программа работает с т.н. платформами (platform) . Платформа это программный пакет, который поставляется соответствующим разработчиком аппаратных средств. Например "AMD Accelerated Parallel Processing" или "Intel OpenCL". При этом несколько платформ могут работать одновременно на одной машине.

Каждая платформа включает в себя ICD (Installable Client Driver) -- программный интерфейс OpenCL для работы с устройствами, которые эта платформа поддерживает.

В среде Linux список ссылок на ICD, присутствующих в системе, обычно хранится в каталоге /etc/OpenCL/vendors/ , а библиотека libOpenCL.so выполняет роль диспетчера (ICD loader) , т.е. она направляет вызовы OpenCL функций на устройства, через соответствующие ICD.

Рис.1: схема работы OpenCL программы с аппаратурой

[ Здесь ] можно увидеть пример конфигурации оборудования. На машине с процессором Intel Core2 CPU 6300 и графическим ускорителем nVidia Quadro FX1700, развернуты три платформы OpenCL: nVidia, AMD, Intel. При этом платформа nVidia поддерживает только GPU Quadro FX1700, а платформы AMD и Intel - только Intel Core2 CPU 6300 в качестве вычислительного устройства. Таким образом каждая платформа содержит одно вычислительное устройство.

Для определения возможностей OpenCL можно воспользоваться утилитой clinfo, также [ здесь ] можно скачать исходник программы, считывающую эту информацию.

Тут надо ещё сказать о версиях OpenCL SDK. При развёртывании платформы на компьютере необходимо уточнить версию пакета. Например, в процессе проведения экспериментов выяснилось, что Intel OpenCL SDK 2014 не поддерживает работу с процессором Intel Core2 CPU 6300, который был установлен на машине. Проблема решилась установкой более старого пакета Intel OpenCL SDK 2012.

3. Структура OpenCL-программы

В OpenCL (аналогично CUDA), программа разделяется на две части: первая часть - управляющая, вторая - вычислительная. В роли управляющего устройства ( host ) выступает центральный процессор (CPU), вычислительное устройство ( device ) выбираем из списка платформ и их устройств. Обычно используется GPU, но не обязательно, это может быть и тот же CPU.

Код, который должен выполняться на device, оформляется специальным способом. Поскольку device могут быть от разных производителей и разных типов, то скомпилировать один бинарник для всех не получится. Решается эта проблема следующим образом.

Текст ядра (kernel) , т.е. части программы выполняемой на device, включается в основную часть программы (выполняемой на host) в "чистом" виде т.е. в виде текстовой строки. Этот исходник компилируется средствами OpenCL непосредственно в процессе работы программы (runtime) для выбранного в данный момент вычислительного устройства, это происходит каждый раз при запуске OpenCL-программы.

Так же есть возможность, скомпилированный таким образом, код ядра сохранить и при инициализации вычислительной программы загружать готовый бинарник. Но в этом случае мы теряем универсальность, этот код будет работать только на одном (данном) типе устройств.

1. получить информацию о платформах и устройствах
   clGetPlatformIDs(),clGetPlatformInfo(), clGetDeviceIDs(),clGetDeviceInfo()
2. выбрать устройства и создать для них контекст
   clCreateContext()
3. создать ядро из текста программы
   clCreateProgramWithSource(), clBuildProgram(), clCreateKernel()
4. выделить память для данных на устройствах
   clCreateBuffer()
5. создать очередь комманд для устройтва
   clCreateCommandQueue(), clCreateCommandQueueWithProperties( )
6. скопировать данные с host на device
   clEnqueueWriteBuffer()
7. назначить параметры выполнения ядра
   clSetKernelArg()
8. запуск ядра
   clEnqueueNDRangeKernel()
9. скопировать результат с device на host
   clEnqueueReadBuffer()
10. обработка результата
11. завершение работы, освобождение ресурсов
    clReleaseMemObject(), clReleaseKernel(), clReleaseProgram(), clReleaseCommandQueue(), clReleaseContext()

4. Пример OpenCL-программы

В качестве первого примера рассмотрим простую вычислительную задачу: С := d * A + B, где d - константа, А,В и С векторы заданного размера. __kernel void kernel1(const float alpha, __global float *A, __global float *B, __global float *C) < int idx = get_global_id(0); C[idx] = alpha* A[idx] + B[idx]; > Листинг 1: код ядра для задачи сложения векторов

Программа запускает много (по количеству элементов векторов) параллельных процессов (тредов) на device, каждый тред получает свой номер (get_global_id), исходя из него считывает свою часть данных из глобальной (__global) памяти device, выполняет вычисления и записывает свою часть результата обратно в память.

Код программы можно скачать [ здесь ]. Для сравнения напишем ещё простую (последовательную) программу и сравним время затраченное на вычисления. Ниже в таблице представлены результаты скорости выполнения программы на разных платформах.

платформа	устройство	время(ms)
NVIDIA	Quadro FX 1700	3.2
AMD	Intel Core2 CPU 6300	19.1
Intel	Intel Core2 CPU 6300	11.0
-	Intel Core2 CPU 6300	9.3

Таблица 1: результаты работы программы saxpy для разных платформ

По результатам, представленным в таб.1, видно, что выполнение примера на GPU почти в три раза быстрее чем на CPU. Для CPU платформа Intel показала лучший результат чем AMD, что очевидно для Intel Core2 CPU 6300. Наилучший результат для CPU показала простая программа, что можно объяснить эффективной работой компилятора и отсутствием необходимости выполнять дополнительные процедуры OpenCL.

5. Группы процессов и их конфигурация

Каждое вычислительное устройство обладает своими ограничениями по количеству одновременно выполняемых тредов, и хотя их общее количество в программе может быть велико, выполняться они будут частями или группами, максимальный размер группы зависит от конкретного устройства.

Иногда треды удобно формировать в виде решетки. Рассмотрим пример задачи умножения матриц.

Листинг 2: код ядра для задачи умножения матриц

Имеем на входе матрицы A[M*K], B[K*N] и соответственно буфер для результата C[M*N]. Программа создаёт M*N тредов в виде решетки MxN, каждый тред [r,c] отрабатывает одну ячейку в матрице результата.

Код программы можно скачать [ здесь ].

Для сравнения напишем ещё простую (последовательную) программу и сравним время затраченное на вычисления. Ниже в таблице представлены результаты скорости выполнения программы на разных платформах.

платформа	устройство	время(ms)
NVIDIA	Quadro FX 1700	5169.5
AMD	Intel Core2 CPU 6300	2279.0
Intel	Intel Core2 CPU 6300	920.1
-	Intel Core2 CPU 6300	1340.6

Таблица 2: результаты работы программы gemm1 для разных платформ

Результаты несколько расстраивают, поскольку из таб.2 видно, что GPU показала наихудший результат. Проблема в крайне неэффективном использовании памяти GPU. Далее мы исправим это затруднение.

6. Модели памяти и синхронизация процессов

В статье [1] была приведена схема организации GPU. Из этой схемы видно, что этот тип устройств обладает сложно организованной памятью, которая может работать с разной скоростью.

• Память global - основная память уcтройства, самая большая по размеру (512MB для FX1700) и самая медленная, она является общей для всех тредов.
• Память local - общая память для одной группы тредов (shared в терминах CUDA), этот тип быстрее global но существенно меньше по размеру (16KB для FX1700)
• Память private - память треда, быстрая но маленькая (8KB для FX1700)

Вернёмся к предыдущему примеру с умножением матриц. В процессе работы ядро выполняет много повторных чтений из памяти global. Попробуем сократить количество чтений с помощью организации быстрого кэша [5].

Листинг 3: код ядра для задачи умножения матриц с кэшированием

Имеем на входе матрицы A[M*K], B[K*N] и соответственно результат C[M*N]. Программа создаёт M*N тредов в виде решетки MxN, группами размера TSxTS, каждый тред отрабатывает одну ячейку в матрице результата, группа тредов кэширует блоки исходных матриц и далее выполняет операции с этим кэшем.

Для того, что бы кэш корректно был заполнен необходима синхронизация группы тредов (barrier) , достигая barrier программа ждет, пока все треды группы соберутся в этой точке и только после этого продолжает выполнение.

Код программы можно скачать [ здесь ].

Ниже в таблице представлены результаты скорости выполнения программы на разных платформах.

платформа	устройство	время(ms)
NVIDIA	Quadro FX 1700	169.5
AMD	Intel Core2 CPU 6300	2301.1
Intel	Intel Core2 CPU 6300	1402.7
-	Intel Core2 CPU 6300	1340.6

Таблица 3: результаты работы программы gemm2 для разных платформ

Из таблицы 3 видно, что модифицированная программа умножения матриц (gemm2) показывает вполне удовлетворительный результат производительности, работая гораздо быстрее первого варианта (gemm1).

Читайте также:

  
• Как отклонить эсф от поставщика в 1с
  
• Как удалить файл max в 3ds max
  
• Браузер ps4 не воспроизводит видео
  
• Как удалить сохраненные пароли в яндекс браузере
  
• Как убрать зелень с волос в фотошопе