Что такое findslide.org?

FindSlide.org - это сайт презентаций, докладов, шаблонов в формате PowerPoint.


Для правообладателей

Обратная связь

Email: Нажмите что бы посмотреть 

Яндекс.Метрика

Презентация на тему Вопросы программирования и оптимизации приложений на CUDA

Содержание

СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыРазное
Вопросы программирования и оптимизации приложений на CUDA. Лекторы:Обухов А.Н. (Обухов А.Н. (Nvidia)Боресков СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыРазное СодержаниеПроцесс разработки программ CUDAПортирование части приложенияОбщие рекомендации по оптимизации ИнструментарийРабота с различными Процесс разработки программ CUDAПортирование части приложенияОпределение класса портируемой задачиУровень параллелизма. SIMDКлассы задач, Процесс разработки программ CUDAПортирование части приложения СодержаниеПроцесс разработки программ CUDAПортирование части приложенияОбщие рекомендации по оптимизации ИнструментарийРабота с различными Переосмысление задачи в терминах параллельной обработки данныхВыявляйте параллелизмМаксимизируйте интенсивность вычисленийИногда выгоднее пересчитать Процесс разработки программ CUDA Общие рекомендации по оптимизацииOccupancyПокрытие латентностей: инструкции потока выполняются Процесс разработки программ CUDA Общие рекомендации по оптимизацииOccupancyУвеличение занятости приводит к лучшему СодержаниеПроцесс разработки программ CUDAПортирование части приложенияОбщие рекомендации по оптимизации ИнструментарийРабота с различными Процесс разработки программ CUDAИнструментарий: Компилятор Процесс разработки программ CUDAИнструментарий: КомпиляторСтатическая компиляция: IDE(MS Visual Studio + cuda.rules), Makefile, CLPTX JIT-компиляция Процесс разработки программ CUDAИнструментарий: Компилятор Процесс разработки программ CUDAИнструментарий: ОтладчикGPU debuggerWednesday, April 08: Today NVIDIA announces an Достоинства эмуляцииИсполняемый файл, скомпилированный в режиме эмуляции работает целиком на CPUНе требуется Недостатки эмуляцииЧасто работает очень медленноНеумышленное разыменование указателей GPU на стороне CPU или Visual Profiler Profiler Counter Plot Процесс разработки программ CUDAИнструментарий: ПрофилировщикCUDA Profiler, позволяет отслеживать:Время исполнения на CPU и ОптимизацияOccupancy Calculator Spreadsheet СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиКонстантнаяТекстурнаяГлобальнаяРазделяемаяПаттерны программирования на CUDAСтратегии распределения работыРазное Работа с константной памятьюБыстрая, кешируемая, только для чтенияДанные должны быть записаны до СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиКонстантнаяТекстурнаяГлобальнаяРазделяемаяПаттерны программирования на CUDAСтратегии распределения работыРазное Работа с текстурной памятьюБыстрая, кешируемая в 2-х измерениях, только для чтенияДанные должны СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиКонстантнаяТекстурнаяГлобальнаяРазделяемаяПаттерны программирования на CUDAСтратегии распределения работыРазное Работа с глобальной памятьюМедленная, некешируемая (G80), чтение/записьЗапись данных с/на хост через cudaMemcpy*Транзакции Работа с глобальной памятьюCoalescing, Compute Capability 1.0, 1.116 потоков. Типы транзакций:4-байтовые слова, Работа с глобальной памятьюCoalescing, Compute Capability 1.0, 1.1CoalescingNo coalescing Работа с глобальной памятьюCoalescing, Compute Capability 1.2, 1.3Объединенная транзакция получается, если все Работа с глобальной памятьюCoalescing, Compute Capability 1.2, 1.3 Работа с глобальной памятьюCoalescing. РекомендацииИспользуйте cudaMallocPitch для работы с 2D-массивамиКонфигурируйте блоки с Коалесинг СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиКонстантнаяТекстурнаяГлобальнаяРазделяемаяПаттерны программирования на CUDAСтратегии распределения работыРазное Работа с разделяемой памятьюБыстрая, некешируемая, чтение/записьОбъявление при помощи слова __shared__Доступ из device Память разделена на 16 банков памяти, по числу потоков в варпеКаждый банк Работа с разделяемой памятьюБанки памятиПрямой доступСмешанный доступ 1:1Доступ без конфликтов банков Работа с разделяемой памятьюБанки памяти2-кратный конфликт8-кратный конфликтДоступ с конфликтами банков СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAПриоритеты оптимизацииСценарий Паттерны программирования на CUDAПриоритеты оптимизацииОбъединение запросов к глобальной памятиУскорение до 20 разСтремление СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAПриоритеты оптимизацииСценарий Паттерны программирования на CUDAСценарий работы с shared памятьюЗагрузка данных из глобальной памяти СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAПриоритеты оптимизацииСценарий Паттерны программирования на CUDAКопирование global shared: 32-bitdim3 block(64);__shared__ float dst[64];__global__ void kernel(float Паттерны программирования на CUDAКопирование global shared: 8-bitdim3 block(64);__shared__ byte dst[64];__global__ void kernel_bad(byte СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAПриоритеты оптимизацииСценарий Паттерны программирования на CUDAОбработка в shared памятиНезависимая обработка элементов. Прямой доступ будет Паттерны программирования на CUDAОбработка в shared памяти__device__ int permute64by4(int t){  return Паттерны программирования на CUDAОбработка в shared памяти (2)Независимая обработка элементов. Прямой доступ Паттерны программирования на CUDAОбработка в shared памяти (2)Одно из решений:__shared__ int buf[16][17];dim3 СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное Стратегии распределения работыЗадачи с нерегулярным параллелизмомПеременное кол-во итераций Большое кол-во ветвлений Стратегии распределения работы: C & CРазделить ядра на более простыеПозволяет выявить bottleneckУвеличивает СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное Стратегии распределения работы: Uber-kernelUber-kernelif ( A ){	  Exec_A();}Else if ( B ){  Exec_B();}… Стратегии распределения работы: Uber-kernel (2)timeBlocks 0Blocks 1Blocks 2Blocks 3kernel1kernel2 Стратегии распределения работы: Uber-kernel (3)timeBlocks 0Blocks 1Blocks 2Blocks 3kernel1kernel2 Стратегии распределения работы: Uber-kernel (3)timeBlocks 0Blocks 1Blocks 2Blocks 3if (A) kernel1if (B) СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное Стратегии распределения работы Стратегии распределения работы: Persistent threadstimeWarp 0Warp 1Warp 2Warp 3Block 0Block 1 Стратегии распределения работы: Persistent threadstimeWarp 0Warp 1Warp 2Warp 3Block 0Block 1 Стратегии распределения работы: Persistent threads (2)timeWarp 0Warp 1Warp 2Warp 3Block 0Block 1 Стратегии распределения работы: Persistent threads (3) timeWarp 0Warp 1Warp 2Warp 3Block 0 СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыРазное ВетвлениеЕсли происходит ветвление внутри варпа, то разные ветви исполнения сериализуютсяУвеличивается общее количество Ветвление Инструкции ОптимизацияPTX PTXПромежуточный ассемблер может показать много интересного--ptxas-options=-v PTXПромежуточный ассемблер может показать много интересного--ptxas-options=-v PTXПромежуточный ассемблер может показать много интересного--keep PTXПромежуточный ассемблер может показать много интересного--keep ИнструкцииСледить за ветвлениемЗаменить часть вычислений на look-up таблицуИнтринсики__sinf(); __cosf(); expf()__[u]mul24()__fdividef()__[u]sad() __mul24 и __umul24 работают быстрее, чем *Возможно увеличение числа регистров после примененияНа Конфигурация gridDim и blockDim возможно во время исполнения:void callKernel(dim3 grid, dim3 threads){ Исользование templatetemplate __global__ void kernel(){  int x = threadIdx.x + blockIdx.x РазноеМатематика FPU (на GPU в частности) не ассоциативна(x+y)+z не всегда равно x+(y+z)Например Ресурсы нашего курса	CUDA.CS.MSU.SUМесто для вопросов и дискуссийМесто для материалов нашего курсаМесто для Вопросы Спасибо!Александр ГужваАнтон ОбуховВладимир ФроловДмитрий ВатолинДмитрий МикушинЕвгений ПерепелкинМихаил СмирновНиколай Сахарных
Слайды презентации

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

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

программирования на CUDA
Стратегии распределения работы
Разное


Слайд 3 Содержание
Процесс разработки программ CUDA
Портирование части приложения
Общие рекомендации по

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

оптимизации
Инструментарий
Работа с различными типами памяти
Паттерны программирования на CUDA
Стратегии

распределения работы
Разное

Слайд 4 Процесс разработки программ CUDA
Портирование части приложения
Определение класса портируемой

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

задачи
Уровень параллелизма. SIMD
Классы задач, которые в общем случае невозможно

распараллелить

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










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




Слайд 6 Содержание
Процесс разработки программ CUDA
Портирование части приложения
Общие рекомендации по

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

оптимизации
Инструментарий
Работа с различными типами памяти
Паттерны программирования на CUDA
Стратегии

распределения работы
Разное

Слайд 7 Переосмысление задачи в терминах параллельной обработки данных
Выявляйте параллелизм
Максимизируйте

Переосмысление задачи в терминах параллельной обработки данныхВыявляйте параллелизмМаксимизируйте интенсивность вычисленийИногда выгоднее

интенсивность вычислений
Иногда выгоднее пересчитать чем сохранить
Избегайте лишних транзакций по

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

Процесс разработки программ CUDA

Общие рекомендации по оптимизации


Слайд 8 Процесс разработки программ CUDA Общие рекомендации по оптимизации
Occupancy
Покрытие латентностей:

Процесс разработки программ CUDA Общие рекомендации по оптимизацииOccupancyПокрытие латентностей: инструкции потока

инструкции потока выполняются последовательно
Исполнение других потоков необходимо для

покрытия латентностей
Занятость: отношение активных варпов к максимально возможному
В архитектуре Tesla 32 варпа на SM



Слайд 9 Процесс разработки программ CUDA Общие рекомендации по оптимизации
Occupancy
Увеличение занятости

Процесс разработки программ CUDA Общие рекомендации по оптимизацииOccupancyУвеличение занятости приводит к

приводит к лучшему покрытию латентностей
После определенной точки (~50%), происходит

насыщение
Занятость ограничена достыпными ресурсами:
Регистры
Разделяемая память


Слайд 10 Содержание
Процесс разработки программ CUDA
Портирование части приложения
Общие рекомендации по

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

оптимизации
Инструментарий
Работа с различными типами памяти
Паттерны программирования на CUDA
Стратегии

распределения работы
Разное

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

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

Слайд 12 Процесс разработки программ CUDA
Инструментарий: Компилятор
Статическая компиляция: IDE(MS Visual

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

Studio + cuda.rules), Makefile, CL
PTX JIT-компиляция


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

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

Слайд 14 Процесс разработки программ CUDA
Инструментарий: Отладчик
GPU debugger
Wednesday, April 08:

Процесс разработки программ CUDAИнструментарий: ОтладчикGPU debuggerWednesday, April 08: Today NVIDIA announces

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

Слайд 15 Достоинства эмуляции
Исполняемый файл, скомпилированный в режиме эмуляции работает

Достоинства эмуляцииИсполняемый файл, скомпилированный в режиме эмуляции работает целиком на CPUНе

целиком на CPU
Не требуется драйвер CUDA и GPU
Каждый поток

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

Процесс разработки программ CUDA

Инструментарий: Отладчик


Слайд 16 Недостатки эмуляции
Часто работает очень медленно
Неумышленное разыменование указателей GPU

Недостатки эмуляцииЧасто работает очень медленноНеумышленное разыменование указателей GPU на стороне CPU

на стороне CPU или наоборот
Результаты операций с плавающей точкой

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

Процесс разработки программ CUDA

Инструментарий: Отладчик


Слайд 17 Visual Profiler

Visual Profiler

Слайд 18 Profiler Counter Plot

Profiler Counter Plot

Слайд 19 Процесс разработки программ CUDA
Инструментарий: Профилировщик
CUDA Profiler, позволяет отслеживать:
Время

Процесс разработки программ CUDAИнструментарий: ПрофилировщикCUDA Profiler, позволяет отслеживать:Время исполнения на CPU

исполнения на CPU и GPU в микросекундах
Конфигурацию grid и

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

Слайд 20 Оптимизация
Occupancy Calculator Spreadsheet

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

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

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

программирования на CUDA
Стратегии распределения работы
Разное


Слайд 22 Работа с константной памятью
Быстрая, кешируемая, только для чтения
Данные

Работа с константной памятьюБыстрая, кешируемая, только для чтенияДанные должны быть записаны

должны быть записаны до вызова кернела (например при помощи

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

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

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

программирования на CUDA
Стратегии распределения работы
Разное


Слайд 24 Работа с текстурной памятью
Быстрая, кешируемая в 2-х измерениях,

Работа с текстурной памятьюБыстрая, кешируемая в 2-х измерениях, только для чтенияДанные

только для чтения
Данные должны быть записаны при помощи cudaMemcpyToArray,

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

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

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

программирования на CUDA
Стратегии распределения работы
Разное


Слайд 26 Работа с глобальной памятью
Медленная, некешируемая (G80), чтение/запись
Запись данных

Работа с глобальной памятьюМедленная, некешируемая (G80), чтение/записьЗапись данных с/на хост через

с/на хост через cudaMemcpy*
Транзакции по PCI-e медленные: макс. 4GB/s

vs. 80 GB/s при копировании device-device
Возможность асинхронных транзакций
Ускорение транзакций путем выделения host page-locked памяти (cudaMallocHost)
Объявление при помощи слова __global__
Доступ простой индексацией
Время доступа от 400 до 600 тактов на транзакцию – высокая латентность

Слайд 27 Работа с глобальной памятью
Coalescing, Compute Capability 1.0, 1.1
16

Работа с глобальной памятьюCoalescing, Compute Capability 1.0, 1.116 потоков. Типы транзакций:4-байтовые

потоков. Типы транзакций:
4-байтовые слова, одна 64-байтовая транзакция
8-байтовые слова, одна

128-байтовая транзакция
16-байтовые слова, две 128-байтовых транзакции
Все 16 элементов должны лежать в едином сегменте, размер и выравнивание которого совпадает с размером транзакции
Строгий порядок доступа: k-й поток обращается к k-му элементу в сегменте
При нарушении порядка вместо одной транзакции получается 16
Некоторые из потоков могут не участвовать

Слайд 28 Работа с глобальной памятью
Coalescing, Compute Capability 1.0, 1.1
Coalescing
No

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

coalescing


Слайд 29 Работа с глобальной памятью
Coalescing, Compute Capability 1.2, 1.3
Объединенная

Работа с глобальной памятьюCoalescing, Compute Capability 1.2, 1.3Объединенная транзакция получается, если

транзакция получается, если все элементы лежат в сегментах:
размера 32

байта, потоки обращаются к 1-байтовым элементам
размера 64 байта, потоки обращаются к 2-байтовым элементам
размера 128 байт, потоки обращаются к 4- и 8-байтовым элементам
Нестрогий порядок доступа. Возможно обращение несколькими потоками к одному адресу
При выходе за границы сегмента число транзакций увеличивается минимально

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

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

Слайд 31 Работа с глобальной памятью
Coalescing. Рекомендации
Используйте cudaMallocPitch для работы

Работа с глобальной памятьюCoalescing. РекомендацииИспользуйте cudaMallocPitch для работы с 2D-массивамиКонфигурируйте блоки

с 2D-массивами
Конфигурируйте блоки с большей протяженностью по x
Параметризуйте конфигурацию,

экспериментируйте
В сложных случаях используйте привязку сегмента глобальной памяти к текстуре в случае если Compute Capability < 1.2
cudaBindTexture, tex1Dfetch
cudaBindTexture2D, tex2D

Слайд 32 Коалесинг

Коалесинг

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

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

программирования на CUDA
Стратегии распределения работы
Разное


Слайд 34 Работа с разделяемой памятью
Быстрая, некешируемая, чтение/запись
Объявление при помощи

Работа с разделяемой памятьюБыстрая, некешируемая, чтение/записьОбъявление при помощи слова __shared__Доступ из

слова __shared__
Доступ из device кода при помощи индексирования
Самый быстрый

тип памяти после регистров, низкая латентность доступа
Можно рассматривать как полностью открытый L1-кеш
При работе с разделяемой памятью следует помнить о ее разбиении на банками памяти

Слайд 35 Память разделена на 16 банков памяти, по числу

Память разделена на 16 банков памяти, по числу потоков в варпеКаждый

потоков в варпе
Каждый банк может обратиться к одному адресу

за 1 такт
Максимальное число адресов, к которым может обратиться память одновременно совпадает с числом банков
Одновременное обращение нескольких потоков из одного полуварпа к одному банку приводит к конфликту банков и сериализации запросов (кроме broadcast)

Работа с разделяемой памятью

Банки памяти


Слайд 36 Работа с разделяемой памятью
Банки памяти
Прямой доступ
Смешанный доступ 1:1
Доступ

Работа с разделяемой памятьюБанки памятиПрямой доступСмешанный доступ 1:1Доступ без конфликтов банков

без конфликтов банков


Слайд 37 Работа с разделяемой памятью
Банки памяти
2-кратный конфликт
8-кратный конфликт
Доступ с

Работа с разделяемой памятьюБанки памяти2-кратный конфликт8-кратный конфликтДоступ с конфликтами банков

конфликтами банков


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

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

программирования на CUDA
Приоритеты оптимизации
Сценарий работы с shared памятью
Копирование global

<-> shared
Обработка в shared памяти
Стратегии распределения работы
Разное

Слайд 40 Паттерны программирования на CUDA
Приоритеты оптимизации
Объединение запросов к глобальной

Паттерны программирования на CUDAПриоритеты оптимизацииОбъединение запросов к глобальной памятиУскорение до 20

памяти
Ускорение до 20 раз
Стремление к локальности
Использование разделяемой памяти
Высокая скорость

работы
Удобство взаимодействия потоков
Эффективное использование параллелизма
GPU не должен простаивать
Преобладание вычислений над операциями с памятью
Много блоков и потоков в блоке
Банк-конфликты
Если избавление от 4-кратных конфликтов банков влечет увеличение числа инструкций, то данный вид оптимизации можно не делать

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

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

программирования на CUDA
Приоритеты оптимизации
Сценарий работы с shared памятью
Копирование global

<-> shared
Обработка в shared памяти
Стратегии распределения работы
Разное

Слайд 42 Паттерны программирования на CUDA
Сценарий работы с shared памятью
Загрузка

Паттерны программирования на CUDAСценарий работы с shared памятьюЗагрузка данных из глобальной

данных из глобальной памяти в разделяемой
__syncthreads();
Обработка данных в разделяемой

памяти
__syncthreads(); //если требуется
Сохранение результатов в глобальной памяти

Шаги 2–4 могут быть обрамлены в условия и циклы
Шаг 4 может быть ненужен в случае если выходные данные независимы между собой

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

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

программирования на CUDA
Приоритеты оптимизации
Сценарий работы с shared памятью
Копирование global

<-> shared
Обработка в shared памяти
Стратегии распределения работы
Разное

Слайд 44 Паттерны программирования на CUDA
Копирование global shared: 32-bit
dim3

Паттерны программирования на CUDAКопирование global shared: 32-bitdim3 block(64);__shared__ float dst[64];__global__ void

block(64);
__shared__ float dst[64];


__global__ void kernel(float *data)
{//coalescing, no bank conflicts

dst[threadIdx.x] = data[threadIdx.x];
}

Слайд 45 Паттерны программирования на CUDA
Копирование global shared: 8-bit
dim3

Паттерны программирования на CUDAКопирование global shared: 8-bitdim3 block(64);__shared__ byte dst[64];__global__ void

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));
}
}


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

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

программирования на CUDA
Приоритеты оптимизации
Сценарий работы с shared памятью
Копирование global

<-> shared
Обработка в shared памяти
Стратегии распределения работы
Разное

Слайд 47 Паттерны программирования на CUDA
Обработка в shared памяти
Независимая обработка

Паттерны программирования на CUDAОбработка в shared памятиНезависимая обработка элементов. Прямой доступ

элементов. Прямой доступ будет вызывать 4-кратный конфликт банков.
Задача: переформировать

потоки в 4 группы по 16 индексов так, чтобы при новой косвенной адресации не было конфликтов банков.

__shared__ byte buf[64];
dim3 block(64);


Слайд 48 Паттерны программирования на CUDA
Обработка в shared памяти
__device__ int

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

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


Слайд 49 Паттерны программирования на CUDA
Обработка в shared памяти (2)
Независимая

Паттерны программирования на CUDAОбработка в shared памяти (2)Независимая обработка элементов. Прямой

обработка элементов. Прямой доступ будет вызывать 16-кратный конфликт банков.
Задача:

свести число банк-конфликтов до нуля.

__shared__ int buf[16][16];
dim3 block(16,16);


Слайд 50 Паттерны программирования на CUDA
Обработка в shared памяти (2)
Одно

Паттерны программирования на CUDAОбработка в shared памяти (2)Одно из решений:__shared__ int

из решений:
__shared__ int buf[16][17];
dim3 block(16,16);
Bank Indices without Padding
Bank Indices

with Padding

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

СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное

программирования на CUDA
Стратегии распределения работы
Command & Conquer
Uber-kernel
Persistent threads
Разное


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

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


Большое кол-во ветвлений


Слайд 53 Стратегии распределения работы: C & C
Разделить ядра на

Стратегии распределения работы: C & CРазделить ядра на более простыеПозволяет выявить

более простые
Позволяет выявить bottleneck
Увеличивает Occupancy
Возможность перераспределять работу между ядрами


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

СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное

программирования на CUDA
Стратегии распределения работы
Command & Conquer
Uber-kernel
Persistent threads
Разное


Слайд 55 Стратегии распределения работы: Uber-kernel
Uber-kernel
if ( A )
{

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

Exec_A();
}
Else if ( B )
{
Exec_B();
}


Слайд 56 Стратегии распределения работы: Uber-kernel (2)
time






Blocks 0
Blocks 1
Blocks 2
Blocks

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

3



kernel1
kernel2


Слайд 57 Стратегии распределения работы: Uber-kernel (3)
time






Blocks 0
Blocks 1
Blocks 2
Blocks

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

3



kernel1
kernel2


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

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

3

if (A) kernel1
if (B) kernel2






Blocks 0
Blocks 1
Blocks 2
Blocks 3



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

СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное

программирования на CUDA
Стратегии распределения работы
Command & Conquer
Uber-kernel
Persistent threads
Разное


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

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

Слайд 61 Стратегии распределения работы: Persistent threads
time






Warp 0
Warp 1
Warp 2
Warp

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

3



Block 0
Block 1


Слайд 62 Стратегии распределения работы: Persistent threads
time






Warp 0
Warp 1
Warp 2
Warp

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

3



Block 0
Block 1


Слайд 63 Стратегии распределения работы: Persistent threads (2)
time






Warp 0
Warp 1
Warp

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

2
Warp 3


Block 0
Block 1


Слайд 64 Стратегии распределения работы: Persistent threads (3)
time






Warp 0
Warp

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

1
Warp 2
Warp 3



Block 0


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

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

программирования на CUDA
Стратегии распределения работы
Разное


Слайд 66 Ветвление
Если происходит ветвление внутри варпа, то разные ветви

ВетвлениеЕсли происходит ветвление внутри варпа, то разные ветви исполнения сериализуютсяУвеличивается общее

исполнения сериализуются

Увеличивается общее количество инструкций

Если ветвление происходит между варпами,

то штраф минимальный

Слайд 67 Ветвление

Ветвление

Слайд 68 Инструкции

Инструкции

Слайд 69 Оптимизация
PTX

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

Слайд 70 PTX
Промежуточный ассемблер может показать много интересного
--ptxas-options=-v





PTXПромежуточный ассемблер может показать много интересного--ptxas-options=-v

Слайд 71 PTX
Промежуточный ассемблер может показать много интересного
--ptxas-options=-v





PTXПромежуточный ассемблер может показать много интересного--ptxas-options=-v

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





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

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




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

Слайд 74 Инструкции
Следить за ветвлением
Заменить часть вычислений на look-up таблицу
Интринсики
__sinf();

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

__cosf(); expf()
__[u]mul24()
__fdividef()
__[u]sad()



Слайд 75 __mul24 и __umul24 работают быстрее, чем *
Возможно увеличение

__mul24 и __umul24 работают быстрее, чем *Возможно увеличение числа регистров после

числа регистров после применения
На будущих архитектурах ситуация может развернуться

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

Разное


Слайд 76 Конфигурация gridDim и blockDim возможно во время исполнения:


void

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

callKernel(dim3 grid, dim3 threads)
{
kernel();
}
Размеры CTA и

GRID

Слайд 77 Исользование template


template
__global__

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

void kernel()
{
int x = threadIdx.x + blockIdx.x

* tx;
}

void callKernel(dim3 grid)
{
kernel<16, 16, 1><<>>();
}

Шаблоны


Слайд 78 Разное
Математика FPU (на GPU в частности) не ассоциативна

(x+y)+z

РазноеМатематика FPU (на GPU в частности) не ассоциативна(x+y)+z не всегда равно

не всегда равно x+(y+z)

Например при x = 10^30, y

= -10^30, z = 1

Слайд 79 Ресурсы нашего курса
CUDA.CS.MSU.SU
Место для вопросов и дискуссий
Место для

Ресурсы нашего курса	CUDA.CS.MSU.SUМесто для вопросов и дискуссийМесто для материалов нашего курсаМесто

материалов нашего курса
Место для ваших статей!
Если вы нашли какой-то

интересный подход!
Или исследовали производительность разных подходов и знаете, какой из них самый быстрый!
Или знаете способы сделать работу с CUDA проще!
www.steps3d.narod.ru
www.nvidia.ru


Слайд 80 Вопросы

Вопросы

  • Имя файла: voprosy-programmirovaniya-i-optimizatsii-prilozheniy-na-cuda.pptx
  • Количество просмотров: 127
  • Количество скачиваний: 0
- Предыдущая Обмен белков – 4