Вопросы программирования и оптимизации приложений на CUDA. zЛекторы: yОбухов А.Н. (Nvidia)Обухов А.Н. (Nvidia) yБоресков А.В. (ВМиК МГУ)Боресков А.В. (ВМиК МГУ) yХарламов А.А. (Nvidia)Харламов А.А. (Nvidia)

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

Содержание zПроцесс разработки программ CUDA yПортирование части приложения yОбщие рекомендации по оптимизации yИнструментарий zРабота с различными типами памяти zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

Процесс разработки программ CUDA Портирование части приложения zОпределение класса портируемой задачи yУровень параллелизма. SIMD yКлассы задач, которые в общем случае невозможно распараллелить

Процесс разработки программ CUDA Портирование части приложения

Содержание zПроцесс разработки программ CUDA yПортирование части приложения yОбщие рекомендации по оптимизации yИнструментарий zРабота с различными типами памяти zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

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

Процесс разработки программ CUDA Общие рекомендации по оптимизации zOccupancy yПокрытие латентностей: инструкции потока выполняются последовательно yИсполнение других потоков необходимо для покрытия латентностей yЗанятость: отношение активных варпов к максимально возможному xВ архитектуре Tesla 32 варпа на SM

Процесс разработки программ CUDA Общие рекомендации по оптимизации zOccupancy yУвеличение занятости приводит к лучшему покрытию латентностей yПосле определенной точки (~50%), происходит насыщение yЗанятость ограничена достыпными ресурсами: xРегистры xРазделяемая память

Содержание zПроцесс разработки программ CUDA yПортирование части приложения yОбщие рекомендации по оптимизации yИнструментарий zРабота с различными типами памяти zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

Процесс разработки программ CUDA Инструментарий: Компилятор

Процесс разработки программ CUDA Инструментарий: Компилятор zСтатическая компиляция: IDE(MS Visual Studio + cuda.rules), Makefile, CL zPTX JIT-компиляция

Процесс разработки программ CUDA Инструментарий: Компилятор

Процесс разработки программ CUDA Инструментарий: Отладчик zGPU debugger yWednesday, 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. zGPU emulation y-deviceemu D_DEVICEEMU yЗапускает по одному host-процессу на каждый CUDA-поток yРаботоспособность в режиме эмуляции не всегда кореллирует с работоспособностю на GPU zДва инструмента не конкурируют, а дополняют друг друга yОдин из интересных сценариев: Boundchecker + Emulation

zДостоинства эмуляции yИсполняемый файл, скомпилированный в режиме эмуляции работает целиком на CPU xНе требуется драйвер CUDA и GPU xКаждый поток GPU эмулируется потоком CPU yПри работе в режиме эмуляции можно: xИспользовать средства отладки CPU (точки останова и т.д.) xОбращаться к любым данным GPU с CPU и наоборот Делать любые CPU-вызовы из код GPU и наоборот (например printf() ) Выявлять ситуации зависания, возникающие из-за неправильного применения __syncthreads() Процесс разработки программ CUDA Инструментарий: Отладчик

zНедостатки эмуляции yЧасто работает очень медленно yНеумышленное разыменование указателей GPU на стороне CPU или наоборот yРезультаты операций с плавающей точкой CPU и «настоящего» GPU почти всегда различаются из-за: xРазного порядка выполняемых операций xРазных допустимых ошибок результатов xИспользования большей точности при расчёте промежуточных результатов на CPU Процесс разработки программ CUDA Инструментарий: Отладчик

Visual Profiler

Profiler Counter Plot

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

Оптимизация Occupancy Calculator Spreadsheet

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти yКонстантная yТекстурная yГлобальная yРазделяемая zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

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

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти yКонстантная yТекстурная yГлобальная yРазделяемая zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

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

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти yКонстантная yТекстурная yГлобальная yРазделяемая zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

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

Работа с глобальной памятью Coalescing, Compute Capability 1.0, 1.1 z16 потоков. Типы транзакций: y4-байтовые слова, одна 64-байтовая транзакция y8-байтовые слова, одна 128-байтовая транзакция y16-байтовые слова, две 128-байтовых транзакции zВсе 16 элементов должны лежать в едином сегменте, размер и выравнивание которого совпадает с размером транзакции zСтрогий порядок доступа: k-й поток обращается к k-му элементу в сегменте zПри нарушении порядка вместо одной транзакции получается 16 zНекоторые из потоков могут не участвовать

Работа с глобальной памятью Coalescing, Compute Capability 1.0, 1.1 CoalescingNo coalescing

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

Работа с глобальной памятью Coalescing, Compute Capability 1.2, 1.3

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

Коалесинг

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти yКонстантная yТекстурная yГлобальная yРазделяемая zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

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

zПамять разделена на 16 банков памяти, по числу потоков в варпе zКаждый банк может обратиться к одному адресу за 1 такт zМаксимальное число адресов, к которым может обратиться память одновременно совпадает с числом банков zОдновременное обращение нескольких потоков из одного полуварпа к одному банку приводит к конфликту банков и сериализации запросов (кроме broadcast) Работа с разделяемой памятью Банки памяти Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 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 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 zДоступ без конфликтов банков

