Opencl library

Как ускорить вычисления и повысить производительность программ с помощью принципов массивного параллелизма и OpenCL — Разработка на vc.ru

Opencl library

Стандартный подход к написанию программ является линейным – операция 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 и т.д.) для выполнения используя ресурсы, например, видеокарты — она имеет десятки, сотни, тысячи вычислительных единиц (ядер), которые могут производить операции параллельно, независимо друг от друга.

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

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

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

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

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

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

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

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

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

В общих чертах, схема работы с OpenCL выглядит следующим образом:

Начнём с установки и настройки OpenCL на нашем компьютере.

Первым делом, скачайте и обновите драйвера вашей видеокарты. Это очень важно, так-как OpenCL не будет работать если ваши драйвера его не поддерживают.

Если вы пользуетесь Apple MacOS, всё что вам нужно сделать — убедиться что у вас установлена новейшая версия ОС и XCode. OpenCL поставляется с вашей системой.

В противном случае, вам требуется скачать и установить подходящую имплементацию OpenCL с сайта производителя вашей видеокарты.

В качестве простого примера, давайте возьмём задачу векторного сложения.

Допустим, у нас есть два массива, A и B, равного размера. Наша цель — найти сумму элементов А и B и записать её в элемент нового массива, С, такого-же размера.

Стандартный метод решения данной задачи подразумевает использование цикла for или while, для последовательной итерации через массив, выполняя операцию:

for (int i = 0; i < РАЗМЕР_МАССИВА; i++){ C[i] = A[i] + B[i];}

Алгоритм очень простой, но имеет линейную временную сложность, O(n), где n — размер массива.

Так-как каждая итерация этого цикла не зависит от других итераций, мы можем переформулировать наш алгоритм под параллелизм: каждая итерация этого цикла может быть выполнена параллельно и одновременно.

Таким образом, если мы имеем n ядер в нашем вычислительном девайсе, временная сложность становится константной O(1). Это как-раз то, о чём мы уже говорили раньше.

Начать написание нашей программы стоит именно с кернеля — функции, которую мы отправим на видеокарту для параллельного вычисления. Давайте рассмотрим как это делается.

Функцию кернеля нужно прописывать в отдельном файле с расширением .cl

Давайте поместим этот файл в корневую директорию нашего проекта и назовём vector_add.cl

Так-как язык OpenCL это слегка модифицированный С, проблем с написанием у нас не возникнет, и данный код выполнит нашу задачу:

