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

Нажмите для полного просмотра!

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

Содержание ▲

Вы можете ознакомиться и скачать Вопросы программирования и оптимизации приложений на CUDA. Лекторы: Обухов А.Н. (Nvidia) Боресков А.В. (ВМиК МГУ) Харламов А.А. (Nvidia) . Презентация содержит 81 слайдов. Презентации для любого класса можно скачать бесплатно. Если материал и наш сайт презентаций Вам понравились – поделитесь им с друзьями с помощью социальных кнопок и добавьте в закладки в своем браузере.

Слайды и текст этой презентации

Слайд 1

Описание слайда:

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

Слайд 2

Описание слайда:

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

Слайд 3

Описание слайда:

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

Слайд 4

Описание слайда:

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

Слайд 5

Описание слайда:

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

Слайд 6

Описание слайда:

Слайд 7

Описание слайда:

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

Слайд 8

Описание слайда:

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

Слайд 9

Описание слайда:

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

Слайд 10

Описание слайда:

Слайд 11

Описание слайда:

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

Слайд 12

Описание слайда:

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

Слайд 13

Описание слайда:

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

Слайд 14

Описание слайда:

Процесс разработки программ CUDA 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

Слайд 15

Описание слайда:

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

Слайд 16

Описание слайда:

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

Слайд 17

Описание слайда:

Слайд 18

Описание слайда:

Слайд 19

Описание слайда:

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

Слайд 20

Описание слайда:

Оптимизация

Слайд 21

Описание слайда:

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

Слайд 22

Описание слайда:

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

Слайд 23

Описание слайда:

Слайд 24

Описание слайда:

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

Слайд 25

Описание слайда:

Слайд 26

Описание слайда:

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

Слайд 27

Описание слайда:

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

Слайд 28

Описание слайда:

Работа с глобальной памятью

Слайд 29

Описание слайда:

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

Слайд 30

Описание слайда:

Работа с глобальной памятью

Слайд 31

Описание слайда:

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

Слайд 32

Описание слайда:

Коалесинг

Слайд 33

Описание слайда:

Слайд 34

Описание слайда:

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

Слайд 35

Описание слайда:

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

Слайд 36

Описание слайда:

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

Слайд 37

Описание слайда:

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

Слайд 38

Описание слайда:

Слайд 39

Описание слайда:

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

Слайд 40

Описание слайда:

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

Слайд 41

Описание слайда:

Слайд 42

Описание слайда:

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

Слайд 43

Описание слайда:

Слайд 44

$Паттерны программирования на CUDA dim3 block(64); __shared__ float dst[64]; __global__ void kernel(float *data) {//coalescing, no bank conflicts dst[threadIdx.x] = data[threadIdx.x]; }$

Описание слайда:

Паттерны программирования на CUDA dim3 block(64); __shared__ float dst[64]; __global__ void kernel(float *data) {//coalescing, no bank conflicts dst[threadIdx.x] = data[threadIdx.x]; }

Слайд 45

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

Описание слайда:

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

Слайд 46

Описание слайда:

Слайд 47

Описание слайда:

Паттерны программирования на CUDA

Слайд 48

$Паттерны программирования на CUDA __device__ int permute64by4(int t) { return (t >> 4) + ((t & 0xF) << 2); }$

Описание слайда:

Паттерны программирования на CUDA __device__ int permute64by4(int t) { return (t >> 4) + ((t & 0xF) << 2); }

Слайд 49

Описание слайда:

Паттерны программирования на CUDA

Слайд 50

Описание слайда:

Паттерны программирования на CUDA

Слайд 51

Описание слайда:

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

Слайд 52

Описание слайда:

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

Слайд 53

Описание слайда:

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

Слайд 54

Описание слайда:

Слайд 55

Описание слайда:

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

Слайд 56

Описание слайда:

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

Слайд 57

Описание слайда:

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

Слайд 58

Описание слайда:

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

Слайд 59

Описание слайда:

Слайд 60

Описание слайда:

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

Слайд 61

Описание слайда:

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

Слайд 62

Описание слайда:

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

Слайд 63

Описание слайда:

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

Слайд 64

Описание слайда:

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

Слайд 65

Описание слайда:

Слайд 66

Описание слайда:

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

Слайд 67

Описание слайда:

Ветвление

Слайд 68

Описание слайда:

Инструкции

Слайд 69

Описание слайда:

Оптимизация

Слайд 70

Описание слайда:

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

Слайд 71

Описание слайда:

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

Слайд 72

Описание слайда:

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

Слайд 73

Описание слайда:

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

Слайд 74

Описание слайда:

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

Слайд 75

Описание слайда:

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

Слайд 76

Описание слайда:

Размеры CTA и GRID Конфигурация gridDim и blockDim возможно во время исполнения:

Слайд 77

Описание слайда:

Шаблоны Исользование template

Слайд 78

Описание слайда:

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

Слайд 79

Описание слайда:

Ресурсы нашего курса CUDA.CS.MSU.SU Место для вопросов и дискуссий Место для материалов нашего курса Место для ваших статей! Если вы нашли какой-то интересный подход! Или исследовали производительность разных подходов и знаете, какой из них самый быстрый! Или знаете способы сделать работу с CUDA проще! www.steps3d.narod.ru www.nvidia.ru