Слайд 2
ОСОБЕННОСТИ CUDA APIs
Виды CUDA APIs и возможности CUDA-устройств
Слайд 3
Виды CUDA APIs
CUDA Driver API
Ручная инициализация контекста GPU
Отсутствуют
CUDA-расширения для C++
Код CPU может компилироваться без nvcc
CUDA Runtime
API
Автоматическая инициализация контекста GPU
Наличие CUDA-расширений для C++
Слайд 4
Выбор CUDA API
CUDA Driver API
Больше гибкости ( +
)
Сложность кода ( – )
CUDA Runtime API
Меньше гибкости (
– )
Простота кода ( + )
Слайд 5
Совместимость CUDA API
Имеется обратная совместимость версий
Слайд 6
Вычислительные возможности GPU
Capability – это версия архитектуры CUDA
GPU, которая указывает на его вычислительные возможности и особенности
cudaGetDeviceProperties()
Слайд 7
ОЦЕНКА ПРОИЗВОДИТЕЛЬНОСТИ
Способы оценки эффективности приложений CUDA
Слайд 8
Время выполнения
Общее время вычислений на GPU
Время выполнения участка
кода GPU
Слайд 9
Таймеры CPU
Таймеры CPU позволяют замерять общее время выполнения
вычислений на GPU
Слайд 10
Таймеры CUDA
Таймеры CUDA позволяют замерять время выполнения участка
кода GPU
Слайд 11
Скорость передачи данных
Теоретическая пропускная способность
FDDRAM * (RDDRAM/8) *
sizeof(float),
где FDDRAM – частота, RDDRAM – разрядность шины
Эффективная пропускная
способность
(BR+ BW) / time,
где BR и BW – объем прочитанной/записанной информации
Реальная пропускная способность
Слайд 12
ОПТИМИЗАЦИЯ РАБОТЫ С ПАМЯТЬЮ
Способы оптимизации работы с памятью
CUDA GPU
Слайд 14
Типы памяти устройства
Streaming Multiprocessor
Регистровая память
Разделяемая память
Память констант
Texture Processing
Cluster
Память текстур
DDRAM
Локальная память
Глобальная память
Слайд 15
Передача данных Host/Device
Является дорогостоящей операцией
Возможна асинхронная передача
cudaMemcpy()
cudaMemcpyAsync()
Слайд 16
Асинхронная передача данных
Копирование данных и выполнение ядра можно
осуществлять параллельно
Слайд 17
Возможная оптимизация
Синхронная передача данных в GPU
Асинхронная передача данных
в GPU
Слайд 18
Нулевое копирование (Zero Copy)
Прямое обращение к памяти Host’а
Встроенные
видеокарты
Использование кэша CPU
Слайд 19
Объединенное чтение DDRAM
Выравнивание исходных данных по границе слова
Потоки
warp’а должны осуществлять одновременное чтение DDRAM
Слайд 20
Разделяемая память и конфликты
Общая для всех потоков блока
Распределяется
между блоками
Разбивается на банки (32-битные слова)
Слайд 21
Регистровое давление
Регистры жестко распределяются между потоками мультипроцессора
При большом
количестве потоков возникает конфликт доступа к регистрам
Слайд 22
ВЫБОР ОПТИМАЛЬНОЙ ТОПОЛОГИИ
Методы оценки топологии вычислений CUDA
Слайд 23
Степень покрытия
Степень покрытия мультипроцессора – это отношение числа
активных warp'ов к максимально возможному числу активных warp'ов
По количеству
используемых регистров
С учетом топологии вычислений
Без учета топологии вычислений
По размеру используемой разделяемой памяти
Слайд 24
Определение степени покрытия
CUDA GPU:
8192 регистра
768 потоков на мультипроцессор
Топология:
12
регистров на ядро
128 потоков в блоке
Tmax = 8192 регистров
/ 12 регистров = 682 потока
Treal = int(682 / 128) * 128 = 640 потоков
С = Тreal / Tmax = 83%
Слайд 25
ОПТИМИЗАЦИЯ КОДА
Оптимизация инструкций CUDA
Слайд 26
Регистровая зависимость
Инструкция использует регистр, значение которого было получено
при выполнении предыдущей инструкции
register = instruction1();
instruction2(register);
Слайд 27
Float vs Double
Арифметические операции с float-числами осуществляются быстрей,
чем с double-числами
Рекомендуется использовать суффикс «f» при объявлении числовых
констант, например, 3.14f
Слайд 28
Деление чисел
При делении чисел на степень двойки рекомендуется
использовать оператор сдвига
X / N → X >> log2(N)
X
% N → X & (N-1)
Слайд 29
Степень числа
Для известных целых значений степеней рекомендуется использовать
явное умножение вместо вызова pow()
pow(X, 2) → X *
X
pow(X, 3) → X * X * X
Слайд 30
Часто используемые функции
Обратный квадратный корень
rsqrtf() / rsqrt()
Прочие арифметические
операции
expf2() / exp2() – экспонента во 2-й степени
expf10() / exp10() –
экспонента в 10-й степени
cbrtf() / cart() – экспонента в степени 1/3
rcbrtf() / rebut() – экспонента в степени -1/3
Слайд 31
Точность vs Скорость
Аппаратные аналоги функций
__sinf() / sinf()
__cosf() /
cosf()
__expf() / expf()
Совмещенные функции
sincosf() / sincos()
Слайд 32
УПРАВЛЕНИЕ ПОТОКОМ КОМАНД
Общие рекомендации по написанию кода
Слайд 33
Операторы ветвления
Инструкции управления потоком команд (if, switch, for,
while, do-while) отрицательно сказываются на производительности
В идеале все потоки
warp'а должны идти по одному пути, иначе увеличивается количество выполняемых инструкций и возможно последовательное выполнение
Слайд 35
ОТЛАДКА И ПРОФИЛИРОВАНИЕ
Отладка и профилирование приложений CUDA
Слайд 36
Существующие утилиты
Linux
CUDA-GDB
http://developer.nvidia.com/cuda-gdb
Windows Vista & Windows 7
NVIDIA Parallel
Nsight
http://developer.nvidia.com/nvidia-parallel-nsight
Слайд 37
АППАРАТНЫЕ ОСОБЕННОСТИ GPU
Краткий обзор архитектурных особенностей GPU
Слайд 38
Причины рассогласования
Основные причины рассогласования результатов вычислений на GPU
и CPU
Усечение double чисел до float при отсутствии аппаратной
поддержки double
Неассоциативность арифметических операций с дробными числами
Небольшие отклонения от стандарта IEEE 754
Особенности архитектуры процессоров x86
Слайд 39
Литература
NVIDIA Developer Zone
http://developer.nvidia.com/cuda
NVIDIA Parallel Nsight
http://developer.nvidia.com/cuda-gdb
CUDA C Best Practices
Guide
http://developer.download.nvidia.com/compute/cuda/4_0/toolkit/docs/CUDA_C_Best_Practices_Guide.pdf