__kernel void addition(__global const int *A, __global const int *B, __global int *C){ int i = get_global_id(0); // получаем индекс обрабатываемого элемента, он-же позиция элемента в массиве C[i] = A[i] + B[i]; // выполняем сложение}

  • Здесь, всё что мы делаем, это избавляемся от цикла for. Мы представляем что нужную нам операцию требуется выполнить всего один раз. Делается это потому что наш кернель будет отправлен на какое-то количество ядер видеокарты одновременно, и каждое ядро параллельно посчитает один элемент под индексом i из нашего массива. Стоит заметить что OpenCL не гарантирует выполнение операций по порядку — он может посчитать сначала двадцатый элемент, а потом третий.
  • Ключевое слово __kernel даст компилятору понять что функция addition — именно функция кернеля. Каждая декларация основной функции кернеля должна с него начинаться. Поскольку кернели ничего не возвращают, мы всегда указываем void как тип return функции (речь идёт об основных функциях, не вспомогательных. С ними вы можете работать так-же как и с обычными функциями С)
  • __global перед типом аргумента означает что данная переменная будет обработана в глобальной памяти (для обращения к локальной или приватной памяти, используются __local и __private соответственно)
  • Указатели на массивы A и B имеют модификатор constant, так как в этих массивах мы будем передавать информацию в кернель. В массив С мы будем записывать результат, по этому такого модификатора нет.
  • Функция get_global_id() позволяет получить уникальный идентификатор рабочей единицы (work-item) для данного измерения. Мы отправляем 0, потому-что у нас одномерный массив, соответственно, мы используем 0 для указания первого (и единственного) измерения. Да, с помощью OpenCL мы можем обрабатывать 2-х, 3-х, …-мерные объекты.

Программа-хост контролирует запуск и выполнение кернелей на вычислительных девайсах. Хост пишется на языке С, но существуют различные библиотеки-обвязки для работы с другими языками, например С++ или Python.

#include #include #ifdef __APPLE__# include // для компьютеров на MacOsX#else# include // для компьютеров на Win\Linux указывайте путь к файлу cl.h #endif #define MAX_SRC_SIZE (0x100000) // максимальный размер исходного кода кернеля

Всё приведённое ниже находится в теле функции main(). Для наглядности, я не стал это прописывать.

int main(void){ …код… return(0);}

int i;const int arr_size = 1024; // размер наших массивов int *A = (int *)malloc(sizeof(int) * arr_size); // выделяем место под массивы А и Вint *B = (int *)malloc(sizeof(int) * arr_size); for(i = 0; i < arr_size; i++) // наполняем массивы данными для отправки в кернель OpenCL{ A[i] = i; B[i] = arr_size - i;}

int fd;char *src_str; // сюда будет записан исходный код кернеляsize_t src_size; src_str = (char *)malloc(MAX_SRC_SIZE); // выделяем память для исходного кодаfd = open(“vector_add.cl”, O_RDONLY); // открываем файл с кодом кернеляif (fd

Источник: https://vc.ru/dev/159198-kak-uskorit-vychisleniya-i-povysit-proizvoditelnost-programm-s-pomoshchyu-principov-massivnogo-parallelizma-i-opencl

OpenCL™ Runtimes for Intel® Processors

Opencl library

Obtain runtimes to execute or develop OpenCL™ applications on Intel­® Processors

Important Change

There is a change in OpenCL™ CPU runtime for Windows* distribution in the 2020 February release to be consistent with Linux* distribution. The OpenCL CPU runtime is removed from the OpenCL driver for Windows starting in the 2020 February release version “igfx_win10_100.7870.exe”.

Intel® Graphics Technology Runtimes

Execute OpenCL™ applications on Intel® Processors with Intel® Graphics Technology.

  • Specifically target Intel® HD Graphics, Intel® Iris® Graphics, and Intel® Iris® Pro Graphics if available on Intel® Processors.
  • Runtimes for Intel® Graphics Technology are often deployed in tandem with an Intel® CPU runtime.
  • Consider graphics runtimes when developing OpenCL™ applications with the Intel® SDK for OpenCL™ Applications or Intel® System Studio.

Check release notes to ensure supported targets include your target device. For Intel® processors older than supported targets, please see the legacy deployment page.

Linux* OS

Repository Install Guidance *Easy* | Manual Download and Install | Build | README | FAQ

Note: The latest OpenCL runtime for CPU requires GNU* gcc version 7.3 or newer.

Intel® Graphics Compute Runtime for OpenCL™ Driver is deployed with package managers for multiple distributions. Please see the documentation on the GitHub* portal for deployment instructions.

Considerations for deployment:

  • Ensure the deployment system has the (libOpenCL.so) ICD loader runtime from either:
  • The Intel® Graphics Compute Runtime for OpenCL™ Driver depends on the i915 kernel driver. Necessary i915 features are available with relatively recent Linux* OS kernels. The recommended kernel is the validation kernel cited in documentation. In general, deployments after the 4.11 kernel should be OK. Make sure to review the release notes and documentation for more specifics.

Windows* OS

  • Intel® Graphics Compute Runtime for OpenCL™ Driver is included with the Intel® Graphics Driver package for Windows* OS.
  • Download Options
    • System Vendor
      • See your vendor website for a graphics or video driver download for the system
    • Intel® Download Center
      • Navigate to “Graphics Drivers” for recent releases.
      • Try the system vendor first in consideration of vendor support. System vendors may disable Intel® Graphics Driver install.
    • Intel® Driver Update Utility
    • The graphics driver package is built in with Windows* 10 OS install.

      However, the built-in default deployment may not contain latest features.

  • Release Notes
    • In the Download Center navigate to “Graphics Drivers” for Release Notes.

Intel® Xeon® Processor OR Intel® Core™ Processor (CPU) Runtimes

Execute OpenCL™ kernels directly on Intel® CPUs as OpenCL™ target devices.

  • Consider an OpenCL™ CPU implementation for Intel® systems without Intel® Graphics Technology.
  • Systems with Intel® Graphics Technology can simultaneously deploy runtimes for Intel® Graphics Technology and runtimes for Intel® CPU (x86-64).
  • For application developers, the CPU-only runtime is pre-included with the Intel® SDK for OpenCL™ Applications or Intel® System Studio: OpenCL™ Tools component.

Check release notes to ensure supported targets include your target device. For Intel® processors older than supported targets, see the legacy deployment page.

Intel® CPU Runtime for OpenCL™ Applications 18.1 for Linux* OS (64bit only)

​Download

  • Size 125 MB
  • See supported platform details in the Release Notes.
  • Ubuntu* install uses an rpm translator
  • The Linux* OS CPU runtime package also includes the ICD loader runtime (libOpenCL.so). The runtime installer should set the deployment system to see this ICD loader runtime by default. When examining system libraries, administrators may observe ICD loader runtimes obtained from other places. Examples include the system package manager (for example with ocl-icd) or as part of the Intel® SDK for OpenCL™ Applications.
  • Maintenance and updates are now provided in the Experimental Intel® CPU Runtime for OpenCL™ Applications with SYCL support implementation. This implementation is listed later in this article.
  • MD5 83c428ab9627268fc61f4d8219a0d670
  • SHA1 5f2fa6e6bc400ca04219679f89ec289f17e94e5d

Intel® CPU Runtime for OpenCL™ Applications 18.1 for Windows* OS (64bit or 32bit)

Download

  • Size 60 MB
  • CPU-only deployments should use the .msi installer linked in the Download button, and consider removal of the Intel® Graphics Technology drivers where applicable.
  • CPU & Graphics deployments should use the Intel® Graphics Technology driver package, which contains both CPU (x86-64) and Intel® Graphics Technology implementations.
  • See supported operating system details in the Release Notes
  • Maintenance and updates are now provided in the Experimental Intel® CPU Runtime for OpenCL™ Applications with SYCL support implementation. This implementation is listed later in this article.
  • MD5 8e2404800146ed6921d658dd71b8ff
  • SHA1 451d96d37259cb111fe8832d5513c5562efa3e56

Experimental Intel® CPU Runtime for OpenCL™ Applications with SYCL support

Download from Intel staging area for llvm.org contribution: prerequisites.

Installation Guide on Github*

  • This OpenCL™ implementation for Intel® CPUs is actively maintained. It is currently in *beta* as of article publication date.
    • OpenCL 1.2, 2.0, and 2.1 programs can use this runtime.
    • The DPC++/SYCL implementation can use this runtime. This runtime additionally supports the SYCL runtime stack. OpenCL™ developers are highly encouraged to explore Intel® DPC++ compiler and SYCL.
  • Deployments with the Intel® CPU Runtime for OpenCL™ Applications 18.1 and this Experimental runtime are not jointly validated at article publication time. Use one or the other implementation, but not both.
  • Feedback can be provided at the Intel® oneAPI Data Parallel C++ forum. Issues are also communicated at the Intel staging area for llvm.org contribution.

Develop OpenCL™ Applications

Tools to develop OpenCL™ applications for Intel® Processors

Intel® System Studio

  • For compilation, cross-platform, IoT, power considerate development, and performance analysis.
    • OpenCL™ development tools component:
      • Develop OpenCL™ applications targeting Intel® Xeon® Processors, Intel® Core™ Processors, and/or Intel® Graphics Technology.
      • Develop applications with expanded IDE functionality, debug, and analysis tools.
        • Note: Some debug and analysis features have been removed from recent versions of the SDK.
      • Earlier versions of the SDK contain an experimental OpenCL™ 2.1 implementation. Intel® CPU Runtime for OpenCL™ Applications 18.1 was intended as a replacement for the experimental implementation.

  • Visit the Intel® System Studio portal

Intel® SDK for OpenCL™ Applications

  • Standalone distribution of Intel® System Studio: OpenCL™ Tools component.
  • Develop OpenCL™ Applications targeting Intel® Xeon® Processors, Intel® Core™ Processors, and/or Intel® Graphics Technology.
  • Develop applications with expanded IDE functionality, debug, and analysis tools.
    • Note: Some debug and analysis features have been removed from recent versions of the SDK.
    • Earlier versions of the SDK contain an experimental OpenCL™ 2.1 implementation suitable for development testing on CPU OpenCL™ targets. Intel® CPU Runtime for OpenCL™ Applications 18.1 was intended as a replacement for that experimental implementation.
  • See release notes, requirements, and download links through the Intel® SDK for OpenCL™ Applications portal.

Intel® FPGA SDK for OpenCL™ Software Technology

  • Build OpenCL™ Applications and OpenCL™ kernels for Intel® FPGA devices.
  • See release notes, requirements, and download links through the SDK’s portal webpage.
  • For OpenCL™ runtimes and required system drivers, visit Download Center for FPGAs.

Intel® Distribution of OpenVINO™ toolkit

  • The Intel® Distribution of OpenVINO™ toolkit is available for vision and deep learning inference. It benefits from OpenCL™ acceleration for each of these components:
    • Intel® Deep Learning Deployment Toolkit
    • OpenCV
    • OpenVX*
  • For a developer oriented overview, see videos on the techdecoded.intel.io training hub.

Intercept Layer for Debugging and Analyzing OpenCL™ Applications

  • The Intercept Layer for Debugging and Analyzing OpenCL™ Applications (clIntercept) can intercept, report, and modify OpenCL™ API calls.
  • No application-level modifications nor OpenCL™ implementation modifications are necessary.
  • clIntercept functionality can supplement removed functionality from recent releases of the Intel® SDK for OpenCL™ Applications.

Additional resources

Getting Started with Intel® SDK for OpenCL™ Applications

Download Intel® SDK for OpenCL™ Applications

Intel® OpenCL™ Forum

*OpenCL and the OpenCL logo are trademarks of Apple Inc. used by permission by Khronos.

Источник: https://software.intel.com/content/www/us/en/develop/articles/opencl-drivers.html

OpenCL. Как начать

Opencl library

Всем привет! Какое-то время назад я начал копать тему с OpenCL под C#. Но наткнулся на трудности, связанные с тем, что не то, что под C#, а вообще по этой теме очень мало материала. Какую-то вводную по OpenCL можно почерпнуть здесь. Так же простой, но работающей старт OpenCL описан вот тут.

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

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

Первоначально я хотел написать статью-минисамоучитель OpenCL, которая содержала в себе информацию, о там что это, как устроено, как писать код и какие-то рекомендации, основанные на моем опыте. Но в процессе понял, что если даже быть кратким, то уткнусь в ограничения объема статьи.

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

Вопросы с устройством памяти, архитектурой и прочим будут описаны в следующих статьях. Друзья, сразу хотел бы сказать, что мой опыт в OpenCL пока, к сожалению, далек от гуру/йода уровня, но на вопросы постараюсь отвечать изо всех сил. А если что-то не знаю, то буду делиться ресурсами и видением того как это на самом деле должно работать.

Что. Где. Как

OpenCL это технология связанная с параллельными компьютерными вычислениями на различных типах графических и центральных процессоров. Тема с параллельным вычислениями на GPU совсем недавно широко продвигалась вместе с технологией CUDA. Данное продвижение в основном обеспечивалось усилиями компании Nvidia. Отличия OpenGL и CUDA уже широко обсуждались.

OpenСL позволяет работать как с CPU так и с GPU, но думаю нам более интересно будет сосредоточиться на работе с GPU. Для использования данной технологии понадобиться мало мальски современная видеокарта. Главное это проверить, что устройство функционирует нормально. На всякий случай напоминаю что это можно сделать в диспетчере устройств.

Если в данном окне вы видите какие-то фейлы или ворнинги, то вам прямая дорога на сайт производителя вашей видеокарты. Свежие драйвера должны решить проблему с функционированием железа и как следствие дать доступ к мощностям OpenCL.

Первоначально я планировал использовать OpenCL по C#.

Но наткнулся на проблему, что все существующие фреймворки типа Cloo или OpenCLNet являются самописными и Khronos не имеет к ним никакого отношения и следовательно не гарантирует их стабильную работу. Ну и все мы помним главную проблему — очень мало примеров.

Исходя из этого вначале я бы хотел представить примеры написанные на C++, а уже потом, получив подтверждение того что OpenCL ведет себя так как мы ожидаем, привинтить прокси в виде C# фреймворка.

Итак, чтобы использовать OpenCL через С++ необходимо найти его API. Для этого открывайте переменные среды, и ищете там переменную со страшным названием, намекающую своим названием на производителя вашей видеокарты. У меня данная переменная называется «AMDAPPSDKROOT». После этого можете посмотреть что лежит по указанному пути. Там ищите папочку include\CL. Кстати, обычно в папочки include, рядом с папкой CL лежит и папка GL, предоставляющая доступ к знаменитой графической библиотеки. Теперь создаем проект в Visual Studio, подключаем в свойствах проекта папку include (в моем случае $(AMDAPPSDKROOT)\include\) и в бой!

Инфраструктура

Мы должны помнить, что мы будем работать с OpenCL не через API, а при помощи API. Вроде бы кажется, что эти две фразы практически идентичны, но это не так. К примеру, вспомните OpenGL.

Как происходит там работа (урощённый вариант) — вначале мы настраиваем какие-то общие параметры, а потом прямо из кода вызываем методы типа «нарисовать сферу», «изменить параметры источника света» и т.д.

Так вот в OpenCL сценарий другой:

  1. При помощи API получаем доступ к устройствам, которые поддерживают OpenCL. Это часть приложения обычно называется хостом;
  2. Пишем код который будет выполняться на устройстве. Этот код называется kernel.

    Данный код о хосте вообще ничего не знает. Его может дернуть любой хост

  3. При помощи API прогружаем код kernel и запускаем его выполнение на выбранном устройстве.

Как видите наше приложение будет иметь комплексную инфраструктуру.

Давайте займемся ее настройкой!

Так как на предыдущем шаге мы предусмотрительно подсоединили папочку include, то теперь вы можем просто добавить ссылку на заголовочный файл cl.h, который даст доступ к API. При добавление cl.h, стоит добавить проверку выбора платформы:

#ifdef __APPLE__#include #else#include #endif Теперь необходимо выбрать устройство на котором будет отрабатывать наш код и создать контекст в котором будут жить наши переменные. Как это сделать показано ниже: /* получить доступные платформы */ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms); /* получить доступные устройства */ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices); /* создать контекст */context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret); /* создаем команду */command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
Обращаю внимание на переменную ret. Это переменная, которая содержит числовое значение которое возвращает та или иная функция. Если ret== 0, то функция выполнилась корректно, если нет, то значит произошла ошибка.
Так же заслуживает внимание константа CL_DEVICE_TYPE_DEFAULT, она запрашивает устройство, которое используется для вычислений на OpenCL по умолчанию. Вместо данной константы могут быть использованы другие. К примеру:

  • CL_DEVICE_TYPE_CPU — запросит существующие CPU.
  • CL_DEVICE_TYPE_GPU — запросит существующие GPU.

Kernel

Отлично. Настроили инфраструктуру. Теперь возьмемся за kernel. Kernel — это просто функция объявление, которой начинается с ключевого слова __kernel. Синтаксис языка программирования OpenCL базируется на стандарте C99, но имеет ряд специфических и очень важных изменений. Об этом будет (я очень надеюсь) отдельная статья.

Пока базовая информация:

  1. Код который, будет дергаться с хостовой части, для исполнения, должен начинаться с ключевого слова __kernel;
  2. Функция с ключевым словом __kernel всегда возвращает void;
  3. Существуют квалификаторы типов памяти: __global, __local, __constant, __private, которые будут определять, в какой памяти будут храниться переменные. Если квалификатора перед переменной нет, то она является __private;
  4. «Общение» между хостом и kernel будет через параметры kernel. Чтобы kernel мог что-то передать хосту через параметр, параметр должен быть с квалификатором __global (пока будем использовать только __global);
  5. Код kernel принято хранить в файле с расширением cl. Но по сути подобный код может генерироваться и на лету. Это позволяет обойти некоторые ограничения. Но об этом в другой раз 🙂

Простейший пример kernel приведен ниже: __kernel void test(__global int* message){ // получаем текущий id. int gid = get_global_id(0); message[gid] += gid;}

Что делает данный код. Первое — получает глобальный id work-item который сейчас выполняется. Work-item — это то что и выполняет наш kernel. Так как мы имеем дела с параллельными вычислениями, то для каждого work-item создается свой kernel который ничего не знает о других. И никто не может гарантировать в каком порядке все work-item отработают. Но об этом подробней будет в отдельной статье (уже утал это повторять). В нашем примере это по сути индекс элемента в массиве, потому что мы будем каждый элемент массива обрабатывать в отдельном work-item. Думаю вторую строчку строчку в kernel комментировать излишни 🙂

Формируем kernel

Следующий шаг скомпилировать, то что лежит в файле *.cl. Делается это следующим образом: cl_program program = NULL;cl_kernel kernel = NULL; FILE *fp;const char fileName[] = “../forTest.

cl”;size_t source_size;char *source_str;int i; try { fp = fopen(fileName, “r”); if (!fp) { fprintf(stderr, “Failed to load kernel.

“); exit(1); } source_str = (char *)malloc(MAX_SOURCE_SIZE); source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp); fclose(fp);} catch (int a) { printf(“%f”, a);} /* создать бинарник из кода программы */program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret); /* скомпилировать программу */ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); /* создать кернел */kernel = clCreateKernel(program, “test”, &ret);
Типы cl_program и cl_kernel определены в cl.h. Сам сценарий довольно прост — загружаем файл, создаем бинарник (clCreateProgramWithSource) и компилируем. Если переменная ret по прежнему содержит 0, то вы все сделали правильно. И останется только создать сам kernel. Важно, чтобы имя передаваемое в команду clCreateKernel, совпадало с именем kernel в файле cl. В нашем случае это «test».

