- Главная
- Разное
- Бизнес и предпринимательство
- Образование
- Развлечения
- Государство
- Спорт
- Графика
- Культурология
- Еда и кулинария
- Лингвистика
- Религиоведение
- Черчение
- Физкультура
- ИЗО
- Психология
- Социология
- Английский язык
- Астрономия
- Алгебра
- Биология
- География
- Геометрия
- Детские презентации
- Информатика
- История
- Литература
- Маркетинг
- Математика
- Медицина
- Менеджмент
- Музыка
- МХК
- Немецкий язык
- ОБЖ
- Обществознание
- Окружающий мир
- Педагогика
- Русский язык
- Технология
- Физика
- Философия
- Химия
- Шаблоны, картинки для презентаций
- Экология
- Экономика
- Юриспруденция
Что такое findslide.org?
FindSlide.org - это сайт презентаций, докладов, шаблонов в формате PowerPoint.
Обратная связь
Email: Нажмите что бы посмотреть
Презентация на тему Вопросы программирования и оптимизации приложений на CUDA
Содержание
- 2. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыРазное
- 3. СодержаниеПроцесс разработки программ CUDAПортирование части приложенияОбщие рекомендации
- 4. Процесс разработки программ CUDAПортирование части приложенияОпределение класса
- 5. Процесс разработки программ CUDAПортирование части приложения
- 6. СодержаниеПроцесс разработки программ CUDAПортирование части приложенияОбщие рекомендации
- 7. Переосмысление задачи в терминах параллельной обработки данныхВыявляйте
- 8. Процесс разработки программ CUDA Общие рекомендации по
- 9. Процесс разработки программ CUDA Общие рекомендации по
- 10. СодержаниеПроцесс разработки программ CUDAПортирование части приложенияОбщие рекомендации
- 11. Процесс разработки программ CUDAИнструментарий: Компилятор
- 12. Процесс разработки программ CUDAИнструментарий: КомпиляторСтатическая компиляция: IDE(MS Visual Studio + cuda.rules), Makefile, CLPTX JIT-компиляция
- 13. Процесс разработки программ CUDAИнструментарий: Компилятор
- 14. Процесс разработки программ CUDAИнструментарий: ОтладчикGPU debuggerWednesday, April
- 15. Достоинства эмуляцииИсполняемый файл, скомпилированный в режиме эмуляции
- 16. Недостатки эмуляцииЧасто работает очень медленноНеумышленное разыменование указателей
- 17. Visual Profiler
- 18. Profiler Counter Plot
- 19. Процесс разработки программ CUDAИнструментарий: ПрофилировщикCUDA Profiler, позволяет
- 20. ОптимизацияOccupancy Calculator Spreadsheet
- 21. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиКонстантнаяТекстурнаяГлобальнаяРазделяемаяПаттерны программирования на CUDAСтратегии распределения работыРазное
- 22. Работа с константной памятьюБыстрая, кешируемая, только для
- 23. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиКонстантнаяТекстурнаяГлобальнаяРазделяемаяПаттерны программирования на CUDAСтратегии распределения работыРазное
- 24. Работа с текстурной памятьюБыстрая, кешируемая в 2-х
- 25. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиКонстантнаяТекстурнаяГлобальнаяРазделяемаяПаттерны программирования на CUDAСтратегии распределения работыРазное
- 26. Работа с глобальной памятьюМедленная, некешируемая (G80), чтение/записьЗапись
- 27. Работа с глобальной памятьюCoalescing, Compute Capability 1.0,
- 28. Работа с глобальной памятьюCoalescing, Compute Capability 1.0, 1.1CoalescingNo coalescing
- 29. Работа с глобальной памятьюCoalescing, Compute Capability 1.2,
- 30. Работа с глобальной памятьюCoalescing, Compute Capability 1.2, 1.3
- 31. Работа с глобальной памятьюCoalescing. РекомендацииИспользуйте cudaMallocPitch для
- 32. Коалесинг
- 33. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиКонстантнаяТекстурнаяГлобальнаяРазделяемаяПаттерны программирования на CUDAСтратегии распределения работыРазное
- 34. Работа с разделяемой памятьюБыстрая, некешируемая, чтение/записьОбъявление при
- 35. Память разделена на 16 банков памяти, по
- 36. Работа с разделяемой памятьюБанки памятиПрямой доступСмешанный доступ 1:1Доступ без конфликтов банков
- 37. Работа с разделяемой памятьюБанки памяти2-кратный конфликт8-кратный конфликтДоступ с конфликтами банков
- 39. СодержаниеПроцесс разработки программ CUDAРабота с различными типами
- 40. Паттерны программирования на CUDAПриоритеты оптимизацииОбъединение запросов к
- 41. СодержаниеПроцесс разработки программ CUDAРабота с различными типами
- 42. Паттерны программирования на CUDAСценарий работы с shared
- 43. СодержаниеПроцесс разработки программ CUDAРабота с различными типами
- 44. Паттерны программирования на CUDAКопирование global shared: 32-bitdim3
- 45. Паттерны программирования на CUDAКопирование global shared: 8-bitdim3
- 46. СодержаниеПроцесс разработки программ CUDAРабота с различными типами
- 47. Паттерны программирования на CUDAОбработка в shared памятиНезависимая
- 48. Паттерны программирования на CUDAОбработка в shared памяти__device__
- 49. Паттерны программирования на CUDAОбработка в shared памяти
- 50. Паттерны программирования на CUDAОбработка в shared памяти
- 51. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное
- 52. Стратегии распределения работыЗадачи с нерегулярным параллелизмомПеременное кол-во итераций Большое кол-во ветвлений
- 53. Стратегии распределения работы: C & CРазделить ядра
- 54. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное
- 55. Стратегии распределения работы: Uber-kernelUber-kernelif ( A ){ Exec_A();}Else if ( B ){ Exec_B();}…
- 56. Стратегии распределения работы: Uber-kernel (2)timeBlocks 0Blocks 1Blocks 2Blocks 3kernel1kernel2
- 57. Стратегии распределения работы: Uber-kernel (3)timeBlocks 0Blocks 1Blocks 2Blocks 3kernel1kernel2
- 58. Стратегии распределения работы: Uber-kernel (3)timeBlocks 0Blocks 1Blocks
- 59. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыCommand & ConquerUber-kernelPersistent threadsРазное
- 60. Стратегии распределения работы
- 61. Стратегии распределения работы: Persistent threadstimeWarp 0Warp 1Warp 2Warp 3Block 0Block 1
- 62. Стратегии распределения работы: Persistent threadstimeWarp 0Warp 1Warp 2Warp 3Block 0Block 1
- 63. Стратегии распределения работы: Persistent threads (2)timeWarp 0Warp 1Warp 2Warp 3Block 0Block 1
- 64. Стратегии распределения работы: Persistent threads (3) timeWarp 0Warp 1Warp 2Warp 3Block 0
- 65. СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыРазное
- 66. ВетвлениеЕсли происходит ветвление внутри варпа, то разные
- 67. Ветвление
- 68. Инструкции
- 69. ОптимизацияPTX
- 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 работают быстрее, чем *Возможно
- 76. Конфигурация gridDim и blockDim возможно во время
- 77. Исользование templatetemplate __global__ void kernel(){ int
- 78. РазноеМатематика FPU (на GPU в частности) не
- 79. Ресурсы нашего курса CUDA.CS.MSU.SUМесто для вопросов и дискуссийМесто
- 80. Вопросы
- 81. Скачать презентацию
- 82. Похожие презентации
СодержаниеПроцесс разработки программ CUDAРабота с различными типами памятиПаттерны программирования на CUDAСтратегии распределения работыРазное
![Паттерны программирования на CUDAКопирование global shared: 32-bitdim3 block(64);__shared__ float dst[64];__global__ void kernel(float](/img/tmb/12/1108332/edc81fd4eb3c7168b707b00d32a1cd29-720x.jpg "Вопросы программирования и оптимизации приложений на CUDA")
![Паттерны программирования на CUDAКопирование global shared: 8-bitdim3 block(64);__shared__ byte dst[64];__global__ void kernel_bad(byte](/img/tmb/12/1108332/3177a0774a890394155d915b688b30af-720x.jpg "Вопросы программирования и оптимизации приложений на CUDA")
![Паттерны программирования на CUDAОбработка в shared памяти (2)Одно из решений:__shared__ int buf[16][17];dim3](/img/tmb/12/1108332/1c755a7380d2a4794d3d2a4fb1680642-720x.jpg "Вопросы программирования и оптимизации приложений на CUDA")
![ИнструкцииСледить за ветвлениемЗаменить часть вычислений на look-up таблицуИнтринсики__sinf(); __cosf(); expf()__[u]mul24()__fdividef()__[u]sad()](/img/tmb/12/1108332/c34768b6574f1d763c74b73184829dee-720x.jpg "Вопросы программирования и оптимизации приложений на CUDA")
Слайд 2
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Паттерны
программирования на CUDA
Слайд 3
Содержание
Процесс разработки программ CUDA
Портирование части приложения
Общие рекомендации по
оптимизации
Инструментарий
Работа с различными типами памяти
Паттерны программирования на CUDA
Стратегии
распределения работыРазное
Слайд 4
Процесс разработки программ CUDA
Портирование части приложения
Определение класса портируемой
задачи
Уровень параллелизма. SIMD
Классы задач, которые в общем случае невозможно
распараллелить
Слайд 6
Содержание
Процесс разработки программ CUDA
Портирование части приложения
Общие рекомендации по
оптимизации
Инструментарий
Работа с различными типами памяти
Паттерны программирования на CUDA
Стратегии
распределения работыРазное
Слайд 7
Переосмысление задачи в терминах параллельной обработки данных
Выявляйте параллелизм
Максимизируйте
интенсивность вычислений
Иногда выгоднее пересчитать чем сохранить
Избегайте лишних транзакций по
памятиОсобое внимание особенностям работы с различными видами памяти (об этом дальше)
Эффективное использование вычислительной мощи
Разбивайте вычисления с целью поддержания сбалансированной загрузки SM’ов
Параллелизм потоков vs. параллелизм по данным
Процесс разработки программ CUDA
Общие рекомендации по оптимизации
Слайд 8
Процесс разработки программ CUDA
Общие рекомендации по оптимизации
Occupancy
Покрытие латентностей:
инструкции потока выполняются последовательно
Исполнение других потоков необходимо для
покрытия латентностейЗанятость: отношение активных варпов к максимально возможному
В архитектуре Tesla 32 варпа на SM
Слайд 9
Процесс разработки программ CUDA
Общие рекомендации по оптимизации
Occupancy
Увеличение занятости
приводит к лучшему покрытию латентностей
После определенной точки (~50%), происходит
насыщениеЗанятость ограничена достыпными ресурсами:
Регистры
Разделяемая память
Слайд 10
Содержание
Процесс разработки программ CUDA
Портирование части приложения
Общие рекомендации по
оптимизации
Инструментарий
Работа с различными типами памяти
Паттерны программирования на CUDA
Стратегии
распределения работыРазное
Слайд 12
Процесс разработки программ CUDA
Инструментарий: Компилятор
Статическая компиляция: IDE(MS Visual
Studio + cuda.rules), Makefile, CL
PTX JIT-компиляция
Слайд 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
Достоинства эмуляции
Исполняемый файл, скомпилированный в режиме эмуляции работает
целиком на CPU
Не требуется драйвер CUDA и GPU
Каждый поток
GPU эмулируется потоком CPUПри работе в режиме эмуляции можно:
Использовать средства отладки CPU (точки останова и т.д.)
Обращаться к любым данным GPU с CPU и наоборот
Делать любые CPU-вызовы из код GPU и наоборот (например printf())
Выявлять ситуации зависания, возникающие из-за неправильного применения __syncthreads()
Процесс разработки программ CUDA
Инструментарий: Отладчик
Слайд 16
Недостатки эмуляции
Часто работает очень медленно
Неумышленное разыменование указателей GPU
на стороне CPU или наоборот
Результаты операций с плавающей точкой
CPU и «настоящего» GPU почти всегда различаются из-за:Разного порядка выполняемых операций
Разных допустимых ошибок результатов
Использования большей точности при расчёте промежуточных результатов на CPU
Процесс разработки программ CUDA
Инструментарий: Отладчик
Слайд 19
Процесс разработки программ CUDA
Инструментарий: Профилировщик
CUDA Profiler, позволяет отслеживать:
Время
исполнения на CPU и GPU в микросекундах
Конфигурацию grid и
thread blockКоличество статической разделяемой памяти на блок
Количество регистров на блок
Коэффициент занятости GPU (Occupancy)
Количество объединенных и индивидуальных запросов к глобальной памяти (coalescing)
Количество дивергентных путей исполнения (branching)
Количество выполненных инструкций
Количество запущенных блоков
Вся эта информация собирается с первого SM или TPC. Профилирование Uber-kernel’ов с осторожностью
Слайд 21
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Константная
Текстурная
Глобальная
Разделяемая
Паттерны
программирования на CUDA
Стратегии распределения работы
Разное
Слайд 22
Работа с константной памятью
Быстрая, кешируемая, только для чтения
Данные
должны быть записаны до вызова кернела (например при помощи
cudaMemcpyToSymbol)Всего 64Kb (Tesla)
Объявление при помощи слова __constant__
Доступ из device кода простой адресацией
Срабатывает за 4 такта на один адрес внутри варпа
4 такта на всю транзакцию в случае если все потоки внутри варпа читают один адрес
В худшем случае 64 такта
Слайд 23
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Константная
Текстурная
Глобальная
Разделяемая
Паттерны
программирования на CUDA
Стратегии распределения работы
Разное
Слайд 24
Работа с текстурной памятью
Быстрая, кешируемая в 2-х измерениях,
только для чтения
Данные должны быть записаны при помощи cudaMemcpyToArray,
либо возможно прикрепление к глобальной памяти через cudaBindTexture2DОбъявление при помощи текстурных ссылок
Доступ из device кода при помощи tex1D, tex2D, tex1Dfetch
Лучшая производительность при условии что потоки одного варпа обращаются к локализованной окрестности в 2D
Слайд 25
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Константная
Текстурная
Глобальная
Разделяемая
Паттерны
программирования на CUDA
Стратегии распределения работы
Разное
Слайд 26
Работа с глобальной памятью
Медленная, некешируемая (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
потоков. Типы транзакций:
4-байтовые слова, одна 64-байтовая транзакция
8-байтовые слова, одна
128-байтовая транзакция16-байтовые слова, две 128-байтовых транзакции
Все 16 элементов должны лежать в едином сегменте, размер и выравнивание которого совпадает с размером транзакции
Строгий порядок доступа: k-й поток обращается к k-му элементу в сегменте
При нарушении порядка вместо одной транзакции получается 16
Некоторые из потоков могут не участвовать
Слайд 29
Работа с глобальной памятью
Coalescing, Compute Capability 1.2, 1.3
Объединенная
транзакция получается, если все элементы лежат в сегментах:
размера 32
байта, потоки обращаются к 1-байтовым элементамразмера 64 байта, потоки обращаются к 2-байтовым элементам
размера 128 байт, потоки обращаются к 4- и 8-байтовым элементам
Нестрогий порядок доступа. Возможно обращение несколькими потоками к одному адресу
При выходе за границы сегмента число транзакций увеличивается минимально
Слайд 31
Работа с глобальной памятью
Coalescing. Рекомендации
Используйте cudaMallocPitch для работы
с 2D-массивами
Конфигурируйте блоки с большей протяженностью по x
Параметризуйте конфигурацию,
экспериментируйтеВ сложных случаях используйте привязку сегмента глобальной памяти к текстуре в случае если Compute Capability < 1.2
cudaBindTexture, tex1Dfetch
cudaBindTexture2D, tex2D
Слайд 33
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Константная
Текстурная
Глобальная
Разделяемая
Паттерны
программирования на CUDA
Стратегии распределения работы
Разное
Слайд 34
Работа с разделяемой памятью
Быстрая, некешируемая, чтение/запись
Объявление при помощи
слова __shared__
Доступ из device кода при помощи индексирования
Самый быстрый
тип памяти после регистров, низкая латентность доступаМожно рассматривать как полностью открытый L1-кеш
При работе с разделяемой памятью следует помнить о ее разбиении на банками памяти
Слайд 35 Память разделена на 16 банков памяти, по числу
потоков в варпе
Каждый банк может обратиться к одному адресу
за 1 тактМаксимальное число адресов, к которым может обратиться память одновременно совпадает с числом банков
Одновременное обращение нескольких потоков из одного полуварпа к одному банку приводит к конфликту банков и сериализации запросов (кроме broadcast)
Работа с разделяемой памятью
Банки памяти
Слайд 36
Работа с разделяемой памятью
Банки памяти
Прямой доступ
Смешанный доступ 1:1
Доступ
без конфликтов банков
Слайд 37
Работа с разделяемой памятью
Банки памяти
2-кратный конфликт
8-кратный конфликт
Доступ с
конфликтами банков
Слайд 39
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Паттерны
программирования на CUDA
Приоритеты оптимизации
Сценарий работы с shared памятью
Копирование global
<-> sharedОбработка в shared памяти
Стратегии распределения работы
Разное
Слайд 40
Паттерны программирования на CUDA
Приоритеты оптимизации
Объединение запросов к глобальной
памяти
Ускорение до 20 раз
Стремление к локальности
Использование разделяемой памяти
Высокая скорость
работыУдобство взаимодействия потоков
Эффективное использование параллелизма
GPU не должен простаивать
Преобладание вычислений над операциями с памятью
Много блоков и потоков в блоке
Банк-конфликты
Если избавление от 4-кратных конфликтов банков влечет увеличение числа инструкций, то данный вид оптимизации можно не делать
Слайд 41
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Паттерны
программирования на CUDA
Приоритеты оптимизации
Сценарий работы с shared памятью
Копирование global
<-> sharedОбработка в shared памяти
Стратегии распределения работы
Разное
Слайд 42
Паттерны программирования на CUDA
Сценарий работы с shared памятью
Загрузка
данных из глобальной памяти в разделяемой
__syncthreads();
Обработка данных в разделяемой
памяти__syncthreads(); //если требуется
Сохранение результатов в глобальной памяти
Шаги 2–4 могут быть обрамлены в условия и циклы
Шаг 4 может быть ненужен в случае если выходные данные независимы между собой
Слайд 43
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Паттерны
программирования на CUDA
Приоритеты оптимизации
Сценарий работы с shared памятью
Копирование global
<-> sharedОбработка в shared памяти
Стратегии распределения работы
Разное
Слайд 44
Паттерны программирования на 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];}
Слайд 45
Паттерны программирования на CUDA
Копирование global shared: 8-bit
dim3
block(64);
__shared__ byte dst[64];
__global__ void kernel_bad(byte *data)
{//no coalescing, 4-way bank
conflicts presentdst[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
Приоритеты оптимизации
Сценарий работы с shared памятью
Копирование global
<-> sharedОбработка в shared памяти
Стратегии распределения работы
Разное
Слайд 47
Паттерны программирования на CUDA
Обработка в shared памяти
Независимая обработка
элементов. Прямой доступ будет вызывать 4-кратный конфликт банков.
Задача: переформировать
потоки в 4 группы по 16 индексов так, чтобы при новой косвенной адресации не было конфликтов банков.__shared__ byte buf[64];
dim3 block(64);
Слайд 48
Паттерны программирования на CUDA
Обработка в 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
Слайд 49
Паттерны программирования на CUDA
Обработка в shared памяти (2)
Независимая
обработка элементов. Прямой доступ будет вызывать 16-кратный конфликт банков.
Задача:
свести число банк-конфликтов до нуля.__shared__ int buf[16][16];
dim3 block(16,16);
Слайд 50
Паттерны программирования на CUDA
Обработка в shared памяти (2)
Одно
из решений:
__shared__ int buf[16][17];
dim3 block(16,16);
Bank Indices without Padding
Bank Indices
with Padding
Слайд 51
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Паттерны
программирования на CUDA
Стратегии распределения работы
Command & Conquer
Uber-kernel
Persistent threads
Разное
Слайд 52
Стратегии распределения работы
Задачи с нерегулярным параллелизмом
Переменное кол-во итераций
Большое кол-во ветвлений
Слайд 53
Стратегии распределения работы: C & C
Разделить ядра на
более простые
Позволяет выявить bottleneck
Увеличивает Occupancy
Возможность перераспределять работу между ядрами
Слайд 54
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Паттерны
программирования на CUDA
Стратегии распределения работы
Command & Conquer
Uber-kernel
Persistent threads
Разное
Слайд 55
Стратегии распределения работы: Uber-kernel
Uber-kernel
if ( A )
{
Exec_A();
}
Else if ( B )
{
Exec_B();
}
…
Слайд 56
Стратегии распределения работы: Uber-kernel (2)
time
Blocks 0
Blocks 1
Blocks 2
Blocks
3
kernel1
kernel2
Слайд 57
Стратегии распределения работы: Uber-kernel (3)
time
Blocks 0
Blocks 1
Blocks 2
Blocks
3
kernel1
kernel2
Слайд 58
Стратегии распределения работы: Uber-kernel (3)
time
Blocks 0
Blocks 1
Blocks 2
Blocks
3
if (A) kernel1
if (B) kernel2
Blocks 0
Blocks 1
Blocks 2
Blocks 3
Слайд 59
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Паттерны
программирования на CUDA
Стратегии распределения работы
Command & Conquer
Uber-kernel
Persistent threads
Разное
Слайд 61
Стратегии распределения работы: Persistent threads
time
Warp 0
Warp 1
Warp 2
Warp
3
Block 0
Block 1
Слайд 62
Стратегии распределения работы: Persistent threads
time
Warp 0
Warp 1
Warp 2
Warp
3
Block 0
Block 1
Слайд 63
Стратегии распределения работы: Persistent threads (2)
time
Warp 0
Warp 1
Warp
2
Warp 3
Block 0
Block 1
Слайд 65
Содержание
Процесс разработки программ CUDA
Работа с различными типами памяти
Паттерны
программирования на CUDA
Стратегии распределения работы
Разное
Слайд 66
Ветвление
Если происходит ветвление внутри варпа, то разные ветви
исполнения сериализуются
Увеличивается общее количество инструкций
Если ветвление происходит между варпами,
то штраф минимальный
Слайд 74
Инструкции
Следить за ветвлением
Заменить часть вычислений на look-up таблицу
Интринсики
__sinf();
__cosf(); expf()
__[u]mul24()
__fdividef()
__[u]sad()
Слайд 75
__mul24 и __umul24 работают быстрее, чем *
Возможно увеличение
числа регистров после применения
На будущих архитектурах ситуация может развернуться
наоборот и __mul24 станет медленнееИспользование флагов
В остальном целочисленная арифметика работает примерно с такой же скоростью, как и с плавающей точкой (за исключением целочисленного деления)
Разное
Слайд 76
Конфигурация gridDim и blockDim возможно во время исполнения:
void
callKernel(dim3 grid, dim3 threads)
{
kernel();
}
Размеры CTA и
GRID
Слайд 77
Исользование template
template
__global__
void kernel()
{
int x = threadIdx.x + blockIdx.x
* tx;}
void callKernel(dim3 grid)
{
kernel<16, 16, 1><<
}
Шаблоны
Слайд 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