Работа с разделяемой памятью Банки памяти 2-кратный конфликт8-кратный конфликт zДоступ с конфликтами банков 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

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA yПриоритеты оптимизации yСценарий работы с shared памятью yКопирование global shared yОбработка в shared памяти zСтратегии распределения работы zРазное

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

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA yПриоритеты оптимизации yСценарий работы с shared памятью yКопирование global shared yОбработка в shared памяти zСтратегии распределения работы zРазное

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

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA yПриоритеты оптимизации yСценарий работы с shared памятью yКопирование global shared yОбработка в shared памяти zСтратегии распределения работы zРазное

Паттерны программирования на 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]; }

Паттерны программирования на 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)); }

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA yПриоритеты оптимизации yСценарий работы с shared памятью yКопирование global shared yОбработка в shared памяти zСтратегии распределения работы zРазное

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

Паттерны программирования на CUDA Обработка в shared памяти __device__ int permute64by4(int t) { return (t >> 4) + ((t & 0xF)

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

Паттерны программирования на CUDA Обработка в shared памяти (2) Одно из решений: __shared__ int buf[16][17]; dim3 block(16,16); Bank Indices without Padding Bank Indices with Padding

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA zСтратегии распределения работы yCommand & Conquer yUber-kernel yPersistent threads zРазное

Стратегии распределения работы zЗадачи с нерегулярным параллелизмом zПеременное кол-во итераций zБольшое кол-во ветвлений

Стратегии распределения работы: C & C zРазделить ядра на более простые yПозволяет выявить bottleneck yУвеличивает Occupancy yВозможность перераспределять работу между ядрами

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA zСтратегии распределения работы yCommand & Conquer yUber-kernel yPersistent threads zРазное

Стратегии распределения работы: Uber-kernel zUber-kernel if ( A ) { Exec_A(); } Else if ( B ) { Exec_B(); } …

Стратегии распределения работы: Uber-kernel (2) time Blocks 0 Blocks 1Blocks 2Blocks 3 kernel1 kernel2

Стратегии распределения работы: Uber-kernel (3) time Blocks 0 Blocks 1Blocks 2Blocks 3 kernel1 kernel2

Стратегии распределения работы: Uber-kernel (3) time Blocks 0 Blocks 1Blocks 2Blocks 3 if (A) kernel1 if (B) kernel2 Blocks 0 Blocks 1Blocks 2Blocks 3

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA zСтратегии распределения работы yCommand & Conquer yUber-kernel yPersistent threads zРазное

Стратегии распределения работы

Стратегии распределения работы: Persistent threads time Warp 0 Warp 1Warp 2Warp 3 Block 0 Block 1

Стратегии распределения работы: Persistent threads time Warp 0 Warp 1Warp 2Warp 3 Block 0 Block 1

Стратегии распределения работы: Persistent threads (2) time Warp 0 Warp 1Warp 2Warp 3 Block 0 Block 1

Стратегии распределения работы: Persistent threads (3) time Warp 0 Warp 1Warp 2Warp 3 Block 0

Содержание zПроцесс разработки программ CUDA zРабота с различными типами памяти zПаттерны программирования на CUDA zСтратегии распределения работы zРазное

Ветвление zЕсли происходит ветвление внутри варпа, то разные ветви исполнения сериализуются zУвеличивается общее количество инструкций zЕсли ветвление происходит между варпами, то штраф минимальный

Ветвление

Инструкции

Оптимизация PTX

zПромежуточный ассемблер может показать много интересного y--ptxas-options=-v __global__ void kernel(float *pData) { float id = (float)threadIdx.x; pData[threadIdx.x]= _sinf(id); } __global__ void kernel(float *pData) { float id = (float)threadIdx.x; pData[threadIdx.x]= sinf(id); }

PTX zПромежуточный ассемблер может показать много интересного y--ptxas-options=-v __global__ void kernel(float *pData) { float id = (float)threadIdx.x; pData[threadIdx.x]= _sinf(id); } __global__ void kernel(float *pData) { float id = (float)threadIdx.x; pData[threadIdx.x]= sinf(id); } 2.reg10.reg ! 28 bytes lmem !

PTX zПромежуточный ассемблер может показать много интересного y--keep float3 f3() { return make_float3(0,0,0); } __global__ void kernel(float3 *pData) { pData[threadIdx.x] = f3(); } float4 f4() { return make_float4(0,0,0); } __global__ void kernel(float4 *pData) { pData[threadIdx.x] = f4(); }

PTX zПромежуточный ассемблер может показать много интересного y--keep

Инструкции zСледить за ветвлением zЗаменить часть вычислений на look-up таблицу zИнтринсики y__sinf(); __cosf(); expf() y__[u]mul24() y__fdividef() y__[u]sad()

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

zКонфигурация gridDim и blockDim возможно во время исполнения: void callKernel(dim3 grid, dim3 threads) { kernel >>(); } Размеры CTA и GRID

zИсользование template template __global__ void kernel() { int x = threadIdx.x + blockIdx.x * tx; } void callKernel(dim3 grid) { kernel >>(); } Шаблоны

Разное zМатематика FPU (на GPU в частности) не ассоциативна z(x+y)+z не всегда равно x+(y+z) zНапример при x = 10^30, y = -10^30, z = 1

Ресурсы нашего курса zCUDA.CS.MSU.SUCUDA.CS.MSU.SU yМесто для вопросов и дискуссий yМесто для материалов нашего курса yМесто для ваших статей! xЕсли вы нашли какой-то интересный подход! xИли исследовали производительность разных подходов и знаете, какой из них самый быстрый! xИли знаете способы сделать работу с CUDA проще! z z

Вопросы