Параметры

Я уже упоминал, что «общения» kernel с хостом происходит за счет записи/чтения в параметры, которые передаются в kernel. В нашем случае это параметр message. Параметры, которые позволяют вот так общаться хосту с kernel, называются буферами(buffer).

Давайте создадим такой буфер на стороне хоста и передадим в kernel через API: cl_mem memobj = NULL;int memLenth = 10;cl_int* mem = (cl_int *)malloc(sizeof(cl_int) * memLenth); /* создать буфер */memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, memLenth * sizeof(cl_int), NULL, &ret); /* записать данные в буфер */ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, memLenth * sizeof(cl_int), mem, 0, NULL, NULL); /* устанавливаем параметр */ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);

Важно отметить константу CL_MEM_READ_WRITE, она означает, что мы у нас есть права для буфера на чтение и запись, на стороне kernel. Так же могут быть использованы константы типа CL_MEM_WRITE_ONLY, CL_MEM_READ_ONLY и др. Так же в методе clSetKernelArg, важен второй аргумент, он содержит индекс параметра. В данном случае 0, так как параметр message идет первым в сигнатуре kernel. Если бы он шел вторым, то мы бы написали:

/* устанавливаем параметр */ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobj);

clEnqueueWriteBuffer записывает данные из массива mem в буфер memobj.

Ну что в целом все готово. Осталось только выполнить kernel.

Исполняем kernel

Погнали, отправляем код на исполнение: size_t global_work_size[1] = { 10 }; /* выполнить кернел */ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); /* считать данные из буфера */ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, memLenth * sizeof(float), mem, 0, NULL, NULL);

global_work_size содержит число work-item которые будут созданы. Я уже говорил, что на обработку каждого элемента массива у нас будет свой work-item. Элементов в массиве у нас 10, следовательно work-item содержит 10. clEnqueueNDRangeKernel особых вопросов порождать не должна — просто запускает указанный kernel заданное число раз. clEnqueueReadBuffer считывает данные из буфера с именем memobj и помещает данные в массив mem. Данные в mem и есть наш результат!

Итоги и выводы

Источник: https://habr.com/ru/post/261323/

Поделиться:
Нет комментариев

    Добавить комментарий

    Ваш e-mail не будет опубликован. Все поля обязательны для заполнения.