МОСКОВСКИЙ АВИАЦИОННЫЙ ИНСТИТУТ национальный исследовательский университет Факультет прикладной

МОСКОВСКИЙ АВИАЦИОННЫЙ ИНСТИТУТ (национальный исследовательский университет) Факультет прикладной математики и физики Кафедра вычислительной математики и программирования Лекция 13 «Типы памяти» Выполнил: Семенов С. А. Руководитель: Ревизников Д. Л.

Содержание ØМодель памяти CUDA ØГлобальная память Global memory ØЛокальная память Local memory ØРазделяемая память Shared memory ØПамять констант Constant memory ØТекстурная память Texture memory ØРегистровая память Register memory Московский авиационный институт (национальный исследовательский университет) 2 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Direct. X Compute Ø Microsoft API Ø Тесно интегрирован с Direct 3 D Ø Доступен – CS 4. x: Direct. X 10 HW – CS 5. x: Direct. X 11 HW Ø ID 3 D 11 Device üID 3 D 11 Resource üID 3 D 11 View Ø ID 3 D 11 Device. Context Московский авиационный институт (национальный исследовательский университет) 3 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

CUDA vs Direct. X Спецификаторы функций Ø CUDA C Ø Direct. X – __global__ – __host__ – __device__ Московский авиационный институт (национальный исследовательский университет) – Compute Shader – n/a 4 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

