Лекция. CUDA 2 (Колганов) (Электронные лекции), страница 3

32/512=6% лишнейпамяти будет выделено (Но для больших матриц эта доля будетсущественно меньше) Зато каждая строка будет выровнена по 128 байт Обращение matrix[idy*128+ idx]Pitch in bytes….512Строка 01024Строка 1Адрес начала каждой строки выровнен по128! Обращения варпов в памятьвыполняются за одну транзакцию1536Строка 2Padding (набивка)2048Строка 3Матрицы и глобальная памятьОбращения варпа071 транзакция!1315набивкаВыделение памяти с «паддингом» cudaError_t cudaMallocPitch (void ** devPtr, size_t * pitch,size_t width, size_t height) width – логическая ширина матрицы в байтах Выделяет не менее width * height байтов , может добавить в конец строк набивку, сцелью соблюдения выравнивания начала строк сохраняет указатель на память в (*devPtr) сохраняет фактическую ширину строк в байтах в (*pitch) Адрес элемента (Row, Column) матрицы, выделенной при помощи cudaMallocPitch:T* pElement = (T*)((char*) devPtr + Row * pitch) + Column;Копирование в матрицу с padding-омcudaError_t cudaMemcpy2D ( void* dst, size_t dpitch, constvoid* src, size_t spitch, size_t width,size_t height, cudaMemcpyKind kind ) dst - указатель на матрицу, в которую нужно копировать ,dpitch – фактическая ширина её строк в байтахsrc - указатель на матрицу из которой нужно копировать,spitch – фактическая ширина её строк в байтахwidth – сколько байтов каждой строки нужно копироватьheight – число строкkind – направление копирования (как в обычном cudaMemcpy) Данная функция из начала каждой строки исходной матрицы копируется по widthбайтов.

Всего копируется width*height байтов, при этом Адрес строки с индексом Row определяется по фактической ширине:(char*)src + Row* spitch – в матрице-источнике(char*)dst + Row* dpitch – в матрице-получателеОбращение к матрице по столбцам? Матрица расположена по строкам, а обращение идёт по столбцамОбращения нитей варпа512512 + j10241024+jКаждая нить варпа обращается всвою строку к элементу в столбце j15361536+j20482048+jЕсли матрица имеет размер больше 128 байт, тоэти обращения ни за что не «влезут» в однутранзакцию!Транспонировать! Решение – хранить матрицу в транспонированном виде! В этом случае обращения по столбцам превратятся в обращения кпоследовательным адресам Выделять память под транспонированную матрицу также черезcudaMallocPitchМассивы структур?struct example {int a;int b;int c;}__global__ void kernel(example * arrayOfExamples) {int idx = threadIdx.x + blockIdx.x * blockDim.x;arrayOfExamples[idx].c =arrayOfExamples[idx].b + arrayOfExamples[idx].a;}Обращения нитей варпа….a b c a b c a b c a b c a b cОбращение варпа в памятьбудет выполняться в тритранзакции - 256 лишнихбайтов… a b c a b c a b cСтруктура с массивами!struct example {int *a;int *b;int *c;}__global__ void kernel(example arrayOfExamples) {int idx = threadIdx.x + blockIdx.x * blockDim.x;arrayOfExamples.c[idx] =arrayOfExamples.b[idx] + arrayOfExamples.a[idx];}Обращения нитей варпаa a a a a a a … b b b b b b bОбращение варпа в памятьбудет выполняться за однутранзакцию… c c c c c c cПроблема Косвенной адресации Использование косвенной адресации нежелательно,поскольку требует двух чтений из памяти (сначала A[i], потомA[i][j])float **A;A[i][j] = 1;Проблема Косвенной адресации A[i] для разных нитей варпа скорее всего будут в одной кеш-линии -> одно обращение к кешу Но A[i,j] в общем случае могут быть «разбросаны»float **A;A[i][j] = 1;Проблема Косвенной адресации Принципиального решения нет! Скорее всего придется переработать алгоритмfloat **A;A[i][j] = 1;Выводы Обращения нитей варпа в память должны быть пространственно-локальными Начала строк матрицы должны быть выровнены Массивы структур -> структура с массивами 16КB vs 48KB L1 Избегаем косвенной адресации Избегаем обращений нитей варпа к столбцу матрицы В случае сильно разреженного доступа проверяем работу сотключенным кешем.