Лекция 3. Общая память (Лекции)

Лихогруд Николайn.lihogrud@gmail.comЧасть третьяРазделяемая(общая) памятьDevice Расположена в том же устройстве, что икеш L1 Совместно используется (разделяется)всеми нитями виртуального блока Если на мультипроцессоре работаетнесколько блоков – общая памятьделится между ними поровну У каждого блока своё адресноепространство общей памятиSMSMSMCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreShared&L1 cacheShared&L1 cache Конфигурации: 16КB общая память, 48KB L1L2 cache 48КB общая память, 16KB L1 – поумолчаниюDevice MemoryShared&L1 cacheРазделяемая(общая) памятьDeviceSMВозможные обменымежду устройствамипри обработке обращений вглобальную памятьВозможные обменымежду устройствамипри обработке обращенийв общую памятьSMSMCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreShared&L1 cacheShared&L1 cacheL2 cacheDevice MemoryShared&L1 cacheВыделение общей памяти В GPU коде объявляем статический массив или переменнуюс атрибутом __shared__#define SIZE 1024__global__ void kernel() {__shared__ int array[SIZE]; //массив__shared__ float varSharedMem; //переменная…}Особенности использованияПеременные с атрибутом __shared__ с точки зренияпрограммирования: Существуют только на время жизни блока недоступны с хоста или из других блоков Индивидуальны для каждого блока и привязаны к его личномупространству общей памяти каждый блок нитей видит «своё» значение Не могут быть проинициализированы при объявленииРаздача указателя нитям блока__global__ void kernel() {__shared__ int *memoryOnDevice;if (threadIdx.x == 0) {// выделяет память только первая нитьsize_t size = blockDim.x * sizeof(float);memoryOnDevice = (int *)malloc(size);memset(memoryOnDevice, 0, size);}memoryOnDevice[thredIdx.x] = …;…// использование указателя всеми нитями блока}Раздача указателя нитям блока__global__ void kernel() {__shared__ int *memoryOnDevice;if (threadIdx.x == 0) {// выделяет память только первая нитьsize_t size = blockDim.x * sizeof(float);memoryOnDevice = (int *)malloc(size);memset(memoryOnDevice, 0, size);}Нужна синхронизация!??memoryOnDevice[thredIdx.x] = …;…// использование указателя всеми нитями блока}Синхронизация Рассмотрим пример ядра, запускаемого на одномерном линейномгриде:__global__ void kernel() {__shared__ int shmem[BLOCK_SIZE];shmem[threadIdx.x] = __sinf(threadIdx.x);int a = shmem[(threadIdx.x + 1 )% BLOCK_SIZE];…} Каждая нить Записывает __sinf от своего индекса в соответствующую ей ячейкумассива Читает из массива элемент, записанный соседней нитьюСинхронизация Рассмотрим пример ядра, запускаемого на одномерном линейномгриде:__global__ void kernel() {__shared__ int shmem[BLOCK_SIZE];shmem[threadIdx.x] = __sinf(threadIdx.x);int a = shmem[(threadIdx.x + 1 )% BLOCK_SIZE];…} Варпы выполняются в непредсказуемом порядке Может получиться, что нить ещё не записала элемент, соседняя ужепытается его считать! read-after-write, write-after-read, write-after-write конфликтыСинхронизация Явная синхронизация нитей одного блока void __syncthreads();При вызове этой функции нить блокируется домомента, когда: все нити в блоке достигнут данную точку результаты всех инициированных к данному моментуопераций с глобальной\общей памятью, станут виднывсем нитям блокаСинхронизация __syncthreads() можно вызывать в ветвях условногооператора только если результат его условия одинаков вовсех нитях блока, иначе выполнение может зависнуть илистать непредсказуемымСинхронизация__global__ void kernel() {__shared__ int shmem[BLOCK_SIZE];shmem[threadIdx.x] = __sinf(threadIdx.x);__syncthreads();int a = shmem[(threadIdx.x + 1 )% BLOCK_SIZE];…} Каждая нить Записывает __sinf от своего индекса в соответствующую ей ячейкумассива Ожидает завершения операций в других нитях Читает из массива элемент, записанный соседней нитьюРаздача указателя нитям блока__global__ void kernel() {__shared__ int *memoryOnDevice;if (threadIdx.x == 0) {// выделяет память только первая нитьsize_t size = blockDim.x * 64;memoryOnDevice = (int *)malloc(size);memset(memoryOnDevice, 0, size);}Нужна синхронизация!__syncthreads();…// использование указателя всеми нитями блока}Динамическая общая память Бывают ситуации, когда нужный размер общей памяти неизвестен на этапе компиляции Зависит от размер задачи, блока и т.д. В этом случае выделить память как статическуюпеременную невозможно Можно указать требуемый размер общей памяти призапуске ядраДинамическая общая память В GPU коде объявляем указатель для доступа к общейпамяти:__global__ void kernel() {extern __shared__ int array[];…} В третьем параметре конфигурации запуска указываемсколько общей памяти нужно выделить каждому блокуkernel<<<gridDim, blockDim, SIZE >>>(params)Динамическая общая память Все переменные extern __shared__ type var[]указывают на одно и то же начало динамической общейпамяти, выделенной блоку Ядру может быть одновременно выделена статическая, идинамическая память. Если суммарный объем динамической и статической памятипревышает 48кб на блок – произойдет ошибка запускаСтратегия использования Общая память по смыслу является кешем, управляемымпользователем Имеет низкую латентность - расположена на том жеоборудовании, что и кеш L1, скорость загрузки сопоставима срегистрами Приложение явно выделяет и использует общую память Пользовать сам выбирает что, как и когда в ней хранить Шаблон доступа может быть произвольным, в отличие отL1Стратегия использования Типичная стратегия использования: Нити блока коллективно1.Загружают данные из глобальной памяти в общую Каждая нить делает часть этой загрузки2.

