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

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

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

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

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

5 2011 Портирование части приложения

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

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

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

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

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

11 2011 Инструментарий: Компилятор

12 2011 Инструментарий: Компилятор Статическая компиляция: IDE(MS Visual Studio + cuda.rules), Makefile, CL PTX JIT-компиляция

13 2011 Инструментарий: Компилятор

14 2011 Инструментарий: Отладчик GPU debugger GPU Profiler

15 Visual Profiler

16 Profiler Counter Plot

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

18 Occupancy Calculator Spreadsheet

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

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

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

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

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

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

25 2011 Объединение запросов к глобальной памяти. GPU умеет объединять ряд запросов к глобальной памяти в транзакцию одного сегмента Длина сегмента должна быть 32/64/128 байт Сегмент должен быть выровнен по своему размеру

26 2011 Объединение (coalescing) 1.2/1.3 Нити обращаются к – 8-битовым словам, дающим один 32- байтовый сегмент – 16-битовым словам, дающим один 64- байтовый сегмент – 32-битовым словам, дающим один 128- байтовый сегмент Объединение происходит на уровне полу-варпов

27 2011 Объединение (coalescing) 1.2/1.3 Если хотя бы одно условие не выполнено – объединяет их в набор сегментов – для каждого проводится отдельная транзакция 1 транзакция 64B сегмент 2 транзакции 64B и 32B сегменты 1 транзакция 128B сегмент Легенда: -нить -128B сегмент Легенда: -нить -128B сегмент

28 2011 Объединение (coalescing) 2.x На мультипроцессоре есть L1 кэш – Физически там, где разделяемая память Мультипроцессоры имеют общий L2 кэш Флаги компиляции – Использовать L1 и L2 :-Xptxas -dlcm=ca – Использовать L2 :-Xptxas -dlcm=cg Кэш линия 128B Объединение происходит на уровне варпов

29 2011 Объединение (coalescing) 2.x Если L1 кэш включен: всегда 128B сегменты Если L1 кэш выключен: всегда 32B сегменты 2 транзакция по 128B 4 транзакция по 32B Легенда: -нить -128B сегмент - 32B сегменты Легенда: -нить -128B сегмент - 32B сегменты

30 2011 Объединение (coalescing) Увеличения скорости работы с памятью на порядок Лучше использовать не массив структур, а набор массивов отдельных компонент – Проще гарантировать условия выполнения coalescinga

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

32 2011 Коалесинг

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

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

35 2011 Банки памяти Память разделена на 16 банков памяти, по числу потоков в варпе Каждый банк может обратиться к одному адресу за 1 такт Максимальное число адресов, к которым может обратиться память одновременно совпадает с числом банков Одновременное обращение нескольких потоков из одного полуварпа к одному банку приводит к конфликту банков и сериализации запросов (кроме broadcast) Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0

36 2011 Банки памяти Доступ без конфликтов банков Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 Thread 15 Thread 7 Thread 6 Thread 5 Thread 4 Thread 3 Thread 2 Thread 1 Thread 0 Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 Thread 15 Thread 7 Thread 6 Thread 5 Thread 4 Thread 3 Thread 2 Thread 1 Thread 0 Прямой доступСмешанный доступ 1:1

37 2011 Банки памяти Доступ с конфликтами банков 2-кратный конфликт8-кратный конфликт Thread 11 Thread 10 Thread 9 Thread 8 Thread 4 Thread 3 Thread 2 Thread 1 Thread 0 Bank 15 Bank 7 Bank 6 Bank 5 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 Thread 15 Thread 7 Thread 6 Thread 5 Thread 4 Thread 3 Thread 2 Thread 1 Thread 0 Bank 9 Bank 8 Bank 15 Bank 7 Bank 2 Bank 1 Bank 0 x8

38 2011 Эффективная работа с shared- памятью Tesla 20 Вся shared-память разбита на 32 банка Все нити варпа обращаются в память совместно. Каждый банк работает независимо от других Можно одновременно выполнить до 32 обращений к shared-памяти

39 2011 Эффективная работа с shared- памятью Tesla 20 Банк конфликты: – При обращении >1 нити варпа к разным 32битным словам из одного банка При обращени >1 нити варпа к разным байтам одного 32битного слова, конфликта нет – При чтении: операция broadcast – При записи: результат не определен

40 2011

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

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

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

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

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

46 2011 Копирование 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]; }

47 2011 Копирование global shared: 8-bit dim3 block(64); __shared__ byte dst[64]; __global__ void kernel_bad(byte *data) {//no coalescing, 4-way bank conflicts present dst[threadIdx.x] = data[threadIdx.x]; } __global__ void kernel_good(byte *data) {//coalescing, no bank conflicts, no branching if (threadIdx.x < 16) { int tx = threadIdx.x * 4; *((int *)(dst + tx)) = *((int *)(data + tx)); }

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

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

50 2011 Обработка в shared памяти __device__ int permute64by4(int t) { return (t >> 4) + ((t & 0xF)

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

69 2011 Ветвление

70 2011 Инструкции

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

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

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

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

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

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

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

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

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

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

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

82 2011 Вопросы

83 2011 Спасибо! Александр Гужва Антон Обухов Владимир Фролов Дмитрий Ватолин Дмитрий Микушин Евгений Перепелкин Михаил Смирнов Николай Сахарных