Дополнение: отображение разреженных матриц
Кстати, для разреженных матриц предложены способы их отображения, что позволяет в коллекциях таких матриц (см. например, The University of Florida Sparse Matrix Collection,
http://www.cise.ufl.edu/research/sparse/matrices/
) использовать красивые "изображения" разреженных матриц, встречавшихся в реальных приложениях.
28
Библиотека Boost — на примере решения обыкновенных дифференциальных уравнений Говоря о различных библиотеках для
C++, невозможно обойти молчанием библиотеку, которая с вычислениями на графических картах пока пересекается довольно слабо, но, тем не менее, заслуживает внимательного рассмотрения. Это библиотека
Boost. Она чрезвычайно объёмна, состоит из большого числа компонентов/разделов и охватывает очень многие потребности
(линейная алгебра, геометрия, работа с графами, обработка изображений, псевдослучайные числа, регулярные выражения, многопоточность и пр.). Влияние этой библиотеки на сам язык
C++ так велико, что часть её включена в новый стандарт языка
C++ (так называемый
C++11).
Один из её разделов, где задействуется
GPU, — это решение обыкновенных дифференциальных уравнений:
odeint.
Хотя, вообще говоря, в разных своих разделах библиотека
Boost оказывается либо библиотекой с чисто заголовочными файлами, либо может ещё включать файлы статических (а иногда — и динамических библиотек), в нашем случае (т.е., раздела
odeint) она
является чисто заголовочной, поэтому для компиляции исходного кода, использующего её, достаточно добавить в путь поиска включаемых файлов значение
$(BOOST_ROOT)
. (Здесь предполагается, что эта библиотека установлена в каталоге, на который указывает переменная окружения
BOOST_ROOT
, если такой переменной нет, можно просто использовать вместо неё путь к этому каталогу).
Поскольку библиотека
Boost написана на весьма «продвинутом»
C++, компиляция примеров применения
odeint в
Visual Studio 2008 довольно проблематична. Камнем преткновения является, конечно же, не совсем полная поддержка компилятором новых особенностей языка
C++.
Компилируемый в
VS2008 пример — это файл
phase_osc_chain.cu, поэтому разберём применение
Boost на нём. Файлы
thrust_lorenz_ensemble.cu,
lorenz_parameters.cu,
phase_oscillator_ensemble.cu имеет смысл пробовать с другими компиляторами, например
g++ или более новой версией
Visual Studio.
Следует обратить внимание на то, что автор программы, вероятно, компилировал её под
Linux, поскольку в
VS2008 применённая им функция получения случайных величин drand48()
— отсутствует. По этой причине её пришлось заменить на приблизительный аналог: double(rand())/RAND_MAX
Поскольку вычисления должны проводиться на
GPU, из библиотеки
Thrust подключается файл
thrust/device_vector.h, а также определения итераторов. Из библиотеки
Boost подключаются заголовочные файлы раздела
odeint. Последние содержат довольно глубоко вложенные пространства имён, поэтому для них, а также для пространства имён std используются директивы using
На основе вектора на графическом устройстве (
thrust::device_vector<>
) определены синоним для векторов из value_type
(называется state_type
) и синоним для векторов из size_t
(фактически это вектор целых величин; будет называться index_vector_type
).
Помимо главной функции и задания нескольких констант в программе определяется класс (новый тип) объектов phase_oscillators
7 29
В нём определён также функтор типа sys_functor
— аналог функции, которая должна вызываться при переборе осцилляторов для каждого отдельного осциллятора. Определение функтора содержит шаблонный параметр — тип набора
Tuple
(обратите внимание, что
функтор может существовать и как __host__
-объект, и как
__device__
-объект).
В главной функции создаются вектор начальных условий и вектор частот (оба — одинаковой размерности
N
), оба располагаются в памяти
CPU (т.к. это вектора из
STL); вектор начальных условий заполнен случайными значениями фазы, а вектор частот — линейно убывающими частотами. Эти векторы переносятся в память графической карты и с помощью пошагового вычислителя типа Рунге-Кутты начинается интегрирование с постоянным шагом dt от
0
до
10
По завершении интегрирования полученные значения фаз с помощью thrust::copy()
копируются в стандартный вывод — по одному значению на строке.
Для более подробного ознакомления с
odeint можно посмотреть статьи:
odeint v2 — Solving ordinary differential equations in C++ http://www.codeproject.com/Articles/268589/ Solving ordinary differential equations with OpenCL in C++ http://www.codeproject.com/Articles/429183/ Boost.Compute Кроме решения обыкновенных дифференциальных уравнений в составе
Boost может появиться ещё один раздел: библиотека, известная уже некоторое время под именем
Boost.Compute.
Boost.Compute — GPU/parallel-computing library for C++ based on OpenCL https://github.com/boostorg/compute Документация к библиотеке Boost.Compute http://boostorg.github.io/compute/ Библиотека
Boost.Compute реализует
C++-интерфейс к многоядерным
CPU и
GPGPU-платформам на основе
OpenCL, по существу являясь «обёрткой» над вызовами библиотеки
OpenCL, обеспечивающей доступ к вычислительным устройствам, контекстам, командным очередям, буферам памяти (подробнее об
OpenCL — далее).
Библиотека
Boost.Compute имеет
STL-подобный программный интерфейс, предусматривающий общие алгоритмы (
transform(), accumulate(), sort()
) и контейнеры (
vector
, flat_set
), вводит параллельные вычислительные алгоритмы (
exclusive_scan(), scatter(), reduce()
),
т.н. "воображаемые" итераторы (
transform_iterator<>, permutation_iterator<>, zip_iterator<> и пр.).
Интересно также познакомиться с библиотекой VexCL: (
https://github.com/ddemidov/vexcl
)
VexCL: Vector expression template library for OpenCL
http://www.codeproject.com/Articles/415058/
Это C++-библиотека, генерирующая ядра OpenCL/CUDA на основе векторных выражений.
30
Стандарт OpenCL OpenCL — это открытый стандарт для параллельного программирования и работы с широким набором современных параллельных вычислителей (многоядерных процессоров,
GPU,
FPGA).
Расшифровывается это название как
Open Computing Language (что-то вроде:
Открытый язык для вычислений). Изначально
OpenCL был предложен фирмой
Apple, но впоследствии получил поддержку многих представителей отрасли, в том числе
Intel,
AMD,
IBM,
NVIDIA,
ARM,
Samsung и др. Ожидается, что каждый изготовитель процессоров реализует доступ к своим вычислительным ресурсам с помощью собственной
OpenCL-библиотеки. Есть такая поддержка и в рамках
CUDA (см. заголовочные файлы в подкаталоге
CL/ каталога с заголовочными файлами).
Предполагается (в идеале), что предоставляемая библиотека
OpenCL может использовать все доступные ресурсы (
GPU,
CPU) параллельно, её программная модель основана на
C и максимально абстрагирована от конкретной реализации вычислительных устройств. Она может опрашивать и
выбирать имеющиеся вычислительные ресурсы, инициализировать их, создавать так называемые вычислительные контексты и рабочие очереди. Кроме того, она может компилировать и создавать программы (они тоже называются ядрами), которые затем исполняются на этих ресурсах. Для описания ядер используется
C99-подмножество языка
C с некоторыми расширениями.
Параллельность вычислений обеспечивается независимой работой отдельных вычислителей (или рабочих единиц, в терминологии
OpenCL —
work-item), которые могут быть объединены в рабочие группы (
work-group); полное число вычислителей, работающих параллельно, называется
global work size. Эти вычислители могут взаимодействовать друг с другом, их работа может быть синхронизирована в рабочей группе для координации доступа к памяти. Нетрудно заметить, что понятие рабочей группы в
OpenCL соответствует понятию блока нитей в
CUDA, а т.н. глобальный размер задаёт аналог сетки блоков из
CUDA.
Работа программы, использующей
OpenCL, протекает примерно так: опрашиваются имеющиеся вычислительные ресурсы, выбираются те, что будут далее использоваться, вычислительные ядра создаются из исходного кода и распределяются для запуска по вычислительным ресурсам.
Таким образом, разработка
OpenCL-программы сводится к написанию ядер и
host-приложения для
PC, которое распределяет нужные ядра по доступным устройствам. Такое приложение должно использовать пять структур: cl_device_id
, cl_kernel
, cl_program
, cl_command_queue
, cl_context
. Оно распределяет ядра (
cl_kernel
), полученные из исходного кода (
cl_program
) по устройствам (
cl_device_id
), эти ядра попадают устройствам через очередь команд (
cl_command_queue
); контекст (
cl_context
) позволяет устройствам получать ядра и обмениваться данными.
Более конкретно, подобное приложение (в качестве примера будем иметь в виду простую программу вроде
hello.c, но не полностью совпадающую с ней) должно получить данные об устройстве, которое будет далее исполнять функцию-ядро, например, как-нибудь так: clGetPlatformIDs(1, &platform, NULL); //
перваяобнаруженнаяплатформаclGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); //
первоеу-во (GPU!) Далее приложение создаёт контекст, например, только с одним обнаруженным устройством: context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
31
В последнем вызове и ряде последующих вызовов функциям передаётся адрес переменной err
, в которую записывается код возможной ошибки при выполнении; реальная программа должна анализировать этот код, используя его значение для вывода сообщений об ошибках. В приводимых здесь строках кода — для простоты — этого не делается.
После этого приложение должно получить программу из кода ядра, содержащегося, например, в файле
hello.cl, для чего содержимое этого файла читается в массив, передаваемый функции clCreateProgramWithSource
: program = clCreateProgramWithSource(context, 1,
(const char**)&program_buffer, &program_size, &err); clBuildProgram(program, 0, &device, NULL, NULL, NULL);
Отсутствующие в последнем вызове параметры могут определять варианты компиляции. После неё из заданной функции создаётся ядро (здесь строка "hello"
— название функции ядра): kernel = clCreateKernel(program, "hello", &err);
Для распределения ядер по устройствам необходимо создавать очереди к устройствам: queue = clCreateCommandQueue(context, device, 0, &err);
Поскольку
ядру понадобится память для вывода, необходимо также создать буфер памяти: mem = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(char), NULL, &ret);
Теперь, когда все компоненты окружения (т.е., структуры cl_device_id
, cl_kernel
, cl_program
, cl_command_queue
, cl_context
) созданы, следует подготовить ядру необходимые параметры вызова, например, в случае ядра hello
— это один параметр (адрес буфера памяти для вывода): ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&mem); и можно отправлять ядро в очередь на исполнение: global_size = 8; local_size = 4; clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size,
&local_size, 0, NULL, NULL);
Последняя функция (как и в случае
CUDA) не только обеспечивает запуск ядра на устройстве, но и определяет, как много рабочих единиц должно быть создано (параметр global_size
), а также сколько рабочих единиц будет в рабочей группе (параметр local_size
).
Для чтения полученных результатов из буфера памяти mem в массив (в данном случае — символов)
string вызывается функция clEnqueueReadBuffer
: ret = clEnqueueReadBuffer(queue, mem, CL_TRUE, 0,
MEM_SIZE * sizeof(char), string, 0, NULL, NULL);
Она возвращает значение кода возможной ошибки (или значение
CL_SUCCESS в случае успешного завершения).
32
Синтаксис оформления текста ядер в OpenCL включает в себя спецификаторы: функции ядра
(__kernel
), памяти (
__global
,
__local
) и некоторые другие, встречающиеся более редко.
Для того, чтобы каждый исполняемый экземпляр кода ядра мог иметь доступ к «своим» данным и вообще «знать» параметры конфигурации, используются вызовы специальных функций:
OpenCL
CUDA
get_local_size(0)
// 1, 2
Размер рабочей группы
(work-group)
blockDim.x
// y, z
Размер блока нитей
(thread block) get_group_id(0)
Номер рабочей группы blockIdx.x
Номер блока нитей get_num_groups(0)
Число рабочих групп gridDim.x
Число блоков get_local_id(0)
Номер элемента (work-item)
threadIdx.x
Номер нити в блоке get_global_id(0)
Глобальный номер эл-та blockDim.x*blockIdx.x
+threadIdx.x
Глобальный номер нити get_global_size(0)
Глобальный размер blockDim.x*gridDim.x
Глобальный размер
Более детальную информацию по спецификации OpenCL можно получить в документе "The
OpenCL Specification" (скажем, файл
opencl-1.2.pdf
), хотя не факт, что доступная версия
OpenCL не окажется более ранней, например, 1.1... На момент написания пособия уже одобрена версия 2.0 спецификации.
Из приведённого выше краткого описания становится понятно, что самое главное преимущество
OpenCL — переносимость. Ядра OpenCL могут выполняться на различных GPU и CPU самых разных изготовителей: Intel, AMD, Nvidia, IBM и др., а также на устройствах программируемой логики
(FPGA), причём одно и то же приложение может распределять выполнение ядер на многих таких устройствах одновременно. Некоторым недостатком является необходимость освоения нового подхода со своим набором функций. Избежать этого можно применением какой-нибудь
«обёртки» (C++-библиотеки), например, VexCL.
В качестве первоначального тестового примера работы с OpenCL предлагаются файлы
hello.c
,
hello.cl
. Первый файл — программа, реализующая почти все описанные выше шаги, необходимые OpenCL, а второй из этих файлов — исходный текст довольно примитивного тестирующего ядра. Их можно поместить в каталог обычного C-проекта (даже не CUDA) и в свойства проекта добавить пути поиска заголовочных файлов
$(CUDA_INC_PATH)
, а также пути поиска библиотек
$(CUDA_LIB_PATH)
и саму нужную библиотеку
OpenCL.lib
Если есть желание попрактиковаться в доведении OpenCL-программ до работоспособного состояния в системе Windows, то можно взять весьма интересные примеры с сайта
http://www.caam.rice.edu/timwar/HPC12/
(Tim Warburton, профессор факультета
Computational and Applied Mathematics университета Rice). С одной стороны, эти примеры написаны на C++, причём довольно изобретательно, с другой стороны, они точно работоспособны, правда, под Windows — далеко не сразу.
Тексты с упомянутого сайта написаны, скорее всего, под Mac, поэтому для компиляции их под
Windows (конкретно, файла clhelper.hpp с помощью VS2008) придётся кое-что подправить. Во-первых, лучше сразу удалить подключение заголовочного файла
unistd.h
, поскольку под Windows его всё равно нет, а в программе он (возможно, и некоторые другие) не используется. Во-вторых, макроопределение
CL_CHECK_ERR
, возвращающее значение некоторого заранее неизвестного типа "из-под" пар круглых, а затем фигурных скобок, слишком «авангардно» для компилятора VC 9.0 и не совсем ясно, чем его можно было бы заменить, поэтому проще его не применять вообще. Это касается вызовов функций clCreateContext()
, clCreateCommandQueue()
в конструкторе clhelper и функции clCreateBuffer()
в методе createbuffer
. При этом нужно добавить определение используемой при первых двух вызовах величины
_err
, поскольку в самом конструкторе
33
этого определения — после игнорирования макроса — уже нет. В-третьих, лучше всего объявление и определение функции обратного вызова pfn_notify снабдить пометкой
CL_CALLBACK
, несущественной везде, кроме
Windows, — иначе придётся заменять четвёртый параметр в вызове clCreateContext()
на
NULL
, лишаясь возможности узнать о каких-либо неприятностях при создании контекста
OpenCL.
Ну, и напоследок надо сказать, что для системы
Windows код файла clhelper.hpp
, являющийся работоспособным во всех системах, кроме неё, должен быть скорректирован в методе load_program_source()
класса clfunction
, поскольку без этого возникает трудноуловимая ошибка компиляции ядра, которую невозможно обнаружить даже
«отдельным» компилятором ядер (например,
clcc; см. далее). А всё дело оказывается в том, что при чтении файла в память используется открытие файла в режиме "r"
, что автоматически приводит к «проглатыванию» первого из символов пары
0x0D,0x0A
, завершающих строки текста, из-за чего только (!) исходные тексты ядер, созданные под
Windows, оказываются короче длины файла. Конец текста ядра в памяти при этом оказывается испорчен посторонними символами и ядро просто не компилируется. Причём — это надо подчеркнуть
— никаких ошибок не будет, если ядро содержит только
Скачать 1.54 Mb.