Синхронизуются Чтобы никакая нить не начинала чтение данных, загружаемыхдругой нитью, до завершения их загрузки3. Используют загруженные данные для вычисления результаты Если нити что-то пишут в общую память, то также можетпотребоваться синхронизация4. Записывают результаты обратно в глобальную памятьРедукция Блоку нитей сопоставляем часть массива Каждый блок нитей суммирует элементы из своей части массива Блок нитей Копирует данные в общую память Иерахически суммирует данные в общей памяти Сохраняет результат в глобальной памятиИерархическое суммированиеИерархическое суммированиеЯдро суммирования__global__ void reduce (int *inData, int *outData){__shared__ int data [BLOCK_SIZE];int tid = threadIdx.x;int i = blockIdx.x * blockDim.x + threadIdx.xdata [tid] = inData [i];__syncthreads ();for ( int s = 1; s < blockDim.x; s *= 2) {if (tid % (2*s) == 0) {data [tid] += data [tid + s];}__syncthreads ();}if (tid == 0) {outData [blockIdx.x] = data [0];}}Банки общей памяти Для увеличения полосы пропускания устройство, на которомрасположена общая память, разделено на подмодули («банки») n – число банков m – сколько последовательных байтов может отдать каждый банк зацикл Адресное пространство общей памяти разделено на nнепересекающихся подмножеств, расположенных в разныхбанках Банки работают независимо друг-от-друга и могут вместевыдать максимум n*m байтов за один циклБанки общей памяти на Fermi 32 банка, каждый банк может выдать за 2 такта ядер одно 32-битное слово (4 последовательных байта) Последовательные 32-битные слова располагаются впоследовательных банках Номер банка для слова по адресу addr: (addr / 4) % 32 За два такта ядер общая память может отдать 128 байтБанк 14132260388516Банк 0 Банк 1 Банк 2 Банк 3…04812…128132136140…256260264268………015Банк 0Банк 11044128172256300384428512556………………31Обращения в общую память Обращение выполняется одновременно всеми нитями варпа (SIMT) Банки работаю параллельно Если варпу нитей нужно получить 32 4-байтных слова,расположенных в разных банках, то такой запрос будет выполненодновременно всеми банками Каждый банк выдаст соответствующее слово Пропускная способность = 32 х пропускная способность банка Поддерживается рассылка (broadcast): Если часть нитей (или все) обращаются к одному и тому же 4-хбайтному слову, то нужное слово будет считано из банка и розданосоответствующим нитям (broadcast) без накладных расходовБанк конфликты Если хотя бы два нужных варпу слова расположены в одном банке, тотакая ситуация называется «банк конфликтом» и обращение вглобальную память будет «сериализованно»: Такое обращение аппаратно разбивается на серию обращений, несодержащих банк конфликтов Если число обращений, на которое разбит исходный запрос,равно n, то такая ситуация называется банк-конфликтомпорядка n Пропускная способность при этом падает в n разБанки общей памяти на Kepler 32 банка, каждый банк может выдать за 1 такт ядер 8 байтов На Kepler частота ядер в 2 раза меньше, чем на Fermi Два режима разбиения общей памяти на банки: Последовательные 32-битные слова располагаются впоследовательных банкаx: (addr / 4) % 32 Последовательные 64-битные слова располагаются впоследовательных банках: (addr / 8) % 32 За два такта ядер общая память может отдать 256 байтБанк 14Банк 01320260388516Банк 1128 4Банк 2…256 384 260 388 264 392 268 396…512 640 516 644 520 648 524 652……0…15Банк 0136 12…138…132 8Банк 3………31Банк 11044128172256300384428512556……Последовательные 32битные слова впоследовательныхбанкаxБанк 1812264368520Банк 0Банк 1Банк 2Банк 3…081624…4122028256 260 264 268 272 276 280 284…512 516 520 524 528 532 536 540………0…15Банк 0………31Банк 11088492256344360348512600……Последовательные 64битные слова впоследовательныхбанкаxПоследовательные 32-битные слова в последовательных банкаxБанк 00Могут быть отданыбанком 0 за одинтакт*Банк 1128 4Банк 2…256 384 260 388 264 392 268 396…512 640 516 644 520 648 524 652……136 12…138…132 8Банк 3………Последовательные 64-битные слова в последовательных банкаxМогут быть отданыбанком 0 за одинтакт*Банк 0Банк 1Банк 2Банк 3…081624…4122028256 260 264 268 272 276 280 284…512 516 520 524 528 532 536 540…………*Каждая ячейка соответствует 4-м последовательным байтам……Банк конфликты на Kepler Последовательные 32-битные слова располагаются впоследовательных банкаx:Банк-конфликта между двумя нитями нет, если запрашиваются байты 32битныx слов из разных банков, либо запрашиваемые слова находятся в 32битных словах с адресами i и i + 128, 256*n <= i < 256*n + 128(broadcast) Последовательные 64-битные слова располагаются впоследовательных банкаx:Банк-конфликта между двумя нитями нет, если запрашиваются байты 64битных слов из разных банков, либо байты 64-битного слова из одного банка(broadcast)Установка режима общей памятиcudaError_t cudaDeviceSetSharedMemConfig (cudaSharedMemConfig config ) Глобально для всех запусков ядер cudaSharedMemBankSizeDefault - последовательные 32-битныеслова в последовательных банкаx cudaSharedMemBankSizeFourByte - последовательные 32-битныеслова в последовательных банкаx cudaSharedMemBankSizeEightByte - последовательные 64-битныеслова в последовательных банкаxcudaError_t cudaFuncSetSharedMemConfig ( const void* func,cudaSharedMemConfig config ) Для запусков конкретного ядраЗачем устанавливать режимextern __shared__ double arr[];double res = sin(arr[thredIdx.x * 3]); Нить 0 обращается к байту со смещением 0, нить 16 – 384• cudaSharedMemBankSizeFourByte•Оба обращения попадают в один банк, банк-конфликт второго порядка• cudaSharedMemBankSizeEightByte•Обращения попадают в банки 0 и 16, банк-конфликта нетБанк конфликтывторого порядкаПримеры банк-конфликтовextern __shared__ float char[];float data = shared[BaseIndex + s * threadIx.x];// конфликтызависят от s Нити threadIx.x и (threadIx.x + n) обращаются к элементам из одного итого же банка когда s*n делится на 32 (число банков). S=1:shared[BaseIndex + threadIx.x] // нет конфликта S=2:shared[BaseIndex + 2*threadIx.x] // конфликт 2-го порядкаНапример, между нитями threadIx.x=0 и (threadIx.x = 16) –попадают в один варп!Распространенная проблема Пусть в общей памяти выделена плоская плотная матрицашириной, кратной 32, и соседние нити варпа обращаются ксоседним элементам столбца__shared__ int matrix[32][32]matrix[thredIdx.x][4] = 0;Распространенная проблема Пусть в общей памяти выделена плоская плотная матрицашириной, кратной 32, и соседние нити варпа обращаются ксоседним элементам столбца__shared__ int matrix[32][32]matrix[thredIdx.x][4] = 0;Банк конфликт 32-го порядкаРаспространенная проблема__shared__ int matrix[32][32]matrix[thredIdx.x][4] = 0;Банк конфликт 32-го порядкаРешение: набивка__shared__ int matrix[32][32 + 1]matrix[thredIdx.x][4] = 0; //нет конфликтаРаспространенная проблемаПусть банков 10, матрица 10х10Транспонирование матрицы__global__ void simpleTranspose(ElemType *inputMatrix, ElemType*outputMatrix, int width, int height) {int i = threadIdx.y + blockIdx.y * blockDim.y;int j = threadIdx.x + blockIdx.x * blockDim.x;if ( i < height && j < width) {outputMatrix[j * (height) + i] = inputMatrix[ i * (width) + j];}}Нити варпа записывают элементы столбца32 транзации на одну записьТранспонирование через общую память Считать плитку матрицы в общую память Записать в результат транспонированную плиткуНити варпа читают строкуплитки из исходной матрицыНити варпа читают столбец плитки вобщей памяти и пишут строкув транспонированной матрицеТранспонирование через общую память__global__ void shmemTranspose( ElemType *inputMatrix,ElemType *outputMatrix, int width, int height) {int i = threadIdx.y + blockIdx.y * blockDim.y;int j = threadIdx.x + blockIdx.x * blockDim.x;__shared__ ElemType shmem[32][32];if ( i < height && j < width) {shmem[threadIdx.y][threadIdx.x] = inputMatrix[i * (width) + j];}__syncthreads();Банк-конфликт 32-го порядкаif ( i < width && j < height) {outputMatrix[i * (height) + j] = shmem[threadIdx.x][threadIdx.y];}}Транспонирование через общую память__global__ void correctShmemTranspose( ElemType *inputMatrix,ElemType *outputMatrix, int width, int height) {int i = threadIdx.y + blockIdx.y * blockDim.y;int j = threadIdx.x + blockIdx.x * blockDim.x;__shared__ ElemType shmem[32][32 + 1];Избавились от банк-конфилктаif ( i < height && j < width) {shmem[threadIdx.y][threadIdx.x] = inputMatrix[i * (width) + j];}__syncthreads();if ( i < width && j < height) {outputMatrix[i * (height) + j] = shmem[threadIdx.x][threadIdx.y];}}Тесты• Kepler K20c, матрица 16384x16384 элементаЯдроDoubleFloatПростое53.942ms41.040msС общей памятью52.338ms32.840msС общей памятью безбанк-конфликтов36.057ms21.274msВыводы Общую память можно использовать как управляемый кеш дляреиспользования данных Как в редукции Доступ в общую память может быть произвольным, в отличие от кеша L1 Можно применять пространственные преобразования к данным, используя общуюпамять как буфер (транспонирование - поворот и отражение) Банк-конфликты высокого порядка могут сильно ухудшить пропускнуюспособность общей памяти Доступный объем общей памяти ограничен Влияет на occupancyThe end.