В.А.Антонюк Программирование на видеокартах (GPGPU) Спецкурс кафедры ММИ Москва Физический факультет МГУ им.М.В.Ломоносова Антонюк Валерий Алексеевич Программирование на видеокартах (GPGPU). Спецкурс кафедры ММИ. – М.: Физический факультет МГУ им. М.В. Ломоносова, 2015. – 48 с. Учебное пособие представляет собой краткий конспект вводного курса в программирование с помощью графических карт, прочитанного студентам кафедры (ранее она называлась кафедрой компьютерных методов физики) в 2013-2015 годах. Фактически речь идёт о возможностях и способах программирования графических процессоров, получивших широкое распространение в последние годы. Рассмотрены средства, предлагаемые известнейшим производителем графических карт, – платформа CUDA фирмы NVIDIA: компилятор nvcc, его симбиоз со средой разработки Visual Studio, предложенные расширения языка C/C++ (дополнительные типы данных, ключевые слова, новые математические функции, синтаксис вызова программ, исполняемых на графической карте), а также особенности параллельного исполнения кода вообще, разделения общих ресурсов, синхронизации исполнения отдельных потоков программ. Анализируются параллельные реализации: базовых векторных алгоритмов, умножения матриц, сортировки (bitonic sort). Рассматривается использование предлагаемых фирмой NVIDIA библиотек: cuBLAS (базовые функции линейной алгебры), cuSPARSE (работа с разреженными матрицами), Thrust (объектно-ориентированный программный интерфейс). Альтернативой привязки к конкретному производителю является использование независимых решений, работающих на более широком спектре устройств различных производителей, поэтому излагаются сведения о работе с библиотекой Boost, платформой OpenCL, библиотекой OpenCV. Для отображения результатов может быть выбран кроссплатформенный интерфейс OpenGL, поэтому излагаются основы его применения, тем более, что в рамках OpenGL теперь тоже существует возможность задействовать вычислительную мощь современных графических процессоров (язык GLSL) для отрисовки сложных трёхмерных сцен. Не следует недооценивать и проникновения современных технологий в самое популярное приложение последнего времени – браузер: даже в нём теперь (с помощью WebGL) можно создавать сложные вычислительные и графические программы, существенно ускоряемые наличием графического процессора; читатели смогут познакомиться с наиболее впечатляющими примерами таких программ. Рассчитано на студентов старших курсов физического факультета и аспирантов, но может быть использовано всеми заинтересованными студентами, аспирантами и сотрудниками для первоначального знакомства с современными вычислительными возможностями. Автор – доцент кафедры математического моделирования и информатики (ММИ) физического факультета МГУ. Рецензент: доцент кафедры физической электроники И.К.Гайнуллин. . Объем 3,0 п.л. экз. Заказ № Физический факультет им. М.В.Ломоносова, 119991 Москва, ГСП-1, Ленинские горы, д.1, . Отпечатано в отделе оперативной печати физического факультета МГУ. Физический факультет МГУ им. М.В. Ломоносова, 2015 В.А.Антонюк, 2015 Программирование на видеокартах (GPGPU) Название этого спецкурса может показаться странным. Программирование? Это понятно. Но на видеокартах... Что такого есть в видеокартах, что нам захотелось бы их программировать? И зачем нам изучать, как это сделать? Оказывается, в своём развитии видеокарты не так давно превратились из специализированных устройств для воспроизведения графической информации в универсальные вычислительные устройства, причём с параллельной архитектурой: они способны не только исполнять задаваемую программу, но и выполнять операции одновременно над целым массивом данных. При этом функция воспроизведения никуда не исчезла, просто она стала далеко не основной, на передний план вышли функции формирования и обработки визуальной и любой другой информации. И теперь нам — обыкновенным пользователям программируемых устройств — предстоит осознать, как можно воспользоваться резко возросшими вычислительными возможностями наших устройств, и научиться это делать. Если, конечно, у нас есть вычислительно-трудоёмкие задачи, которые хотелось бы решать быстрее, и они потенциально могут быть (хотя бы частично) распараллелены. В англоязычной литературе даже появился специальный термин для такого рода программирования и вычислений: GPGPU, что принято расшифровывать как General-Purpose computation on Graphics Processing Units ( Мы же используем здесь неформальное выражение "программирование на видеокартах". Теперь о самом спецкурсе. Условно его можно разделить на две части: "прямое" программирование видеокарт, предполагающее некоторое знакомство с их аппаратными свойствами и особенностями, и "косвенное" их программирование, когда используются специальные библиотеки, а пользователю вообще ничего не нужно знать об устройстве видеокарты. В первом случае речь идёт об освоении средств, предлагаемых одним из производителей видеокарт — фирмой NVIDIA. Как известно, для программирования своих графических карт (или обобщённо — GPU) она использует расширения языка C и специальный компилятор nvcc, "понимающий" эти расширения. Задействуется он посредством визуальной среды с компилятором стандартного языка C/C++, например Visual Studio 2008 или 2010, после установки пакета CUDA Toolkit, реализующего так называемую платформу параллельных вычислений CUDA. В рамках спецкурса рассматриваются эти расширения языка C в CUDA: квалификаторы функций (теперь в наших программах будет два разных типа функций, поскольку одни должны исполняться компьютером, т.е., CPU, а другие — аппаратурой видеокарты, т.е., GPU), виды памяти (так как GPU имеет собственную память, причём нескольких видов), а также дополнительные типы данных и математические функции (поскольку данные во время вычислений располагаются в GPU и обрабатываются с помощью GPU). Обсуждаются передача данных между CPU и GPU (в терминологии CUDA — между Host и Device), оформление кода, исполняемого на GPU (так называемых ядер, kernels), передаваемые ядрам параметры, особенности параллельного исполнения кода аппаратными средствами (нити, блоки), включая взаимодействие между нитями и синхронизацию их действий, приводятся примеры простых ядер: для суммирования двух векторов, для вычисления суммы компонент вектора. Подробно разбираются некоторые основополагающие алгоритмы: умножение матриц ("наивная" реализация и блочная, с использованием разделяемой памяти GPU), параллельная сортировка Bitonic Sort. Разумеется, совсем не обязательно детально знакомиться с внутренним устройством GPU, можно воспользоваться более высокоуровневыми библиотеками, тем не менее использующими GPU (и являющимися частью CUDA): cuBLAS (базовые функции линейной алгебры), cuSPARSE (работа с разреженными матрицами), cuFFT (БПФ), cuRAND (случайные величины), Thrust (С++-библиотека c шаблонами). И слушатели спецкурса смогут получить некоторый опыт работы с этими библиотеками с помощью разбираемых примеров программ. Конечно, при этом не следует забывать, что программы, написанные с использованием CUDA, будут работать только на видеокартах NVIDIA, но не других производителей (например, Intel или AMD). С этой точки зрения гораздо более оправданным было бы создание программ, которые смогут работать на видеокартах не только какого-то одного производителя, но и других изготовителей. Для этого можно использовать "программную прослойку" между нашим кодом и самой видеокартой, например, в виде какой-то готовой библиотеки. И такие решения уже есть. Слушатели спецкурса могут поближе познакомиться с ними и попрактиковаться в написании соответствующих программ. Прежде всего, это открытый стандарт для работы с параллельными вычислителями — OpenCL. Каждый изготовитель видеокарт или просто GPU может реализовать (и часто реализует!) интерфейс, предлагаемый OpenCL, что позволяет использовать один и тот же код на видеокартах или GPU разных производителей (сейчас это возможно для видеокарт NVIDIA, Intel, AMD, а также для последних вариантов GPU: Adreno, Mali, Vivante и PowerVR — т.е., на планшетах и даже на некоторых мобильных телефонах). Кроме этого, наиболее солидные библиотеки имеют в своём составе программные средства для работы с GPU. Это и OpenCV (Open Source Computer Vision Library), где часть алгоритмов работы с изображениями имеет варианты как для CPU, так и для GPU, и библиотека Boost, где интерфейс С++-шаблонов применён для решения обыкновенных дифференциальных уравнений. Полноценная работа едва ли возможна без графического отображения информации, поэтому слушатели познакомятся с графической библиотекой и программным интерфейсом OpenGL — как на простейших примерах, так и на демонстрационных программах, входящих в состав CUDA SDK; узнают о существовании ещё одного языка в рамках OpenGL: GLSL (OpenGL Shading Language), применяемого для написания кода GPU, и о том, что OpenGL вместе с GLSL уже "ждут" нас почти в каждом браузере (в том числе и на новых мобильных телефонах!) — в виде WebGL, программного интерфейса к GPU в рамках браузера. Завершающим "десертом" спецкурса будут такие приятные мелочи в работе программиста, как визуализаторы: небольшие фрагменты специализированного кода-описания, позволяющие сделать отладку программ с шаблонными библиотеками или просто сложно организованными данными в Visual Studio более комфортной. Видеокарты и графические процессоры как супервычислители Окружающий нас мир переполнен созданными человеком электронными устройствами. И заметная доля их обладает серьёзными вычислительными возможностями, причём такими, что совсем недавно казались поразительными. Если раньше видеокарта настольного компьютера была просто устройством воспроизведения визуальной информации, то сейчас вполне может оказаться, что имеющаяся в нашем распоряжении видеокарта заметно мощнее в вычислительном смысле, чем центральный (пусть даже и многоядерный) процессор нашего компьютера. Если раньше мобильный телефон был просто миниатюрным вариантом (радио-) телефонной трубки, то сейчас это чаще всего небольшой, но мощный компьютер, содержащий многоядерный графический процессор и при этом «умеющий» совершать телефонные звонки. А ведь есть ещё масса других, в каком-то смысле «промежуточных» между ними устройств: ноутбуки, нетбуки, планшеты, неттопы и т.п. И все они в последнее время тоже обладают солидными вычислительными возможностями. Разумеется, хорошо было бы как-то использовать их. Но что для этого понадобится сделать, что нужно иметь и чего можно достичь — вот те вопросы, которые тогда встают перед нами. Для осознания возможностей графических процессоров полезно ознакомиться с некоторыми демонстрационными программами, «умеющими» эффективно использовать видеокарту. Всё, что необходимо для того, чтобы они могли работать, — это графический процессор на видеокарте и драйвер к ней. Обратите также внимание на то, что некоторые демонстрации (это касается тех из них, что запускаются в рамках браузера) вполне работоспособны даже на современных мобильных телефонах… Синтезированная шахматная доска (автор Thanassis Tsiodras) — традиционное приложение для операционной системы Windows, в котором осуществляется отрисовка трёхмерной сцены: доски с шахматными фигурами на ней (https:// причём и доска, и фигуры имеют почти зеркальную поверхность, так что все они (многократно) отражаются друг в друге. В качестве ускорителя вычислений предполагается видеокарта архитектуры CUDA, при этом синтез картинки осуществляется в реальном времени с приличной скоростью (в зависимости от возможностей видеокарты — от единиц до десятков кадров в секунду), что позволяет наблюдателю интерактивно управлять вращением доски и углом зрения на неё. Если нам необходимо разрабатывать подобные приложения, то понадобится компилятор C/C++ и комплект разработки, именуемый CUDA Toolkit; также стоит установить в дополнение к нему и CUDA SDK. Имеющиеся там многочисленные примеры весьма поучительны и полезны — хотя бы потому, что для них приведён исходный код, — что позволяет использовать их как основу для собственных программ. Эффектные демонстрации существуют и в рамках браузеров (на настольном компьютере или даже на мобильном телефоне), поскольку браузер — пример часто используемого приложения, которое ощутимо выигрывает от возможного ускорения синтезируемого графического вывода. Пожалуй, самая известная демонстрация такого рода — это WebGL Water (автор — Evan Wallace). Большое число аналогичных примеров можно посмотреть также на сайте babylonjs.com. http:// В настоящее время основными изготовителями видеокарт являются фирмы NVIDIA, Intel, AMD. Каждая из них выпускает также и программные средства поддержки своих изделий. Кроме того, выпускается широкий ряд графических процессоров (GPU) для планшетов и мобильных телефонов под брэндами Adreno, Mali, Vivante, PowerVR. Для поддержки разработчиков изготовители этих GPU тоже выпускают специализированные средства разработки программ. Понятно, что в рамках одного спецкурса невозможно охватить всё многообразие существующих (и постоянно появляющихся) вариантов графических процессоров и параллельных вычислителей, поэтому здесь мы ограничимся, с одной стороны, знакомством с программными средствами лишь одного (но практически доминирующего) производителя: CUDA от NVIDIA, а, с другой стороны, постараемся в программировании графических процессоров уйти от привязки к конкретному производителю, разбирая возможные альтернативные решения (OpenCL, OpenCV, Boost, GLSL). От слушателей спецкурса ожидается определённый уровень владения языками C и C++, поскольку большая часть разбираемых примеров использует именно эти языки. Необходимым является также понимание программирования с применением шаблонов — по той простой причине, что рассматриваемые библиотеки их интенсивно используют. В остальном никаких специальных знаний от слушателей не требуется, хотя практика работы с проектами Visual Studio – приветствуется. CUDA – архитектура и программные средства В рамках CUDA (Compute Unified Device Architecture) под Windows написание программ происходит в среде Visual Studio (в нашем случае — версии 2008 года), а их компиляция осуществляется как с помощью компилятора из Visual Studio, так и с помощью специального компилятора фирмы NVIDIA, называемого nvcc. Именно последний начинает разбор любой CUDA-программы, «отделяет» нестандартные дополнения к языку C/C++ и преобразует их (совместно с драйвером графической карты) в программу для процессора этой карты (далее обобщённо именуемого GPU). Компилятору из Visual Studio остаётся для преобразования традиционная часть программы, исполняемая на основном процессоре (CPU). Необходимость применения двух компиляторов достаточно очевидна: вычислительные средства графической карты (GPU) обладают совершенно иной системой команд, о которой стандартный C/C++-компилятор ничего «не знает» и знать не может. По этой причине каждый изготовитель GPU должен создать свой собственный компилятор, «умеющий» «строить» программы для GPU, и драйвер, чтобы загружать их туда для исполнения. В основе высокой производительности графических процессоров (GPU) лежит их способность выполнять программный код параллельно на большом числе одинаковых процессоров (вычислительных ядер) с помощью некоторого количества одновременно исполняемых потоков. Подобная архитектура в литературе называется SIMT — Single Instruction Multiple Threads: код чаще всего предполагается одинаковым, а обрабатываемые данные — разными, что, разумеется, применимо далеко не для всех алгоритмов. К счастью, многие часто используемые алгоритмы в той или иной степени допускают подобного рода распараллеливание. CUDA-программа В программной модели, предлагаемой в рамках CUDA, программист может определять специальные функции (тоже называемые обычно ядрами — kernel), которые затем исполняются параллельно многими различными потоками исполнения (thread). Кстати, несмотря на то, что за понятием thread в нашей литературе закрепился перевод поток, здесь (и в переводной литературе по CUDA) применяется перевод нить, поскольку в рамках CUDA понятие поток (stream) тоже используется. Функции, исполняемые многими потоками-нитями, помечаются в программе спецификаторами __global__ и могут быть вызваны на исполнение с помощью специального нового синтаксиса <<<...>>>: // определение функции-ядра __global__ void MyKernel(...) . . . . . . // вызов ядра на исполнение MyKernel<<>>(...); . . . Легко видеть, что вызов ядер на исполнение инициируется в программе упоминанием их имени с так называемыми параметрами конфигурации запуска ядер (которые указываются в тройных угловых скобках <<<...>>>) наряду с обычными параметрами, необходимыми для работы каждого отдельного ядра. Таким образом, в рамках CUDA-программы в реальности «сосуществуют» два типа кода: для CPU (в терминологии CUDA именуемого host) и для GPU (в терминологии CUDA — device), причём оба типа исходного кода используют язык C/C++, только в device-коде пока употребляется некоторое подмножество C99, а вот host-часть программного кода вполне может быть написана на C++. Подобный симбиоз позволяет программисту оставаться в привычной среде одного языка, правда, «разбавленного» некоторыми дополнительными ключевыми словами и типами величин. Основная работа по анализу такой «смеси» ложится на компилятор nvcc. Иерархия нитей Вызов на исполнение отдельных копий кода ядра сопровождается указанием, сколько нужно запустить экземпляров кода (фактически — нитей исполнения), а также как они должны быть организованы для решения конкретной задачи, поскольку эти нити группируются в так называемые блоки (blocks), а те, в свою очередь, формируют сетку блоков (grid). В зависимости от выбранной при запуске конфигурации все нити, исполняющие код ядра, имеют одно-, двух- или трёхмерную организацию (как правило, соответствующую характеру решаемой задачи), причём каждая нить «имеет информацию» о том, где она расположена в выбранной конфигурации. Виды памяти Для работы программного кода на графическом процессоре видеокарты необходима память, доступная этому процессору, поэтому в host-коде понадобится выполнять дополнительные операции по выделению device-памяти, копированию данных с host на device, а после обработки данных — возвращение результата с device на host. Чаще всего используется так называемая глобальная память графической карты, однако нити в пределах одного и того же блока могут использовать и так называемую разделяемую память с быстрым доступом. Кроме этого возможна работа с так называемой константной памятью, доступной всем нитям сетки, — но только на чтение. Компиляция программ При установке пакета CUDA создаются специальные переменные окружения: базовая (с именем CUDA_PATH) и переменная конкретной версии (например, CUDA_PATH_V4_1) — видимо, для возможности одновременного использования нескольких установленных версий CUDA; на их основе определяются переменные окружения, которые помогают указывать в Visual Studio пути поиска включаемых заголовочных файлов (CUDA_INC_PATH=%CUDA_PATH%include), статических библиотек (CUDA_LIB_PATH=%CUDA_PATH%lib\Win32), а также исполняемых файлов и динамических библиотек (CUDA_BIN_PATH=%CUDA_PATH%bin). Значения этих (да и любых других) переменных окружения могут быть использованы в свойствах проекта Visual Studio для успешной компиляции программ наряду с указанием конкретных путей к файлам, для этого имя переменной окружения заключается в круглые скобки и предваряется символом доллара, например: $(CUDA_INC_PATH). По сути, можно сказать, что в рамках проекта Visual Studio заданы макроопределения с именами переменных окружения. Далее иногда мы будем использовать такие макросы для обобщённого указания конкретных каталогов. Поскольку какие-то проекты могут использовать также NVIDIA GPU Computing SDK, то иногда необходимо добавлять и пути поиска включаемых файлов и библиотек из местоположения этого SDK (соответствующая переменная окружения называется NVSDKCOMPUTE_ROOT), например, указывая путь $(NVSDKCOMPUTE_ROOT)\C\common\inc для включаемых заголовочных файлов и путь $(NVSDKCOMPUTE_ROOT)\C\common\lib\Win32 для статических библиотек. Расширения языка C в архитектуре CUDA Программы в архитектуре CUDA пишутся на «расширенном» языке C, а для компиляции их используется, как уже говорилось, специальный компилятор nvcc. Расширения языка включают спецификаторы функций (они показывают, где будет исполняться функция, и откуда может быть вызвана), дополнительные типы данных, спецификаторы переменных (они определяют, какая память будет использоваться для их размещения), синтаксис запуска ядер, встроенные переменные, хранящие информацию о каждой нити, а также математические функции. Спецификаторы функций и переменных Для обозначения того, что функция является ядром (т.е., должна исполняться на GPU и может быть запущена на исполнение с CPU), используется спецификатор __global__. Спецификаторы __host__ и __device__ помечают функции, исполняемые на CPU и GPU соответственно, причём и вызваны они могут быть только оттуда, где исполняются. Для задания размещения в памяти GPU переменных употребляются спецификаторы __device__ (переменная находится в глобальной памяти и доступна всем нитям), __constant__ (размещена в так называемой константной памяти, откуда она может быть только прочитана любой из нитей), __shared__ (переменная — в разделяемой памяти, где доступна только всем нитям «своего» блока). Дополнительные типы данных CUDA Это 1/2/3/4-мерные векторы из базовых типов языка C (char, short, int, long — unsigned и signed, — longlong, float и double), имена их образуются добавлением цифры с числом компонент: char1, char2, char3, char4, uchar1, uchar2, uchar3, uchar4, short1, short2, short3, short4, ushort1, ushort2, ushort3, ushort4, int1, int2, int3, int4, uint1, uint2, uint3, uint4, long1, long2, long3, long4, ulong1, ulong2, ulong3, ulong4, float1, float2, float3, float4, longlong1, longlong2, double1, double2. Обращение к компонентам осуществляется по именам (x, y, z, w), а создание значенийвекторов — с помощью вызова функций make_, где — один из вышеприведённых типов данных, например: int2 a = make_int2(1,4); float3 u = make_float3(1,2.72f,3.14f); Для задания размерности имеется тип dim3, он обладает также нормальным конструктором: dim3 Blocks(16,16); dim3 Grid(256); // <=> Blocks(16,16); // <=> Grid(256,1,1); Полный синтаксис запуска ядер на исполнение Имея в своём распоряжении две переменные типа dim3, можно задавать предпочтительную конфигурацию запуска ядер на исполнение. В простейших случаях, когда достаточно одномерной организации данных, можно обойтись и двумя целочисленными значениями. Более сложные случаи запуска могут иметь также и дополнительные параметры; вот как выглядит полный синтаксис запуска ядер на исполнение: kernelName<<>>(args); Здесь kernelName — имя __global__-функции, Dg — переменная типа dim3 (параметры сетки блоков), Db — переменная типа dim3 (параметры блока нитей), Ns — дополнительный объём разделяемой памяти в байтах для каждого блока (необязательный параметр, он по умолчанию равен нулю), S — задаёт поток (CUDAstream), в котором должен произойти вызов (поток 0 по умолчанию), args — параметры функции-ядра (их может быть несколько, а общий размер этих параметров пока ограничен 256 байтами). Дополнительные математические функции Поскольку ядра исполняются на GPU, а не на CPU, им необходимы альтернативные реализации как минимум всех тех математических функций, которые были доступны в программах ранее в рамках стандартной математической библиотеки. Кроме double-версий в CUDA также имеются float-аналоги стандартных функций (например, для sin — sinf). Есть ещё набор функций пониженной точности, но более «быстрых» (для sinf, например, __sinf). Для функций, реализующих арифметические операции и преобразование типов, можно также задавать с помощью суффиксов их имени способ округления результата: rn (к ближайшему), rz (к нулю), ru (вверх), rd (вниз). Разделяемая (__shared__) память Основная часть динамической памяти GPU доступна как глобальная __device__-память — всем нитям и блокам, а также CPU. Несмотря на то, что пересылка из памяти CPU в память GPU осуществляется довольно быстро, пересылки в рамках динамической памяти ещё быстрее. Разделяемая (__shared__) память располагается в GPU и потому обладает гораздо большим быстродействием, чем глобальная, но доступна для чтения и записи нитям из одного блока. Первый способ выделения разделяемой памяти в ядре — явное указание размера массива со спецификатором __shared__. Второй способ — при запуске ядра задать дополнительный объём разделяемой памяти в байтах, который необходимо выделить каждому блоку при запуске ядра (третий параметр конфигурации функции-ядра, см. выше). Для доступа в ядрах к такой памяти используется описание массива в функции ядра без явно заданного размера, например: __shared__ float buf[]; При этом предполагается, что сам вызов ядра (с названием kernel) выглядит так: kernel<<< ... , ... , k*sizeof(float) >>>( ... ); Тогда каждый блок ядер будет иметь доступ к массиву из k вещественных значений в разделяемой памяти (см. реализацию алгоритма умножения матриц далее).