Лекция 4. Константная память. Регистры и локальная память (Лекции)

Лихогруд Николайn.lihogrud@gmail.comЧасть четвертаяКонстантная памятьDevice Параметр устройстваtotalConstMem Кешируется в специальномread-only кеше – Unifrom Cache Объём 8 KBCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreShared&L1 cacheShared&L1 cacheL2 cacheDevice MemoryUniform cache Объём до 64KBSMUniform cache Расположена в DRAM GPUSMКонстантная памятьDeviceВозможные обменымежду устройствамипри обработке обращенийв константную памятьCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreShared&L1 cacheShared&L1 cacheL2 cacheDevice MemoryUniform cacheВозможные обменымежду устройствамипри обработке обращенийв общую памятьSMUniform cacheВозможные обменымежду устройствамипри обработке обращений вглобальную памятьSMОбъявление В глобальной области видимости__constant__ int constMem[1024];__constant__ int constVar; Можно ещё дополнительно указать __device__ , чтобыподчеркнуть, что память выделяется на устройстве :__device__ __constant__ int constVar2;Особенность __device__ и __constant__ Переменные с атрибутами __device__ и __constant__находятся в глобальной области видимости и хранятся вобъектном модуле как отдельные символы Память под них выделяется на устройстве автоматически пристарте приложения, освобождается при завершении Работать с ними на хосте можно через функцииcudaMemcpyToSymbol() , cudaMemcpyToSymbolAsync(),cudaGetSymbolAddress(), cudaMemcpyFromSymbol(),cudaMemcpyFromSymbolAsynс(), cudaGetSymbolSize()Особенность __constant__ Доступна на чтение (и только на чтение!) из любой нити любогогрида обычным способом:__constant__ int constMem[32];__global__ void kernel() {…int a = constMem[ threadIdx.x / 32 ];…}Пример__constant__ float constData[256]; На хосте:float data[256];cudaMemcpyToSymbol(constData, data, sizeof(data));cudaMemcpyFromSymbol(data, constData, sizeof(data));Обращения в константную память Обращение выполняется одновременно для всех нитей варпа(SIMT) Исходное обращение разбивается на столько запросов,сколько различных адресов в нём было Каждый запрос выполняется либо через запрос к кешу вслучае кеш-попадания, либо через глобальную память Если их было n, то пропускная способность уменьшается в nразОбращения в константную память__constant__ int constMem[32];__global__ void kernel() {…int a = constMem[ threadIdx.x / 32 ]; // 1 запрос вконстантнуюпамятьint a = constMem[ threadIdx.x]; // 32 запроса вконстантнуюпамять…}Однородные обращения Помимо обработки запросов в константную память, Uniform Cacheобрабатывает «Однородные» обращения (Unifrom Accesses) – когда всенити варпа обращаются в глобальную память по одному адресу При выполнении требований:Доступ только по чтениюАдрес не зависит от индекса нити в блоке (threadIdx)while(k < 100 ) tmp += a[blockIdx.x + k++];Компилятор может заменить в ассемблере обычную инструкцию загрузкииз глобальной памяти на инструкцию однородной загрузки, котораябудет выполнена через Uniform CacheОднородные обращения При выполнении требований:Доступ только по чтениюАдрес не зависит от индекса нити в блоке (threaIdx)while(k < 100 ) tmp += a[blockIdx.x + k++]; Второе требование гарантирует, что все нити варпа обращаются поодному адресу Чтобы помочь компилятору с первым требованием, можно пометитьуказатели атрибутом const __restrictПередача параметров в ядра Параметры передаются в ядра через константную память Параметры передаются в единственном экземпляре для всехнитей грида Это приемлимо, т.к., в основном, нити варпа обращаются к одному и тому жепараметру -> Uniform Access После первого варпа параметры уже будут в кеше Суммарный размер передаваемых параметров долен быть небольше, чем 4 KBПередача грида в ядра Помимо параметров, через константную память передаютсяразмеры грида: gridDim, blockDim threadIdx, blockIdx нить получает из спец.

регистров(заведомо не Uniform) gridDim, blockDim нить считывает из константной памятив самом начале работы (Uniform)Пример: фильтрация с ядром Наложить ядро на окрестность каждого элемента Для элемента (i,j) записать в (i, j) матрицы-результата суммупроизведений элементов окрестности на соответствующиекоэффициенты фильтраКвадратное ядроПример: фильтрация с ядромtmp = 0;for(ik=0..2*r)for(jk=0..2*r)tmp += matrix[i+ik-r][j+jk-r]*kernel[ik][jk];result[i][j]=tmp;Квадратное ядроПример: фильтрация с ядромtmp = 0;for(ik=0..2*r)for(jk=0..2*r)tmp += matrix[i+ik-r][j+jk-r] * kernel[ik][jk];result[i][j]=tmp;Не зависит от i,jВсе нити варпа читают один элементМожно расположить в константнойпамятиПример: преобразование координат Поворот, перенос, сжатие/растяжение объекта в пространствевыполняются путем умножения всех вершин объекта на матрицупреобразования координат0010203001 02 03′11 12 13′21 22 23 × = ′ ,31 31 331 ′ - однородные координаты вершины 1- нормированные однородныекоординаты вершины ′ ′ ′ ′- координаты вершины послепреобразованияПример: преобразование координат При реализации на CUDA каждая нить умножает координатыодной вершины на матрицу преобразования координат − номер нити − массив вершин − матрица преобразования координат = × [] Матрица преобразования координат для всех нитей одна,доступ к ней только на чтение, нити варпа обращаются к одномуэлементу -> Можем разместить её в константной памятиПример: преобразование координат Вспомогательные типы для упрощения ядра:typedef float CoordinateType;typedef CoordinateType TransformMatrixRow[4]; // строка матрицы преобразованияtypedef TransformMatrixRow TransformMatrix[4]; // матрица преобразования как 4строки// вершина фигурыstruct Vertex {CoordinateType x,y,z,w; // однородные координаты__device__ Vertex(CoordinateType x = 0, CoordinateType y = 0,CoordinateType z = 0, CoordinateType w = 1):x(x), y(y), z(z), w(w) {} // конструктор__device__ void normalize() {…} // нормализация – разделить x, y, z на w// умножение на строку матрицы преобразования__device__ CoordinateType operator*(TransformMatrixRow row) {return x * row[0] + y * row[1] + z * row[2] + w * row[3];}};Пример: преобразование координат Вспомогательные типы для упрощения ядра:// вершины фигурыstruct Vertexes {CoordinateType *x, *y, *z; // структура с массивами, а не массив структурint count; // число вершин// получение отдельной вершины__device__ Vertex operator[](int n) {return Vertex(x[n], y[n], z[n]);}// запись вершины__device__ void setVertexAtIndex(Vertex vertex, int index) {x[index] = vertex.x;y[index] = vertex.y;z[index] = vertex.z;}};Пример: преобразование координат Матрица преобразования и ядро:__constant__ TransformMatrix transformMatrix; // матрица преобразования вконстантной памяти__global__ void coordinateTransform(Vertexes vertexes) {int id = blockIdx.x * blockDim.x + threadIdx.x; // индекс нитиif (id >= vertexes.count) { // проверка на выход за границуreturn;}Vertex vertex = vertexes[id]; // считали вершинуVertex newVertex( vertex * transformMatrix[0],vertex * transformMatrix[1],vertex * transformMatrix[2],vertex * transformMatrix[3]); // умножение на матрицуnewVertex.normalize(); // нормализацияvertexes.setVertexAtIndex(newVertex, id); // запись результата}Пример: преобразование координат Результаты тестов:Tesla K20c, 65 миллионов вершинТочностьАтрибут матрицыFloatDouble__constant__12.392ms22.161ms__device__50.546ms53.667msВыводы Использование константной памяти: Позволяет разгрузить кеш L1 Ускоряет выполнение при выполнении требованияоднородности обращенийРегистровая памятьDeviceгенерируются компилятором Каждая нить является эксклюзивнымпользователем своих регистров на всё времяжизни Регистры делятся между всемирезидентными нитями намультипроцессоре Быстрое переключение контекстаCoreCoreCoreCoreCoreCoreCoreShared&L1 cacheCoreCoreCoreCoreCoreCoreCoreCoreShared&L1 cacheL2 cacheDevice MemoryRegisters Инструкции с использованием регистровCoreUniform cache Недоступна явно в программеSMRegisters Самая быстрая памятьSMUniform cache Расположена на мультипроцессореРегистры в Fermi и Kepler Параметр устройства regsPerMultiprocessor - число 32-битных регистровна каждом мультипроцессоре.

Разделяются межу виртуальными блоками,работающими на мультипроцессоре. Для Fermi: Каждый мультипроцессор содержит 32768 32-битных регистров (128 КB) Отдельной нити доступно максимум 63 регистра Максимальное число регистров при occupancy = 1: 32768 / 1536 = 21 Для Kepler: Каждый мультипроцессор содержит 65536 32-битных регистров (256КB) Отдельной нити доступно максимум 255 регистров Максимальное число регистров при occupancy = 1: 65536 / 2048 = 32Особенность выделения регистров Регистры выделяются пачками отдельно для каждого варпа Пусть размер пачки - regAllocUnitSize Пусть каждой нити нужно nRegs регистров.

Тогда общее числорегистров, необходимых блоку из blockSize нитей: ∗ 32∗ ∗32Число пачек на варп− ближайшее сверху целоеЧисло варпов в блокеОсобенность выделения регистров ∗ 32∗ ∗32 Для Fermi = 64 Пусть nRegs = 21 а блок – 512 нитей(16 варпов):21∗3264= 11 пачек, 11 ∗ 64 ∗ 16 = 11264 - столько регистроввыделится одному блоку 11264 ∗ 3 = 33792 > 32768, т.е. на три блока регистров не хватит Получаем, что максимальное число регистров при occupancy=1 иблоке из 512 нитей равно 20 Для Kepler = 256Локальная памятьDeviceSMCoreCoreCoreCoreCoreCore Высокая латентность Доступ осуществляется по тем же правилам,что и запросы в глобальную память Кеширование в L1 Транзакции Недоступна явно в программе Инструкции обращения в локальнуюпамять генерируются компилятором Обладает упрощённой схемой адресации Оптимизирована для минимизацииколичества транзакцийShared&L1 cacheCoreCoreCoreCoreCoreCoreCoreCoreShared&L1 cacheL2 cacheDevice MemoryRegistersCoreUniform cacheCoreRegistersSMUniform cache Расположена в DRAMКогда используется локальная память? Обычно компилятор помещает на регистры все локальныепеременные Но есть исключения, размещаемые в локальной памяти: Массивы, для которые не всегда можно определить к какомуэлементу в какой момент времени идёт доступ (неконстантные индексы) Большие массивы или структуры, которые использовали быслишком много регистров Любая переменная, если превышен лимит регистров на нить(так называемый «спиллинг регистров» register spilling)Когда используется локальная память? Некоторые встроенные математические функции могутиспользовать локальную память Через локальную память передаётся часть операндов привызове функций В т.ч.