М. В. Ломоносова 2015 Антонюк Валерий Алексеевич Программирование на видеокартах (gpgpu). Спецкурс кафедры мми. М. Физический факультет мгу им. М. В. Ломоносова, 2015. 48 с. Учебное пособие
 Скачать 1.54 Mb.
|
|
В.А.Антонюк Программирование на видеокартах (GPGPU) Спецкурс кафедры ММИ Москва Физический факультет МГУ им.М.В.Ломоносова 2015 Антонюк Валерий Алексеевич Программирование на видеокартах (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) можно создавать сложные вычислительные и графические программы, суще- ственно ускоряемые наличием графического процессора; читатели смогут познакомиться с наиболее впечатляющими примерами таких программ. Рассчитано на студентов старших курсов физического факультета и аспирантов, но может быть использовано всеми заинтересованными студентами, аспирантами и сотрудниками для первоначального знакомства с современными вычислительными возможностями. Автор – доцент кафедры математического моделирования и информатики (ММИ) физического факультета МГУ. Рецензент: доцент кафедры физической электроники И.К.Гайнуллин. Подписано в печать 28.12.2015. Объем 3,0 п.л. Тираж 50 экз. Заказ № Физический факультет им. М.В.Ломоносова, 119991 Москва, ГСП-1, Ленинские горы, д.1, стр. 2. Отпечатано в отделе оперативной печати физического факультета МГУ. c Физический факультет МГУ им. М.В. Ломоносова, 2015 c В.А.Антонюк, 2015 Программирование на видеокартах (GPGPU) Название этого спецкурса может показаться странным. Программирование? Это понятно. Но на видеокартах... Что такого есть в видеокартах, что нам захотелось бы их программировать? И зачем нам изучать, как это сделать? Оказывается, в своём развитии видеокарты не так давно превратились из специализированных устройств для воспроизведения графической информации в универсальные вычислительные устройства, причём с параллельной архитектурой: они способны не только исполнять задаваемую программу, но и выполнять операции одновременно над целым массивом данных. При этом функция воспроизведения никуда не исчезла, просто она стала далеко не основной, на передний план вышли функции формирования и обработки визуальной и любой другой информации. И теперь нам — обыкновенным пользователям программируемых устройств — предстоит осознать, как можно воспользоваться резко возросшими вычислительными возможностями наших устройств, и научиться это делать. Если, конечно, у нас есть вычислительно-трудоёмкие задачи, которые хотелось бы решать быстрее, и они потенциально могут быть (хотя бы частично) распараллелены. В англоязычной литературе даже появился специальный термин для такого рода программирования и вычислений: GPGPU, что принято расшифровывать как General-Purpose computation on Graphics Processing Units ( http://gpgpu.org/ ). Мы же используем здесь неформальное выражение "программирование на видеокартах". Теперь о самом спецкурсе. Условно его можно разделить на две части: "прямое" программирование видеокарт, предполагающее некоторое знакомство с их аппаратными свойствами и особенностями, и "косвенное" их программирование, когда используются специальные библиотеки, а пользователю вообще ничего не нужно знать об устройстве видеокарты. В первом случае речь идёт об освоении средств, предлагаемых одним из производителей видеокарт — фирмой 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), передаваемые ядрам параметры, особенности параллельного исполнения кода аппаратными средствами (нити, блоки), включая взаимодействие между нитями и синхронизацию их действий, приводятся примеры 0 3 простых ядер: для суммирования двух векторов, для вычисления суммы компонент вектора. Подробно разбираются некоторые основополагающие алгоритмы: умножение матриц ("наивная" реализация и блочная, с использованием разделяемой памяти 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 более комфортной. 4 Видеокарты и графические процессоры как супервычислители Окружающий нас мир переполнен созданными человеком электронными устройствами. И заметная доля их обладает серьёзными вычислительными возможностями, причём такими, что совсем недавно казались поразительными. Если раньше видеокарта настольного компьютера была просто устройством воспроизведения визуальной информации, то сейчас вполне может оказаться, что имеющаяся в нашем распоряжении видеокарта заметно мощнее в вычислительном смысле, чем центральный (пусть даже и многоядерный) процессор нашего компьютера. Если раньше мобильный телефон был просто миниатюрным вариантом (радио-) телефонной трубки, то сейчас это чаще всего небольшой, но мощный компьютер, содержащий многоядерный графический процессор и при этом «умеющий» совершать телефонные звонки. А ведь есть ещё масса других, в каком-то смысле «промежуточных» между ними устройств: ноутбуки, нетбуки, планшеты, неттопы и т.п. И все они в последнее время тоже обладают солидными вычислительными возможностями. Разумеется, хорошо было бы как-то использовать их. Но что для этого понадобится сделать, что нужно иметь и чего можно достичь — вот те вопросы, которые тогда встают перед нами. Для осознания возможностей графических процессоров полезно ознакомиться с некоторыми демонстрационными программами, «умеющими» эффективно использовать видеокарту. Всё, что необходимо для того, чтобы они могли работать, — это графический процессор на видеокарте и драйвер к ней. Обратите также внимание на то, что некоторые демонстрации (это касается тех из них, что запускаются в рамках браузера) вполне работоспособны даже на современных мобильных телефонах… Синтезированная шахматная доска (автор Thanassis Tsiodras) — традиционное приложение для операционной системы Windows, в котором осуществляется отрисовка трёхмерной сцены: доски с шахматными фигурами на ней ( https://www.thanassis.space/cudarenderer-BVH.html ), причём и доска, и фигуры имеют почти зеркальную поверхность, так что все они (многократно) отражаются друг в друге. В качестве ускорителя вычислений предполагается видеокарта архитектуры CUDA, при этом синтез картинки осуществляется в реальном времени с приличной скоростью (в зависимости от возможностей видеокарты — от единиц до десятков кадров в секунду), что позволяет наблюдателю интерактивно управлять вращением доски и углом зрения на неё. Если нам необходимо разрабатывать подобные приложения, то понадобится компилятор C/C++ и комплект разработки, именуемый CUDA Toolkit; также стоит установить в дополнение к нему и CUDA SDK. Имеющиеся там многочисленные примеры весьма поучительны и полезны — хотя бы потому, что для них приведён исходный код, — что позволяет использовать их как основу для собственных программ. 1 5 Эффектные демонстрации существуют и в рамках браузеров (на настольном компьютере или даже на мобильном телефоне), поскольку браузер — пример часто используемого приложения, которое ощутимо выигрывает от возможного ускорения синтезируемого графического вывода. Пожалуй, самая известная демонстрация такого рода — это WebGL Water (автор — Evan Wallace). Большое число аналогичных примеров можно посмотреть также на сайте babylonjs.com http://madebyevan.com/webgl-water/ http://www.babylonjs.com/cyos/ В настоящее время основными изготовителями видеокарт являются фирмы 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++ и преобразует их 6 (совместно с драйвером графической карты) в программу для процессора этой карты (далее обобщённо именуемого 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-программы в реальности «сосуществуют» два типа кода: для |