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

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

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 3 DirectX Compute Microsoft API Тесно интегрирован с Direct3D Доступен –CS 4.x: DirectX 10 HW –CS 5.x: DirectX 11 HW ID3D11Device ID3D11Resource ID3D11View ID3D11DeviceContext

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

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

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 6 DirectX ID3D11Device –ID3D11Resource –ID3D11View ID3D11DeviceContext ID3D11Asynchronous –ID3D11Query ID3D11ComputeShader ID3DX11Effect

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 7 DirectX ID3D11Device –ID3D11Resource Buffer StructuredBuffer Texture –ID3D11View ShaderResourceView UnorderedAccessView RenderTargetView

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 8 DirectX ID3D11DeviceContext –Dispatch(bx, by, bz) –DispatchIndirect(pBuffer, offset) –End(pQuery) –GetData(g_pQuerry, NULL, 0, 0 )

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 9 DirectX ID3D11ComputeShader ConstantBuffer ShaderResourceView UnorderedAccessView ID3D11Effect ConstantBuffer ShaderResourceView UnorderedAccessView

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 10 DirectX ID3D11ComputeShader pContext->CSSetShader(pCS, NULL, 0); pContext->CSSetUnorderedAccessViews(0, 1, &pRWBufUAV, NULL); ID3D11Effect pEffect->GetVariableByName(tSimple)->AsUnorderedAccessView()- >SetUnorderedAccessView(pRWBufUAV); pEffect->GetTechniqueByName(tSimple)->GetPassByName(pSimple")- >Apply(0, pContext);

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

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

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

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 14 Типы памяти в CUDA Тип памяти ДоступУровень выделения Скорость работы РегистрыR/WPer-thread Высокая(on-chip) ЛокальнаяR/WPer-thread Низкая (DRAM) SharedR/WPer-block Высокая(on-chip) ГлобальнаяR/WPer-grid Очень плохая (DRAM) ConstantR/OPer-grid Высокая(L1 cache) TextureR/OPer-grid Высокая(L1 cache)

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

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

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

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

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

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

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

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

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 23 Работа с глобальной памятью в CUDA float * devPtr; // pointer device memory // allocate device memory cudaMalloc ( (void **) &devPtr, 256*sizeof ( float ); // copy data from host to device memory cudaMemcpy ( devPtr, hostPtr, 256*sizeof ( float ), cudaMemcpyHostToDevice ); // process data // copy results from device to host cudaMemcpy ( hostPtr, devPtr, 256*sizeof( float ), cudaMemcpyDeviceToHost ); // free device memory cudaFree ( devPtr ); Пример кода для работы с памятью

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 24 Работа с глобальной памятью в CUDA cudaError_t cudaMalloc ( void ** devPtr, size_t size ); cudaError_t cudaMallocPitch ( void ** devPtr, size_t * pitch, size_t width, size_t height ); cudaError_t cudaFree ( void * devPtr ); cudaError_t cudaMemcpy ( void * dst, const void * src, size_t count, enum cudaMemcpyKind kind ); cudaError_t cudaMemcpyAsync ( void * dst, const void * src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream ); cudaError_t cudaMemset ( void * devPtr, int value, size_t count ); Функции для работы с глобальной памятью

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

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

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

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

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

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 30 Использование константной памяти Константная память используется тогда, когда в ядро необходимо передать много различных данных, которые будут одинаково использоваться всеми тредами ядра. __constant__ float contsData [256]; - объявление глобальной переменной с именем contsData для использования в качестве константной памяти. cudaMemcpyToSymbol ( constData, hostData, sizeof ( data ), 0, cudaMemcpyHostToDevice ); --копирование данных с центрального процессора в константную память. Использование внутри ядра ничем не отличается от использования любой глобальной переменной на хосте.

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

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

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

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