CUDA vs Direct. X Compute Пространство памяти Ø CUDA C – – Ø Direct. X – – __device__ __shared__ __constant__ local Московский авиационный институт (национальный исследовательский университет) 5 [Structured]Buffer groupshared Constant Buffer n/a Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Direct. X • ID 3 D 11 Device – ID 3 D 11 Resource – ID 3 D 11 View • ID 3 D 11 Device. Context • ID 3 D 11 Asynchronous – ID 3 D 11 Query • ID 3 D 11 Compute. Shader • ID 3 DX 11 Effect Московский авиационный институт (национальный исследовательский университет) 6 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Direct. X • ID 3 D 11 Device – ID 3 D 11 Resource • Buffer • Structured. Buffer • Texture – ID 3 D 11 View • Shader. Resource. View • Unordered. Access. View • Render. Target. View Московский авиационный институт (национальный исследовательский университет) 7 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Direct. X • ID 3 D 11 Device. Context – Dispatch(bx, by, bz) – Dispatch. Indirect(p. Buffer, offset) – End(p. Query) – Get. Data(g_p. Querry, NULL, 0, 0 ) Московский авиационный институт (национальный исследовательский университет) 8 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Direct. X ID 3 D 11 Compute. Shader ID 3 D 11 Effect • Constant. Buffer • Shader. Resource. View • Unordered. Access. View Московский авиационный институт (национальный исследовательский университет) 9 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Direct. X ID 3 D 11 Compute. Shader p. Context->CSSet. Shader(p. CS, NULL, 0); p. Context->CSSet. Unordered. Access. Views(0, 1, &p. RWBuf. UAV, NULL); ID 3 D 11 Effect p. Effect->Get. Variable. By. Name(“t. Simple”)->As. Unordered. Access. View()>Set. Unordered. Access. View(p. RWBuf. UAV); p. Effect->Get. Technique. By. Name(“t. Simple”)->Get. Pass. By. Name(“p. Simple")>Apply(0, p. Context); Московский авиационный институт (национальный исследовательский университет) 10 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Московский авиационный институт (национальный исследовательский университет) 11 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Direct. X Compute Московский авиационный институт (национальный исследовательский университет) 12 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Direct. X Compute Московский авиационный институт (национальный исследовательский университет) 13 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Типы памяти в CUDA Тип памяти Доступ Уровень выделения Per-thread Скорость работы Регистры Локальная R/W Shared R/W Глобальная R/W Per-block Per-grid Constant R/O Per-grid Высокая(on-chip) Очень плохая (DRAM) Высокая(L 1 cache) Texture R/O Per-grid Высокая(L 1 cache) Московский авиационный институт (национальный исследовательский университет) 14 Высокая(on-chip) Низкая (DRAM) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Регистры Ø 32 Кб на SM, используется для хранения локальныхпеременных. Ø Расположены непосредственно на чипе, скорость доступа самая быстрая. Ø Выделяются отдельно для каждого треда. Московский авиационный институт (национальный исследовательский университет) 15 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Локальная Ø используется для хранения локальных переменных когда регистров не хватает, Ø скорость доступа низкая, так как расположена в DRAM партициях. Ø Выделяется отдельно для каждого треда. Московский авиационный институт (национальный исследовательский университет) 16 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Разделяемая • 16 Кб (или 48 Кб на Fermi) на SM, • используется для хранения массивов данных, используемых совместно всеми тредами в блоке. • Расположена на чипе, имеет чуть меньшую скорость доступа чем регистры (около 10 тактов). • Выделяется на блок. Московский авиационный институт (национальный исследовательский университет) 17 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Глобальная Ø основная память видеокарты (на данный момент максимально 6 Гб на Tesla c 2070). Ø Используется для хранения больших массивов данных. Ø Расположена на DRAM партициях и имеет медленную скорость доступа (около 80 тактов). Ø Выделяется целиком на грид. Московский авиационный институт (национальный исследовательский университет) 18 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Константная • память, располагающаяся в DRAM партиции, • Кэшируется специальным константным кэшем. • Используется для передачи параметров в ядро, превышающих допустимые размеры для параметров ядра. • Выделяется целиком на грид. Московский авиационный институт (национальный исследовательский университет) 19 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Текстурная • память, располагающаяся в DRAM партиции, кэшируется. • Используется для хранения больших массивов данных, выделяется целиком на грид. Московский авиационный институт (национальный исследовательский университет) 20 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Типы памяти в CUDA Ø Самая быстрая – shared (on-chip) и регистры Ø Самая медленная – глобальная (DRAM) Ø Для ряда случаев можно использовать кэшируемую константную и текстурную память Ø Доступ к памяти в CUDA идет независимо для каждой половины warp’а (half-warp) Московский авиационный институт (национальный исследовательский университет) 21 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Работа с памятью в CUDA Основа оптимизации – оптимизация работы с памятью: Ø Максимальное использование sharedпамяти Ø Использование специальных паттернов доступа к памяти, гарантирующих эффектный доступ üПаттерны работают независимо в пределах каждого half-warp’а Московский авиационный институт (национальный исследовательский университет) 22 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Работа с глобальной памятью в CUDA Пример кода для работы с памятью float * dev. Ptr; // pointer device memory // allocate device memory cuda. Malloc ( (void **) &dev. Ptr, 256*sizeof ( float ); // copy data from host to device memory cuda. Memcpy ( dev. Ptr, host. Ptr, 256*sizeof ( float ), cuda. Memcpy. Host. To. Device ); // process data // copy results from device to host cuda. Memcpy ( host. Ptr, dev. Ptr, 256*sizeof( float ), cuda. Memcpy. Device. To. Host ); // free device memory cuda. Free ( dev. Ptr ); Московский авиационный институт (национальный исследовательский университет) 23 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Работа с глобальной памятью в CUDA Функции для работы с глобальной памятью cuda. Error_t cuda. Malloc ( void ** dev. Ptr, size_t size ); cuda. Error_t cuda. Malloc. Pitch ( void ** dev. Ptr, size_t * pitch, size_t width, size_t height ); cuda. Error_t cuda. Free ( void * dev. Ptr ); cuda. Error_t cuda. Memcpy ( void * dst, const void * src, size_t count, enum cuda. Memcpy. Kind kind ); cuda. Error_t cuda. Memcpy. Async ( void * dst, const void * src, size_t count, enum cuda. Memcpy. Kind kind, cuda. Stream_t stream ); cuda. Error_t cuda. Memset ( void * dev. Ptr, int value, size_t count ); Московский авиационный институт (национальный исследовательский университет) 24 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Оптимизация работы с глобальной памятью. Ø Обращения идут через 32/64/128 битовые слова Ø При обращении к t[i] üsizeof(t [0]) равен 4/8/16 байтам üt [i] выровнен по sizeof ( t [0] ) Ø Вся выделяемая память всегда выровнена по 256 байт Московский авиационный институт (национальный исследовательский университет) 25 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Условия возникновения объединения СС 1. 0, 1. 1 СС >= 1. 2 Нити обращаются к · 32 -битовым словам, давая 64 -байтовый блок · 64 -битовым словам, давая 128 -байтовый блок Все 16 слов лежат в пределах блока k-ая нить half-warp’а обращается к k-му слову блока Нити обращаются к · 8 -битовым словам, дающим один 32 -байтовый сегмент · 16 -битовым словам, дающим один 64 -байтовый сегмент · 32 -битовым словам, дающим один 128 -байтовый сегмент Получающийся сегмент выровнен по своему размеру Московский авиационный институт (национальный исследовательский университет) 26 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Работа с shared-памятью Ø Самая быстрая (on-chip) Ø Сейчас всего 16 Кбайт на один мультипроцессор Ø Совместно используется всеми нитями блока Ø Отдельное обращение для каждой половины warp’а (half-warp) Ø Как правило, требует явной синхронизации Московский авиационный институт (национальный исследовательский университет) 27 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Эффективная работа с shared-памятью z Для повышения пропускной способности вся shared-память разбита на 16 банков z Каждый банк работает независимо от других z Можно одновременно выполнить до 16 обращений к shared-памяти z Если идет несколько обращений к одному банку, то они выполняются по очереди Московский авиационный институт (национальный исследовательский университет) 28 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Расположение short в разделяемой памяти Московский авиационный институт (национальный исследовательский университет) 29 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Использование константной памяти • Константная память используется тогда, когда в ядро необходимо передать много различных данных, которые будут одинаково использоваться всеми тредами ядра. • __constant__ float conts. Data [256]; - объявление глобальной переменной с именем conts. Data для использования в качестве константной памяти. • cuda. Memcpy. To. Symbol ( const. Data, host. Data, sizeof ( data ), 0, cuda. Memcpy. Host. To. Device ); --копирование данных с центрального процессора в константную память. • Использование внутри ядра ничем не отличается от использования любой глобальной переменной на хосте. Московский авиационный институт (национальный исследовательский университет) 30 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Использование текстурной памяти Ø Объем данных не влезает в shared память Ø Паттерн доступа хаотичный Ø Данные переиспользуются разными потоками Московский авиационный институт (национальный исследовательский университет) 31 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Типы текстурной памяти в CUDA Линейная cuda. Array Можно использовать обычную глобальную память Ограничения: · Только для одномерных массивов · Нет фильтрации · Доступ по целочисленным координатам · Обращение по адресу вне допустимого диапазона возвращает ноль Доступ: tex 1 Dfetch(tex, int) Позволяет организовывать данные в 1 D/ 2 D/3 D массивы данных вида: · 1/2/4 компонентные векторы · 8/16/32 bit signed/unsigned integers · 32 bit float · 16 bit float (driver API) Доступ по семейству функций tex 1 D()/tex 2 D()/tex 3 D() Московский авиационный институт (национальный исследовательский университет) 32 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Общая схема работы с текстурной памятью Московский авиационный институт (национальный исследовательский университет) 33 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования

Московский авиационный институт (национальный исследовательский университет) 34 Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования