МОСКОВСКИЙ АВИАЦИОННЫЙ ИНСТИТУТ (государственный технический университет) (государственный технический университет) Факультет прикладной математики и физики Кафедра вычислительной математики и программирования Выполнил: Семенов С.А. Руководитель: Ревизников Д.Л. Лекция 10 «Вопросы программирования и оптимизации приложений на CUDA»

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 2 Подход CUDA Исходная задача разбивается на подзадачи, которые можно решать независимо друг от друга. Каждая из этих подзадач решается набором взаимодействующих между собой нитей

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 3 Программная модель CUDA Код состоит как из последовательных, так и из параллельных частей Последовательные части кода выполняются на CPU Массивно-параллельные части кода выполняются на GPU как ядра

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 4 Программная модель CUDA Параллельная часть кода выполняется как большое количество нитей (threads) Нити группируются в блоки (blocks) фиксированного размера Блоки объединяются в сеть блоков (grid) Ядро выполняется на сетке из блоков Каждая нить и блок имеют свой уникальный идентификатор

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 5 Программная модель CUDA Потоки в CUDA объединяются в блоки: –Возможна 1D, 2D, 3D топология блока Общее кол-во потоков в блоке ограничено В текущем HW это 512 потоков

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 6 Программная модель CUDA Потоки в блоке могут разделять ресурсы со своими соседями float data[N]; for ( int ix = 0; ix < nx; ix++ ) data[ix] = f(ix, data[ix / n]);

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 7 Программная модель CUDA Блоки могут использовать shared память – Т.к. блок целиком выполняется на одном SM – Объем shared памяти ограничен и зависит от HW Внутри Блока потоки могут синхронизироваться –Т.к. блок целиком выполняется на одном SM

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 8 Программная модель CUDA Блоки потоков объединяются в сетку (grid) потоков –Возможна 1D, 2D топология сетки блоков потоков

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 9 Синтаксис CUDA Спецификаторы Спецификатор Выполняется на Может вызываться из __device__device __global__devicehost __host__host Спецификатор функций Спецификатор переменных Спецификатор НаходитсяДоступна Вид доступа __device__device R __constant__devicedevice / hostR / W __shared__deviceblock RW / __syncthreads()

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 10 Синтаксис CUDA Встроенные переменные Сравним CPU код vs CUDA kernel: __global__ void incKernel ( float * data ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; data [idx] = data [idx] + 1.0f; } float * data; for ( int i = 0; i < n; i++ ) { data [x] = data[i] + 1.0f; } Пусть nx = 2048 Пусть в блоке 256 потоков кол-во блоков = 2048 / 256 = 8 [ ][ == 256][ ]

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 11 Синтаксис CUDA Встроенные переменные В любом CUDA kernele доступны: –dim3 gridDim; –uint3 blockIdx; –dim3 blockDim; –uint3 threadIdx; –int warpSize; dim3 – встроенный тип, который используется для задания размеров kernelа По сути – это uint3.

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 12 Основы CUDA host API Два API Низкоуровневый driver API (cu*) Высокоуровневый runtime API (cuda*) Реализован через driver API Не требуют явной инициализации Все функции возвращают значение типа cudaError_t cudaSuccess в случае успеха

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 13 Основы CUDA API Многие функции API асинхронны: Запуск ядра Копирование при помощи функций *Async Копирование device device Инициализация памяти

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 14 CUDA Compute Capability Возможности GPU обозначаются при помощи Compute Capability, например 1.1 Старшая цифра соответствует архитектуре Младшая – небольшим архитектурным изменениям Можно получить из полей major и minor структуры cudaDeviceProp

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 15 Тип памяти ДоступУровень выделения Скорость работы РегистрыR/WPer-thread Высокая(on-chip) ЛокальнаяR/WPer-thread Низкая (DRAM) SharedR/WPer-block Высокая(on-chip) ГлобальнаяR/WPer-grid Низкая (DRAM) ConstantR/OPer-grid Высокая(L1 cache) TextureR/OPer-grid Высокая(L1 cache) Типы памяти в CUDA

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 16 Типы памяти в CUDA Самая быстрая – shared (on-chip) и регистры Самая медленная – глобальная (DRAM) Для ряда случаев можно использовать кэшируемую константную и текстурную память Доступ к памяти в CUDA идет отдельно для каждой половины warpа (half-warp)

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 17 Объединение (coalescing) для GPU с CC 1.0/1.1 Нити обращаются к 32-битовым словам, давая 64-байтовый блок 64-битовым словам, давая 128-байтовый блок Все 16 слов лежат в пределах блока k-ая нить half-warpа обращается к k-му слову блока

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 18 Объединение (coalescing) для GPU с CC 1.0/1.1 Coalescing

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 19 Объединение (coalescing) для GPU с CC 1.0/1.1 Not Coalescing

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 20 Объединение (coalescing) для GPU с CC 1.2/1.3 Нити обращаются к 8-битовым словам, дающим один 32- байтовый сегмент 16-битовым словам, дающим один 64- байтовый сегмент 32-битовым словам, дающим один 128- байтовый сегмент Получающийся сегмент выровнен по своему размеру

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 21 Объединение (coalescing) Если хотя бы одно условие не выполнено 1.0/1.1 – 16 отдельных транзакций 1.2/1.3 – объединяет их в блоки (2,3,…) и для каждого блока проводится отдельная транзакция Для 1.2/1.3 порядок в котором нити обращаются к словам внутри блока не имеет значения (в отличии от 1.0/1.1)

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 22 Объединение (coalescing) Можно добиться заметного увеличения скорости работы с памятью Лучше использовать не массив структур, а набор массивов отдельных компонент – это позволяет использовать coalescing

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 23 «Особенности программирования под GPU»

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 24 Введение Процесс разработки программ CUDA Портирование части приложения Общие рекомендации по оптимизации Инструментарий Работа с различными типами памяти Паттерны программирования на CUDA Разное

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 25 Портирование части приложения Определение класса портируемой задачи –Уровень параллелизма. SIMD –Классы задач, которые в общем случае невозможно распараллелить: сжатие данных IIR-фильтры другие рекурсивные алгоритмы

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 26 Портирование части приложения

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 27 Общие рекомендации по оптимизации Переосмысление задачи в терминах параллельной обработки данных –Выявляйте параллелизм –Максимизируйте интенсивность вычислений Иногда выгоднее пересчитать чем сохранить –Избегайте лишних транзакций по памяти Особое внимание особенностям работы с различными видами памяти (об этом дальше) Эффективное использование вычислительной мощи –Разбивайте вычисления с целью поддержания сбалансированной загрузки SMов –Параллелизм потоков vs. параллелизм по данным

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 28 Инструментарий: Компилятор

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 29 Инструментарий: Компилятор Статическая компиляция: IDE(MS Visual Studio + cuda.rules), Makefile, CL PTX JIT-компиляция

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 30 Инструментарий: Компилятор

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 31 Инструментарий: Отладчик GPU debugger –Wednesday, April 08: Today NVIDIA announces an industry milestone for GPU Computing. With CUDA 2.2 beta we are including the industries 1st GPU HW Debugger to our developer community. GPU emulation –-deviceemu D_DEVICEEMU –Запускает по одному host-процессу на каждый CUDA-поток –Работоспособность в режиме эмуляции не всегда кореллирует с работоспособностю на GPU Два инструмента не конкурируют, а дополняют друг друга –Один из интересных сценариев: Boundchecker + Emulation

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 32 Инструментарий: Отладчик Недостатки эмуляции –Часто работает очень медленно –Неумышленное разыменование указателей GPU на стороне CPU или наоборот –Результаты операций с плавающей точкой CPU и «настоящего» GPU почти всегда различаются из- за: Разного порядка выполняемых операций Разных допустимых ошибок результатов Использования большей точности при расчёте промежуточных результатов на CPU

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 33 Инструментарий: Профилировщик CUDA Profiler, позволяет отслеживать: –Время исполнения на CPU и GPU в микросекундах –Конфигурацию grid и thread block –Количество статической разделяемой памяти на блок –Количество регистров на блок –Коэффициент занятости GPU (Occupancy) –Количество объединенных и индивидуальных запросов к глобальной памяти (coalescing) –Количество дивергентных путей исполнения (branching) –Количество выполненных инструкций –Количество запущенных блоков Вся эта информация собирается с первого SM или TPC. Профилирование Uber-kernelов с осторожностью

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 34 Работа с константной памятью Быстрая, кешируемая, только для чтения Данные должны быть записаны до вызова кернела (например при помощи cudaMemcpyToSymbol) Всего 64Kb (Tesla) Объявление при помощи слова __constant__ Доступ из device кода простой адресацией Срабатывает за 4 такта на один адрес внутри варпа –4 такта на всю транзакцию в случае если все потоки внутри варпа читают один адрес –В худшем случае 64 такта

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 35 Работа с текстурной памятью Быстрая, кешируемая в 2-х измерениях, только для чтения Данные должны быть записаны при помощи cudaMemcpyToArray, либо возможно прикрепление к глобальной памяти через cudaBindTexture2D Объявление при помощи текстурных ссылок Доступ из device кода при помощи tex1D, tex2D, tex1Dfetch Лучшая производительность при условии что потоки одного варпа обращаются к локализованной окрестности в 2D

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 36 Работа с глобальной памятью Медленная, некешируемая (G80), чтение/запись Запись данных с/на хост через cudaMemcpy* Транзакции по PCI-e медленные: макс. 4GB/s vs. 80 GB/s при копировании device-device Возможность асинхронных транзакций Ускорение транзакций путем выделения host page-locked памяти ( cudaMallocHost ) Объявление при помощи слова __global__ Доступ простой индексацией Время доступа от 400 до 600 тактов на транзакцию – высокая латентность

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 37 Работа с глобальной памятью Coalescing, Compute Capability 1.0, потоков. Типы транзакций: –4-байтовые слова, одна 64-байтовая транзакция –8-байтовые слова, одна 128-байтовая транзакция –16-байтовые слова, две 128-байтовых транзакции Все 16 элементов должны лежать в едином сегменте, размер и выравнивание которого совпадает с размером транзакции Строгий порядок доступа: k-й поток обращается к k-му элементу в сегменте При нарушении порядка вместо одной транзакции получается 16 Некоторые из потоков могут не участвовать

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 38 Работа с глобальной памятью Coalescing, Compute Capability 1.0, 1.1 CoalescingNo coalescing

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 39 Работа с глобальной памятью Coalescing, Compute Capability 1.2, 1.3 Объединенная транзакция получается, если все элементы лежат в сегментах: –размера 32 байта, потоки обращаются к 1-байтовым элементам –размера 64 байта, потоки обращаются к 2-байтовым элементам –размера 128 байт, потоки обращаются к 4- и 8-байтовым элементам Нестрогий порядок доступа. Возможно обращение несколькими потоками к одному адресу При выходе за границы сегмента число транзакций увеличивается минимально

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 40 Работа с глобальной памятью Coalescing, Compute Capability 1.2, 1.3

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 41 Работа с глобальной памятью Coalescing. Рекомендации Используйте cudaMallocPitch для работы с 2D- массивами Конфигурируйте блоки с большей протяженностью по x Параметризуйте конфигурацию, экспериментируйте В сложных случаях используйте привязку сегмента глобальной памяти к текстуре в случае если Compute Capability < 1.2 –cudaBindTexture, tex1Dfetch –cudaBindTexture2D, tex2D

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 42 Работа с разделяемой памятью Быстрая, некешируемая, чтение/запись Объявление при помощи слова __shared__ Доступ из device кода при помощи индексирования Самый быстрый тип памяти после регистров, низкая латентность доступа Можно рассматривать как полностью открытый L1- кеш При работе с разделяемой памятью следует помнить о ее разбиении на банками памяти

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 43 Работа с разделяемой памятью Банки памяти Память разделена на 16 банков памяти, по числу потоков в варпе Каждый банк может обратиться к одному адресу за 1 такт Максимальное число адресов, к которым может обратиться память одновременно совпадает с числом банков Одновременное обращение нескольких потоков из одного полуварпа к одному банку приводит к конфликту банков и сериализации запросов (кроме broadcast) Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 44 Работа с разделяемой памятью Банки памяти Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 Thread 15 Thread 7 Thread 6 Thread 5 Thread 4 Thread 3 Thread 2 Thread 1 Thread 0 Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 Thread 15 Thread 7 Thread 6 Thread 5 Thread 4 Thread 3 Thread 2 Thread 1 Thread 0 Прямой доступ Смешанный доступ 1:1 Доступ без конфликтов банков

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 45 Работа с разделяемой памятью Банки памяти 2-кратный конфликт 8-кратный конфликт Доступ с конфликтами банков Thread 11 Thread 10 Thread 9 Thread 8 Thread 4 Thread 3 Thread 2 Thread 1 Thread 0 Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 Thread 15 Thread 7 Thread 6 Thread 5 Thread 4 Thread 3 Thread 2 Thread 1 Thread 0 Bank 9 Bank 8 Bank 15 Bank 7 Bank 2 Bank 1 Bank 0 x8

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 46 Паттерны программирования на CUDA Приоритеты оптимизации Объединение запросов к глобальной памяти –Ускорение до 20 раз –Стремление к локальности Использование разделяемой памяти –Высокая скорость работы –Удобство взаимодействия потоков Эффективное использование параллелизма –GPU не должен простаивать –Преобладание вычислений над операциями с памятью –Много блоков и потоков в блоке Банк-конфликты –Если избавление от 4-кратных конфликтов банков влечет увеличение числа инструкций, то данный вид оптимизации можно не делать

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 47 Паттерны программирования на CUDA Сценарий работы с shared памятью 1. Загрузка данных из глобальной памяти в разделяемой 2.__syncthreads(); 3. Обработка данных в разделяемой памяти 4.__syncthreads(); //если требуется 5. Сохранение результатов в глобальной памяти Шаги 2–4 могут быть обрамлены в условия и циклы Шаг 4 может быть ненужен в случае если выходные данные независимы между собой

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 48 Паттерны программирования на CUDA Копирование global shared: 32-bit dim3 block(64); __shared__ float dst[64]; __global__ void kernel(float *data) {//coalescing, no bank conflicts dst[threadIdx.x] = data[threadIdx.x]; }

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 49 Паттерны программирования на CUDA Копирование global shared: 8-bit dim3 block(64); __shared__ byte dst[64]; __global__ void kernel_bad(byte *data) {//no coalescing, 4-way bank conflicts present dst[threadIdx.x] = data[threadIdx.x]; } __global__ void kernel_good(byte *data) {//coalescing, no bank conflicts, no branching if (threadIdx.x < 16) { int tx = threadIdx.x * 4; *((int *)(dst + tx)) = *((int *)(data + tx)); }

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 50 Паттерны программирования на CUDA Обработка в shared памяти __shared__ byte buf[64]; dim3 block(64); Независимая обработка элементов. Прямой доступ будет вызывать 4- кратный конфликт банков. Задача: переформировать потоки в 4 группы по 16 индексов так, чтобы при новой косвенной адресации не было конфликтов банков.

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 51 Обработка в shared памяти __device__ int permute64by4(int t) { return (t >> 4) + ((t & 0xF) << 2); } Одно из решений: Thread 63 Thread 32 Thread 31 Thread 16 Thread 15 Thread 1 Thread 0 Bank 15 Bank 0 Bank 15 Bank 0 Bank 15 Bank 1 Bank 0 Index 63 Index 2 Index 61 Index 1 Index 60 Index 4 Index 0

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 52 Обработка в shared памяти Независимая обработка элементов. Прямой доступ будет вызывать 16-кратный конфликт банков. Задача: свести число банк-конфликтов до нуля. __shared__ int buf[16][16]; dim3 block(16,16);

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 53 Обработка в shared памяти Одно из решений: __shared__ int buf[16][17]; dim3 block(16,16); Bank Indices without Padding Bank Indices with Padding

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 54 Разное Конфигурация gridDim и blockDim возможно во время исполнения: void callKernel(dim3 grid, dim3 threads) { kernel >>(); }

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 55 Конфигурация gridDim и blockDim возможно во время исполнения: void callKernel(dim3 grid, dim3 threads) { kernel >>(); } Разное

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 56 Разное __mul24 и __umul24 работают быстрее, чем * Возможно увеличение числа регистров после применения На будущих архитектурах ситуация может развернуться наоборот и __mul24 станет медленнее В остальном целочисленная арифметика работает примерно с такой же скоростью, как и с плавающей точкой (за исключением целочисленного деления)

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 57 Разное Математика FPU (на GPU в частности) не ассоциативна (x+y)+z не всегда равно x+(y+z) Например при x = 10^30, y = -10^30, z = 1

Московский авиационный институт (государственный технический университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 58 Вопросы