Методичка (CUDA) - черновик. Введение cuda (Compute Unified Device Architecture)Ошибка! Источник ссылки не найден
 Скачать 1.06 Mb.
|
|
1 Введение CUDA (Compute Unified Device Architecture)Ошибка! Источник ссылки не найден. – программно-аппаратная архитектура массового параллелизма для графических процессоров NVIDIA©. Данная архитектура обладает одной из наиболее успешных и простых моделей массового параллелизма, необходимого для эффективного программирования GPU. Изначально GPU не имели универсальной архитектуры и использовались только для рендеринга 3D сцен. По мере роста сложности создаваемых сцен и улучшения их качества потребовалось внести урезанные программные возмож- ности по трансформации геометрии и текстур. Эти мини программы получили название вершинные и пиксельные шейдеры. Начиная с версии 2.0, которые вошли в состав Microsoft© DirectX 9.0, была введена поддержка циклов. Имен- но на этих шейдерах впервые были реализованы универсальные алгоритмы, та- кие как, перемножение матриц, вычисления DCT и др. Это послужило началом создания новой области применения графических процессоров – GPGPU (Gen- eral Purpose Computation on GPU). Из-за высокой сложности программирования шейдеров и недостаточного развития средств разработки, развитие GPGPU шло достаточно медленно. Следующим этапом развития GPGPU стало объединение вершинных и пиксельных шейдеров. Именно с этого момента GPU стали универсальными вычислительными устройствами. Осознав высокий потенциал GPU в области универсальных вычислений, компания NVIDIA© в 2008 году представила первую версию технологии CUDA. Эта модель не является исторически первой архитектурой для универ- сальных вычислений на GPU. Чуть раньше компания ATI/AMD© представила архитектуру, разработанную совместно с университетом Стенфорда, которая называлась BrookGPU[2]. Несмотря на некоторые революционные идеи, зало- женные в ней, она обладала меньшей универсальностью и не имела явных средств для работы с разделяемой памятью. Из-за достаточно сырой программ- ной реализации компилятора и средств разработки данная модель не получила широкого распространения. В 2010 году представлена технология параллельного программирования OpenCL (Open Computational Language)Ошибка! Источник ссылки не найден.. Технология является открытым стандартом и поддерживаемая рядом крупных компаний (Apple©, Intel©, AMD©). C точки зрения вычислительной модели и способа выражения параллелизма эта технология идентична CUDA. 2 Из-за того, что OpenCL развивается достаточно медленно и его поддерживает множество компаний, первоначально все новые особенности GPU отражаются в CUDA, а уже потом в OpenCL. В настоящее время компания NVIDIA© являет- ся лидером на рынке GPGPU, ее сопроцессоры устанавливаются в большинство современных суперкомпьютеров. 3 1. Программная модель CUDA Классический процессор, построенный по фон-неймановской архитектуре, не подразумевает наличие параллелизма: есть один поток команд, все команды исполняются последовательно и обрабатывают один поток данных. Степень параллелизма у такого процессора равна 1 (количество одновременно обраба- тываемых элементов данных). Для данной модели рост производительности возможен только за счет роста тактовой частоты. Современные процессоры умеют использовать параллелизм на уровне инструкций за счет суперскалярно- го и внеочередного исполнения. Но это ведет к росту сложности управляющей логики, которая начинает доминировать над исполняющими устройствами. Вся микроархитектура центрального процессора ориентирована на уменьшение ла- тентности, т.к. для скорейшего исполнения единственного потока команд нуж- но как можно быстрее исполнить цепочку зависимых инструкций на критиче- ском пути. Архитектуры массового параллелизма ориентированы на общую произво- дительность системы (пропускную способность), когда однотипные действия выполняется над массивом независимых элементов. GPU изначально являлись архитектурой массового параллелизма, так как задача построения 3D сцены как раз и является задачей с однотипной обработкой массива независимых элемен- тов (вершин или пикселей). Для таких архитектур важным является скорейшее завершение всей задачи, нежели обработка одного элемента, т.к. все элементы множества обладают равным приоритетом. Степень параллелизма таких задач определяется количеством элементов в массиве. Для задач построения 3D сцен степень параллелизма может достигать порядка – Латентность для устройств с архитектурой массового параллелизма особой роли не играет, так как наличие большого количества независимых элементов позволяет замаскировать вычисления одних элементов другими. 1.1. Нить, блок, сетка блоков В упрощенном виде вычислительный блок графического процессора пред- ставляет собой совокупность потоковых мультипроцессоров (streaming multi- processor, SMX), управляемые с помощью GigaThread Engine. Каждый SMX со- стоит из множества вычислительных CUDA-ядер (рисунок 2.1). 4 GigaThread Engine L2 кэш К он тр ол ле ры п ам ят и К он тр ол ле ры п ам ят и SMX CUDA- ядра SMX CUDA- ядра SMX CUDA- ядра SMX CUDA- ядра SMX CUDA- ядра SMX CUDA- ядра Рис. 1.1. Упрощенная структура графического процессора Такая архитектура легла в основу базовых понятий в CUDA – нити (thread), блока (block) и решетки блоков (grid). Нить (thread) – базовая абстракция, представляющая собой легковесный поток, отвечающий за исполнение инструкций. Именно нить соответствует од- ному исполняющему ядру CUDA мультипроцессора. Максимальное число ни- тей, которое можно определить, - 2048 (для архитектуры NVIDIA© Kepler). Легковесность нити заключается в том, что время, необходимое на его созда- ние, разрушение и переключение между нитями, ничтожно мало. Для упроще- ния аппаратной схемы каждые 32 CUDA ядра соединены с одним счетчиком команд, поэтому данные блоки ядер всегда выполняют одну и ту же инструк- цию. Данные 32 нити называются варпом (warp). Множество нитей объединяется в блок (block), который проецируется на мультипроцессор. Это означает, что все нити блока всегда будут выполняться на выделенном для него мультипроцессоре. Решетка блоков (grid) представляет собой самый высокий уровень аб- стракции и, как правило, представляет собой всю задачу, которую требуется решить. 5 Каждая из абстракций является 3-мерной. Для ее определения в NVIDIA© CUDA определена специальная структура dim3 . Инициализация значений структуры осуществляется следующим образом: dim3 blocks (16, 16); // эквивалентно blocks ( 16, 16, 1 ) dim3 grid (256); // эквивалентно grid ( 256, l, 1 ) Для того чтобы определить текущие координаты в решаемой задаче введе- ны следующие константы: // размерность решетки блоков при запуске ядра dim3 gridDim; // координаты текущего блока внутри решетки блоков uint3 blockIdx; // размерность блока при запуске ядра dim3 blockDim; // координаты текущей нити внутри блока uint3 threadIdx; Базовая задача, которую программист должен решить, - правильное разби- ение решаемой задачи между нитями и блоками для обеспечения эффективной нагрузки всех имеющихся вычислительных ресурсов. Например, графический процессор с архитектурой NVIDIA© Kepler со- держит до 15 мультипроцессоров, каждый из которых содержит 192 CUDA- ядра. Всего получается до 2880 CUDA ядер. Для утилизации данной вычисли- тельной мощности требуется как минимум 15 блоков (так как каждый блок проецируется на один мультипроцессор) и по 192 нити в каждом блоке. Данная конфигурация, с одной стороны, сможет сопоставить на каждое CUDA ядро фрагмент решаемой задачи, а, с другой стороны, не является эффективной. Главным преимуществом CUDA является легковесность нитей и ничтожно ма- лое время переключения между ними. Встроенный в графический процессор планировщик следит за загруженностью ядер и, если нити варпа перешли в ре- жим ожидания результатов выполнения инструкции, планировщик может пере- ключиться на выполнение других нитей данного блока, что позволяет снизить простаивание системы и в итоге время выполнения всего кода. Поэтому при проектировании конфигурации нитей и блоков рекомендуется подбирать кон- фигурацию таким образом, чтобы она была заведомо большей, чем имеющиеся вычислительные ресурсы. 6 1.2. Ядро, конфигурирование и запуск ядра Ядро (kernel) – функция, описывающая последовательность операций, вы- полняемых каждой нитью параллельно. Ядро, выполняющее поэлементное сложение двух матриц, описывается следующим образом: __global__ void MatAdd(float *A, float *B, float *C) { int idx = blockIdx.x * blockDim.x + threadIdx.x; C[idx] = A[idx] + B[idx]; } Для определения устройства исполнения функции (ядра) в CUDA вводятся следующие ключевые слова: __host__ – функция может быть вызвана и выполнена только на сто- роне хоста (на центральном процессоре). Если не указывается специфи- катор, то функция считается объявленной как __host__ __global__ – ядро, предназначенное для выполнения на устройстве (графическом процессоре) и может быть вызвано с хоста. На устрой- ствах с Compute Capability 3.5 (способ описания версии архитектуры графического процессора и поддерживаемого CUDA API) и новее до- пускается вызов с устройства. Ядра данного типа не возвращают ника- ких данных (тип void ); __device__ – ядро выполняется на устройстве, вызывается из кода, вы- полняемого на устройстве; Вызов ядра MatAdd выполняется следующим образом: int main(){ const int N = 1024 * 1024; dim3 threadsPerBlock(1024); dim3 numBlocks(N / threadsPerBlock.x); MatAdd<< } Здесь считается, что матрицы d_A , d_B и d_C имеют размер N элементов. Специальное расширение языка C <<<...>>> определяет конфигурацию за- пуска нитей. Первый параметр задает количество блоков по каждому из 3-х из- 7 мерений (тип dim3 ), а второй – количество нитей в блоке, так же по каждому измерению. Все запуски ядер CUDA являются асинхронными: CPU запрашивает раз- решение на запуск ядра путем записи в специальный буфер команд без провер- ки выполнен код или нет, после этого продолжает выполнять код в соответ- ствии с алгоритмом для хоста. 1.2.1. Выделение памяти, копирование на GPU Хост и устройство имеют разные типы памяти. Ядро на момент написания методического пособия может взаимодействовать только с памятью устройства, поэтому перед запуском ядра требуется выделить память на устройстве и ско- пировать туда данные из памяти хоста. После завершения вычислений резуль- тат копируется обратно, а выделенная память должна быть освобождена. Память на устройстве может быть выделена как линейная память или как CUDA массив (CUDA Array) – специальная разметка памяти, оптимизированная для работы с текстурами. Более подробно массивы описываются в Ошибка! Источник ссылки не найден.. Линейная память выделяется с помощью функции cudaMalloc() и осво- бождается с помощью cudaFree() . Передача данных между памятью хоста и девайса чаще всего выполняется с помощью cudaMemcpy() Для рассмотренного ранее ядра MatAdd выделение памяти производится следующим образом: int main() { const int N = 1024 * 1024; size_t size = N * sizeof(float); // выделение памяти h_A и h_B на хосте float* h_A = (float*)malloc(size); float* h_B = (float*)malloc(size); float* h_C = (float*)malloc(size); // Инициализация входных массивов // Выделение векторов в памяти GPU float* d_A, *d_B, *d_C; cudaMalloc(&d_A, size); cudaMalloc(&d_B, size); cudaMalloc(&d_C, size); 8 // Копирование векторов из памяти хоста // в память устройства cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // Вызов ядра dim3 threadsPerBlock(1024); dim3 numBlocks(N / threadsPerBlock.x); MatAdd<< // Копирование результата с устройства на хост cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); // Освобождение памяти устройства cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); free(h_A); free(h_B); free(h_C); } Линейная память может быть выделена так же функциями cudaMallocPitch() и cudaMalloc3D() . Эти функции рекомендуется использо- вать для выделения памяти под 1D и 3D массивы соответственно, при этом начало каждой строки или слоя размещается по адресу кратному 128, что необ- ходимо для быстрого доступа к глобальной памяти. Соседние строки или слои могут размещаться не последовательно (за последним байтом одной строки не гарантируется, что следует первый байт следующей), а с некоторым смещени- ем, которое называется pitch (или stride для 3D). Пример использования функ- ции cudaMallocPitch() будет показан в разделе 4. Для копирования данных в/из 2D памяти устройства используется функция cudaMemcpy2D( void* dst,// адрес приемника данных size_t dpitch,// фактическая ширина одной строки приемника const void* src, // адрес источника данных size_t spitch,// фактическая ширина одной строки источника size_t width, // ширина копируемых данных в байтах size_t height, // высота матрицы копируемых данных cudaMemcpyKind kind // направление копирования 9 ) Параметр cudaMemcpyKind используется во всех функциях передачи дан- ных и может иметь следующие значения: cudaMemcpyHostToHost = 0 // Host -> Host cudaMemcpyHostToDevice = 1 // Host -> Device cudaMemcpyDeviceToHost = 2 // Device -> Host cudaMemcpyDeviceToDevice = 3 // Device -> Device cudaMemcpyDefault = 4 // По-умолчанию 1.3. Конфигурирование проекта (nvcc, -arch, -code, -gencode) Код программы, написанный с использованием CUDA, должен сохранять- ся в файлы с расширением .cu , заголовочные файлы – с расширением .cuh Исходный код CUDA программы компилируется с помощью утилиты nvcc, распространяемой вместе с CUDA Toolkit. Данная утилита является консольной и позволяет управлять процессом компиляции с помощью параметров команд- ной строки. Исходный код, компилируемый с помощью nvcc, как правило, содержит код, выполняемый на CPU, и код, выполняемый на GPU. Процесс компиляции выглядит следующим образом: код, предназначенный для CPU, компилируется с использованием ком- пилятора C/C++, используемого по-умолчанию; компилируется код, предназначенный для GPU, в ассемблерную форму (специальный PTX код) и/или бинарную форму (cubin объект). При этом с помощью опций компилятора возможно настроить версию Com- pute Capability, для которой осуществляется генерация PTX-кода, и вер- сию архитектуры при генерации бинарного кода; полученный объектный файл интегрируется в исполняемый (либо в библиотеку) и добавляется код управления данным объектом. Для определения версии Compute Capability, т.е. для определения версии CUDA API и PTX-кода, используется параметр compute_XX, где XX – версия Compute Capability (допустимые значения – compute_10, compute_11, compute_12, compute_13, compute_20, compute_30, compute_32, compute_35, compute_50). Для определения версии архитектуры, для которой создается би- нарный файл, используется параметр sm_XX, где XX – версия архитектуры (допустимые значения – sm_10, sm_11, sm_12, sm_13, sm_20, sm_21, sm_30, sm_32, sm_35, sm_50). 10 Для управления процессом компиляции используются следующие пара- метры: -arch – определяет архитектуру графического процессора, для которой осуществляется создание PTX-кода. Значение по-умолчанию – compute_10; -code – определяет архитектуру графического процессора, для которой будет создан PTX и/или бинарный код. Значение по-умолчанию – compute_10,sm_10; -gencode – объединяет опции -arch и -code в качестве одного параметра. При использовании параметра sm_XX не гарантируется совместимость между различными версиями архитектур графического процессора, т.е. для би- нарного кода, скомпилированного для архитектуры sm_30, не гарантируется, что он будет исполнен на графическом процессоре с архитектурой sm_20. В тоже время указание, например, compute_20 добавит PTX-код с поддержкой CUDA API и PTX для архитектуры 2.0 в исполняемый файл. При этом поведение драйвера можно описать следующим образом: бинарная версия кода для текущей архитектуры есть в исполняемом файле? o да – исполняем бинарную версию кода; o нет – проверяем PTX-код: есть PTX код и версия API в PTX-коде меньше или равна поддерживаемой текущей архитектурой? да – компилируем код с помощью драйвера и исполня- ем; нет – прерываем работу с генерацией ошибки. PTX-код отсутствует – прерываем работу с генерацией ошибки. Также компилятором поддерживается множественное задание различных архитектур (например, для версий 1.0, 2.0, 3.0). При этом компилятором будет созданы версии кода для каждой из архитектур и поведение драйвера в резуль- тате усложняется, т.к. требуется из всего множества версий найти наиболее подходящую. Пример использования опции -gencode: nvcc x.cu -gencode arch=compute_20,code=sm_20 -gencode arch=compute_30,code=\'compute_30,sm_30\' Примечание [D1]: Оставить так? Или сделать блок-схему 11 В коде программы можно определить участок кода для различных архи- тектур, используя макрос __CUDA_ARCH__ . Данный макрос будет работать толь- ко в ядре. Например, при компиляции с параметром arch=compute_20 , значе- ние __CUDA_ARCH__ равно 200 При использовании CUDA Toolkit 6.0 и новее не рекомендуется использо- вание архитектуры версии 1.0, т.к. данная версия считается устаревшей. 1.4. Использование событий в CUDA При работе с технологией CUDA вместо стандартных методов измерения времени (функции GetTickCount , time и другие) рекомендуется использовать CUDA-события (events). При этом будут задействованы счетчики, расположен- ные на графическом процессоре, что обеспечит более качественное измерение времени при работе с CUDA. Следующий пример демонстрирует создание, использование и уничтоже- ние событий для замера времени выполнения передачи данных и работы ядра: cudaEvent_t start, stop; // создание события для точки старта cudaEventCreate(&start); // создание события для точки завершения cudaEventCreate(&stop); // точка начала замера времени cudaEventRecord(start, 0); for (int i = 0; i < 2; ++i) { cudaMemcpyAsync(inputDev, inputHost, size, cudaMemcpyHostToDevice, 0); MyKernel<<<100, 512>>> (outputDev, inputDev, size); cudaMemcpyAsync(outputHost, outputDev, size, cudaMemcpyDeviceToHost, 0); } // точка завершения замера времени cudaEventRecord(stop, 0); // ожидание завершения выполнения задач на GPU cudaEventSynchronize(stop); float elapsedTime; // elapsedTime - затраченное время в миллисекундах cudaEventElapsedTime(&elapsedTime, start, stop); 12 // уничтожение объектов событий cudaEventDestroy(start); cudaEventDestroy(stop); При исполнении кода функция cudaEventCreate создает событие и с по- мощью функции cudaEventRecord фиксируются интересуемые точки, между которыми планируется измерить время работы (переменные start и stop ). Так как ядро вызывается асинхронно и используются асинхронные версии функций копирования, требуется осуществить синхронизацию для гарантии того, что яд- ро завершило свою работу. Для этого используется функция cudaEventSynchronize После синхронизации можно определить время между двумя событиями с помощью функции cudaEventElapsedTime , которая возвращает время в мил- лисекундах. |