Программирование на видеокартах GPGPU (1184391), страница 2
Текст из файла (страница 2)
Кроме того,выпускается широкий ряд графических процессоров (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<<<M,N>>>(...);. . .Легко видеть, что вызов ядер на исполнение инициируется в программе упоминанием их имени стак называемыми параметрами конфигурации запуска ядер (которые указываются в тройныхугловых скобках <<<...>>>) наряду с обычными параметрами, необходимыми для работыкаждого отдельного ядра.Таким образом, в рамках CUDA-программы в реальности «сосуществуют» два типа кода: для CPU(в терминологии CUDA именуемого host) и для GPU (в терминологии CUDA — device), причём обатипа исходного кода используют язык C/C++, только в device-коде пока употребляется некотороеподмножество C99, а вот host-часть программного кода вполне может быть написана на C++.7Подобный симбиоз позволяет программисту оставаться в привычной среде одного языка, правда,«разбавленного» некоторыми дополнительными ключевыми словами и типами величин.Основная работа по анализу такой «смеси» ложится на компилятор 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 для статических библиотек.82Расширения языка 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_<TypeName>, где <TypeName> — один извышеприведённых типов данных, например: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, можно задавать предпочтительнуюконфигурацию запуска ядер на исполнение.
![](https://s.studizba.com/z.php?f=/templates/si/i/noavatar.png&w=100&h=100&t=1"){kind=avatar}
