ТЕХНОЛОГИЯ ПРОГРАММИРОВАНИЯ CUD. Учебное пособие казань 2017 2 удк 004. 4, 681. 3 Ббк 32. 81 Печатается по постановлению Редакционноиздательского совета
 Скачать 1.28 Mb.
|
|
КАЗАНСКИЙ ФЕДЕРАЛЬНЫЙ УНИВЕРСИТЕТ Д.Н. Тумаков, Д.Е. Чикрин, А.А. Егорчев, С.В. Голоусов ТЕХНОЛОГИЯ ПРОГРАММИРОВАНИЯ CUDA Учебное пособие КАЗАНЬ 2017 2 УДК 004.4, 681.3 ББК 32.81 Печатается по постановлению Редакционно-издательского совета ФГАОУВПО «Казанский (Приволжский) федеральный университет» Научный редактор доктор физ.-мат. наук, проф. Н. Б. Плещинский Рецензенты кандидат физ.-мат.наук, доцент И.Е. Плещинская кандидат физ.-мат.наук, доцент Е.В. Рунг Тумаков Д.Н. Технология программирования CUDA: учебное пособие / Д.Н. Тумаков, Д.Е. Чикрин, А.А. Егорчев, С.В. Голоусов. – Казань: Казанский государственный университет, 2017. – 112 с. ISBN 978-5-00019-913-8 В пособии описаны технология CUDA и основные принципы работы с ней. Рассмотрены вопросы использования разделяемой, константной и текстурной памяти. Приведены примеры реализации различных алгоритмов. Описаны пакеты для работы с линейной алгеброй: cuBLAS и cuBLAS-Xt. В конце глав содержатся упражнения для лабораторных и домашних занятий. Пособие предназначено для студентов, магистрантов и аспирантов ВУЗов, специализирующихся в области параллельного программирования и численных методов. © Казанский университет, 2017 © Тумаков Д.Н., Чикрин Д.Е., 2017 ISBN 3 Оглавление Оглавление ..................................................................................................................................................... 3 Введение .......................................................................................................................................................... 5 Глава 1. Основы работы с CUDA ............................................................................................................... 7 Основные понятия .................................................................................................................................... 7 Сетки, блоки, нити .................................................................................................................................... 8 Варпы ......................................................................................................................................................... 10 Первая программа ................................................................................................................................... 13 Версии Compute Capability .................................................................................................................... 14 Глава 2. Свойства устройства ................................................................................................................... 16 Архитектура видеокарт NVidia ............................................................................................................. 16 Типы видеокарт, поддерживающих CUDA ........................................................................................ 19 Функция cudaGetDeviceProperties() ..................................................................................................... 21 Глава 3. Простейшая программа на CUDA ............................................................................................ 27 Функция копирования ........................................................................................................................... 28 Замер времени работы части кода программы ................................................................................. 30 Обработка ошибок................................................................................................................................... 31 Глава 4. Атомарные операции в CUDA .................................................................................................. 34 Атомарные арифметические операции ............................................................................................... 34 Атомарные побитовые операции ......................................................................................................... 36 Дополнительные возможности Compute Capability 6.x .................................................................... 37 Пример использования атомарных операций ................................................................................... 38 Глава 5. Работа с векторами. Математические функции .................................................................... 40 Сложение векторов ................................................................................................................................. 40 Вычисление математических функций ............................................................................................... 42 Глава 6. Работа с матрицами .................................................................................................................... 47 Создание матриц ...................................................................................................................................... 48 Транспонирование матрицы ................................................................................................................. 51 Сложение матриц .................................................................................................................................... 57 Умножение матриц .................................................................................................................................. 58 Глава 7. Типы памяти. Разделяемая память ......................................................................................... 64 Банки данных ........................................................................................................................................... 67 Пример. Оптимизация программы по перемножению двух матриц ............................................ 69 Глава 8. Константая и текстурная память ............................................................................................. 75 Константная память ............................................................................................................................... 75 Текстурная память .................................................................................................................................. 76 Примеры использования одномерной текстуры ............................................................................... 78 4 Глава 9. Пакеты для работы с векторами и матрицами ..................................................................... 83 Пакет cuBLAS .......................................................................................................................................... 83 Пакет cuBLAS-Xt ................................................................................................................................... 102 Литература .................................................................................................................................................. 111 5 Введение Параллельные вычисления давно уже перестали быть специальными технологиями, которыми пользовались лишь в исследовательских центрах. Сегодня методы параллельных вычислений в той или иной степени используются при решении практически всех сложных задач. Совершенствование технической базы компьютеров в сочетании с параллельной обработкой данных позволило существенно повысить производительность вычислений. В данном пособии рассмотрим одну из технологий высокопроизводительных вычислений, использующую графический процессор (GPU). Изначально этот класс устройств разрабатывался для обработки графики. Техника совершенствовалась, GPU наращивали производительность, и в какой-то момент оказалось, что GPU можно успешно использовать не только для задач компьютерной графики, но и как математический сопроцессор для CPU, получая при этом существенный прирост в производительности. Этот класс технологий получил название GPGPU (General-purpose computing for graphics processing units) – использование графического процессора для различных вычислительных задач. Популярный представитель этого класса – CUDA (Compute Unified Device Architecture), который был впервые представлен компанией NVIDIA в 2007 году. CUDA может использоваться на GPU производства NVIDIA, таких как GeForce, Quadro, Tesla, Fermi, Kepler, Maxwell, Pascal и Volta. Некоторые из этих устройств не имеют видеовыхода, и эти видеокарты предназначены исключительно для высокопроизводительных вычислений. CUDA работает только с устройствами производства NVIDIA, но отметим, что помимо CUDA существуют и другие аналогичные технологии, например OpenCL, AMD FireStream, CUDAfy и PyCUDA. 6 В качестве дополнительной литературы можно порекомендовать [1-3]. В [1] разобрана как сама технология CUDA, так и вопросы оптимизации кода с использованием ptx-ассемблера. В книге также рассмотрены задачи по моделированию ряда физических процессов. Задачи моделирования и использование стандартных библиотек рассмотрены и в [2]. В [3] подробно описана работа с несколькими графическими устройствами. 7 Глава 1. Основы работы с CUDA В первой главе рассмотрим базовые понятия технологии CUDA – хост, устройство, ядро, сетка, блок и нить. Основные понятия Хост (host) – это центральный процессор (CPU) компьютера, устройство (deviсe) – графический процессор (GPU), расположенный на видеокарте. Каждый процессор имеет свою отдельную память, память хоста и память устройства физически разделены. Хост управляет вычислительным процессом – с хоста вызываются функции, организующие обмен данными между различными видами памяти, и запускаются задачи на устройстве. Каждой такой задаче соответсвует ядро (kernel) – функция, оформленная по особым правилам. В настоящее время технологию CUDA поддерживают такие языки программирования, как C/C++, Fortran и Phyton. В программу, написанную для CPU, можно добавить дополнительные элементы для работы с GPU. Компилятор CUDA работает как препроцессор. После предварительной обработки исходного кода вызывается стандартный компилятор и генерируется исполняемый код отдельно для хоста и отдельно для устройства. Обычно в CUDA-программах при работе с устройством содержатся следующие этапы: • выделяется память под данные на устройстве; • копируются данные из памяти хоста в память устройства; • на устройстве выполняются ядра; • результаты пересылаются из памяти устройства в память хоста. В расширении языка C/C++ используют три спецификатора функций, которые определяют, откуда функции вызываются и где выполняются: __host__ – вызываются с хоста, выполняются на хосте (можно не указывать действует по умолчанию); __global__ – вызываются с хоста, выполняются на устройстве; 8 __device__ – вызываются с устройства, выполняются на устройстве. Сетки, блоки, нити При вычислениях на GPU одновременно запускается большое число параллельных процессов, их принято называть нитями. Все запущенные на выполнение нити объединены в сложную структуру – сетку (grid). Сетка представляет собой одномерный, двухмерный или трехмерный массив блоков (block). Каждый блок – это одномерный, двухмерный или трехмерный массив нитей (thread). При этом все блоки, образующие сетку, имеют одинаковую размерность и размер (Рис. 1). Рис. 1. Иерархия нитей в CUDA Каждый блок в сетке имеет свой адрес, состоящий из одного или двух неотрицательных целых чисел (индекс блока в сетке). Аналогично каждая нить 9 внутри блока также имеет также свой адрес – одно, два или три неотрицательных целых числа, задающих индекс нити внутри блока. Поскольку одно и то же ядро выполняется одновременно очень большим числом нитей, то для того, чтобы ядро могло однозначно определить номер нити (а значит, и элемент данных, который нужно обрабатывать), используются встроенные переменные threadIdx (индекс текущей нити в вычислении на GPU, имеет тип uint3) и blockIdx (индекс текущего блока в вычислении на GPU, имеет тип uint3). Каждая из этих переменных является трехмерным целочисленным вектором. Обратим внимание, что они доступны только для функций, выполняемых на GPU; для функций, выполняющихся на CPU, они не имеют смысла. Исходная задача Подзадача Подзадача Подзадача Подзадачи Нити Рис. 2. Разбиение исходной задачи на набор независимо решаемых подзадач Также ядро может получить размеры сетки и блока через встроенные переменные gridDim (размерность сетки, имеет тип dim3) и blockDim (размерность блока, также имеет тип dim3). Подобное разделение всех нитей является еще одним общим приемом использования CUDA: исходная задача разбивается на набор отдельных 10 подзадач, решаемых независимо друг от друга (Рис. 2). Каждой такой подзадаче соответствует блок нитей. При этом каждая подзадача совместно решается всеми нитями своего блока. Объединение нитей в варпы происходит отдельно для каждого блока; таким образом, все нити одного варпа всегда принадлежат одному блоку. При этом нити могут взаимодействовать между собой только в пределах блока. Нити разных блоков взаимодействовать между собой не могут. Подобный подход является удачным компромиссом между необходимостью обеспечить взаимодействие нитей между собой и стоимостью подобного взаимодействия – обеспечить возможность взаимодействия каждой нити с каждой было бы слишком сложно и дорого. Однако есть способы косвенного взаимодействия нитей в блоке. Они могут взаимодействовать друг с другом через разделяемую (shared) память (об использовании разделяемой памяти речь пойдет в Главе 7). Каждый блок получает в свое распоряжение определенный объем быстрой разделяемой памяти, которую все нити блока могут совместно использовать. Поскольку нити блока необязательно выполняются физически параллельно (то есть мы имеем дело не с чистой SIMD-архитектурой, а имеет место прозрачное управление нитями), то для того, чтобы не возникало проблем с одновременной работой с shared-памятью, необходим некоторый механизм синхронизации нитей блока. CUDA предлагает довольно простой способ синхронизации – это так называемая барьерная синхронизация. Для ее осуществления используется вызов встроенной функции __syncthreads(), которая блокирует вызывающие нити блока до тех пор, пока все нити блока не войдут в эту функцию. Таким образом, при помощи __syncthreads() можно организовать «барьеры» внутри ядра, гарантирующие, что если хотя бы одна нить прошла такой барьер, то не осталось ни одной за барьером (не прошедшей его) (Рис. 3) Варпы Варп (Warp) – группировка потоков по 32 штуки, которые оказываются 11 частями более крупных образований – блоков (blocks). Все потоки каждого блока запускаются строго на одном потоковом мультипроцессоре (SM, Streaming Multiprocessors), поэтому имеют доступ только к его ресурсам. Однако, на одном SM может запускаться более одного блока, и тогда ресурсы будут разделяться между ними поровну. syncthreads syncthreads syncthreads Рис. 3. Барьерная синхронизация В каждом SM существует блок управления, который занимается распределением ресурса процессорного времени. Делается это так, что в каждый момент времени все ядра одного SM исполняют строго один варп. По его завершению работы варпа оптимальным способом выбирается следующий варп, приписанный к данному SM. Таким образом, оказывается, что потоки одного варпа синхронизируются за счет аппаратной особенности CUDA и исполняются по еще более близкому к SIMD (Single Instruction, Multiple Data) методу. Но потоки даже одного блока из разных варпов могут оказаться заметно рассинхронизированными. Можно сделать вывод, что взаимодействие между потоками одного блока следует стараться осуществлять через их общую быструю разделяемую память, а между потоками двух различных блоков – только с использованием глобальной памяти. Тут возникает проблема: слежение за актуальностью данных в общедоступной разным потокам для чтения и записи области памяти. Иначе говоря – проблема синхронизации потоков. Как уже отмечалось, в 12 рамках одного блока потоки каждого варпа синхронизированы между собой. Для синхронизации потоков блока вне зависимости от принадлежности к варпам существует несколько команд барьерного типа. Как была сказано выше, основная из них – __syncthreads(). Эта функция заставляет каждый поток ждать, пока все остальные потоки этого блока достигнут этой точки и все операции по доступу к разделяемой и глобальной памяти, совершенные потоками этого блока, завершатся и станут видны потокам этого блока. Не желательно размещать эту команду внутри условного оператора if, а также следует обеспечивать безусловный вызов этой функции всеми потоками блока. Выпишем остальные «барьерные команды»: • __threadfence_block() заставляет ждать вызвавший её поток, пока все совершенные операции доступа к разделяемой и глобальной памяти завершатся и станут видны потокам этого блока. • __threadfence() заставляет ждать вызвавший её поток, пока все совершенные операции доступа к разделяемой памяти станут видны потокам этого блока, а операции с глобальной памятью – всем потокам на девайсе. • __threadfence_system() подобна __threadfence(), но включает синхронизацию с потоками на CPU («хосте»), при использовании page-locked памяти. Параметры запуска ядра служат для описания количества блоков, нитей и памяти, которые необходимо выделить при расчете на GPU. Например, запуск ядра может быть осуществлен следующим образом: myKernelFunc<< 13 • gridSize – размерность сетки блоков (тип dim3), выделенную для расчетов, • blockSize – размер блока (тип dim3), выделенного для расчетов, а myKernelFunc – функция ядра (спецификатор __global__). Дополнительные типы переменных и их спецификаторы будут рассмотрены непосредственно в примерах работы с памятью. Первая программа В любом учебном пособии первая программа, которую предлагают вам написать - это программа, выводящая на консоль «Hello, world!». Данная программа на языке программирования C++ выглядит следующим образом: #include "cuda_runtime.h" #include "device_launch_parameters.h" #include __global__ void HelloWorld() { printf("Hello world, %d, %d\n", blockIdx.x, threadIdx.x); } int main() { HelloWorld << <1, 1 >> >(); // ожидаем нажатия любой клавиши getchar(); return 0; } Первые две строки в программе (два #include) не обязательно. В среде Visual Studio они подключаются по умолчанию. 14 В отличие от файлов с программами на языке программирования C++, имеющих расширения .cpp, файлы с программами, которые используют технологию CUDA, должны иметь расширение .cu. Версии Compute Capability Перечислим поддержку возможностей в версиях вычислительной производительности (Compute Capability): • с версии 1.1 – атомарные операции над 32-битными целыми числами и atomicExch() над 32-битными числами с плавающей точкой в глобальной памяти; • с версии 1.2 – тоже в разделяемой памяти, также в глобальной над 64-разрядными целыми, функции голосования в варпе (warp vote functions); • с версии 1.3 – операции над числами с плавающей точкой двойной точности; • с версии 2.0 – атомарные операции над 64-битными целыми числами в разделяемой памяти и атомарное сложение 32-битных чисел с плавающей точкой, функции _ballot(), _threadfence_system(), _syncthreads_count(), _syncthreads_and(), _syncthreads_or(), поверхностные функции, трехмерная сетка блоков; • с версии 3.0 – функции варп-перемешивания (warp shuffle functions); • с версии 3.2 – сдвиги __funnelshift_lc() и __funnelshift_rc() (funnel shift); • c версии 3.5 – динамический параллелизм; • с версии 5.3 – операции над числами половинной точности; • с версии 6.0 - атомарное сложение 64-битных чисел с плавающей точкой в глобальной и разделяемой памяти. Заметим, что версия производительности – не то же самое, что версия CUDA SDK; она зависит он архитектуры конкретной видеокарты. 15 • CUDA SDK 6.5 – последняя версия, поддерживающая производительность 1.x; • CUDA SDK 7.5 поддерживает 2.0 – 5.x; • CUDA SDK 8.0 поддерживает 2.0 – 6.x; • В версии 9.0 поддержка производительности 2.x удалена. Вопросы к разделу 1. Какие основные части содержит CUDA-программа? 2. Что такое ядро? Опишите все возможные спецификаторы функций и параметры ядра. 3. Как организована иерархия нитей в CUDA? 4. Какие типы синхронизации потоков существуют? Лабораторная работа 1. Определите характеристики используемой видеокарты: число SM, объем памяти, частоту работы видеопроцессоров и памяти, ширину полосы пропускания памяти. Сделайте выводы о производительности видеокарты. 2. Напишите программу, которая выводит на экран результат сложения двух чисел. 16 |