Лекция 2. Программирования GPU CPU. Компиляция. Иерархия_ доступ к памяти (1265183), страница 2

Просмтор этого файла доступен только зарегистрированным пользователям. Но у нас супер быстрая регистрация: достаточно только электронной почты!

Текст из файла (страница 2)

Т.к. конструкцию запуска ядра в *.cpp применятьнельзя, вынесем её в отдельную функцию, определённую в каком-нибудь*.cu#include <cuda_runtime.h> // здесь объявления функций тулкитаvoid launchKernel(params); // определить эту функцию в каком-нибудь *.cuint main() {… // обычный хост-кодcudaSetDevice(0); // Можно! Обычная функция, потом слинкуется… // обычный хост-код,kernel<<1,1>>>(params) здесь нельзя, только в *.cu!launchKernel(params); // Внутри этой функции будет вызвано ядро// Определена в некотором *.cu… // обычный хост-код}Компиляция хост-кодаВ файле test.cpp : Основной хост-код.

Т.к. конструкцию запуска ядра в *.cpp применятьнельзя, вынесем её в отдельную функцию, определённую в каком-нибудь*.cu Компиляция:$g++ -I /toolkit_install_dir/include test.cpp –c –o test.oУказали путь, по которому следует искать инклюд-файлы для CUDAПопросили сделать объектникИли же через nvcc, без указания директории с инклюдами:$nvcc test.cpp –c –o test.oКомпиляция device-кодаВ файле kernel.cu Определяем ядро и функцию для его запуска.

В функции запуска рассчитываетсяконфигурация и дополнительно выводится на экран время работы ядра__global __ void kernel(params)….; // код ядра}void launchKernel(params){{…;// расчёт конфигурации запуска в зависимости от параметров,создание событийfloat time;cudaEventRecord( start, 0 );kernel<<< конфигурация >>> (params); // запуск ядраcudaEventRecord( stop, 0 );cudaEventSynchronize( stop );cudaEventElapsedTime( &time, start, stop );// Время между двумя событиямиprintf(%4.4f, time);}Компиляция device-кодаВ файле kernel.cu Определяем ядро и функцию для его запуска. В функции запускарассчитывается конфигурация и дополнительно выводится на экранвремя работы ядра Компиляция:$nvcc –arch=sm_20 –Xptxas –v kernel.cu -c –o kernel.oЛинковка проекта$g++ -L/toolkit_install_dir/lib64 –lcudart test.o kernel.o –o test Попросили слинковаться с libcudart.so, указали где она можетлежать$nvcc test.o kernel.o –o test nvcc –v test.o kernel.o –o test покажет какая конкретно командавызваласьТакже можно расположить весь код в *.cu файле и не пользоваться*.cpp вообщеЗапуск В результате компиляции и линковки получаемобычный исполняемый файл Запускаем из командной строки обычным способом$./test 1024Глобальная памятьDevice Расположена в DRAM GPU Объём до 6Gb Параметр устройства totalGlobalMem Кешируется в: L2 – на устройствемаксимальный размер 1536 KBПараметр устройства l2CacheSize L1 (только на Fermi) – на каждоммультипроцессоремаксимальный размер 48KBминимальный размер 16KBSMSMSMCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreL1 cacheL1 cacheL2 cacheDevice MemoryL1 cacheВыделение Динамически с хоста через cudaMalloc() Динамически из ядер через malloc() Статически - глобальная переменная с атрибутом__device__Динамически с хоста__global__ void kernel(int *arrayOnDevice) {arrayOnDevice[threaIdx.x] = threaIdx.x;}int main() {size_t size = 0;void *devicePtr = NULL;int hostMem[512];cudaMalloc(&devicePtr, sizeof(hostMem));cudaMemcpy(devicePtr, hostMem, size, cudaMemcpyHostToDevice);kernel<<<1,512>>>(devicePtr);}Статически__device__ int arrayOnDevice[512]__global__ void kernel() {arrayOnDevice[threaIdx.x] = threaIdx.x;}int main() {size_t size = 0;void *devicePtr = NULL;int hostMem[512];cudaGetSymbolSize(&size, arrayOnDevice);cudaMemcpyToSymbol(arrayOnDevice, localMem, size);kernel<<<1,512>>>(devicePtr);}Статически__device__ int arrayOnDevice[512]__global__ void kernel() {arrayOnDevice[threaIdx.x] = threaIdx.x;}int main() {size_t size = 0;int hostMem[512];void *devicePtr = NULL;cudaGetSymbolSize(&size, arrayOnDevice);cudaGetSymbolAddress(&devicePtr, arrayOnDevice);cudaMemcpy(devicePtr, hostMem, size, cudaMemcpyHostToDevice);kernel<<<1,512>>>();}сuda*Symbol* Переменные с атрибутами __device__ и __constant__находятся в глобальной области видимости и хранятсяобъектном модуле как отдельные символы Память под них выделяется статически при стартеприложения, как и под обычные глобальные переменные Работать с ними на хосте можно через функцииcudaMemcpyToSymbol , cudaMemcpyToSymbolAsync,cudaGetSymbolAddress, cudaMemcpyFromSymbol,cudaMemcpyFromSymbolAsynс, cudaGetSymbolSizeДинамически из ядер#include <stdlib.h>__global__ void kernel() {size_t size = 1024 * sizeof(int);int *ptr = (int *)malloc(size);memset(ptr, 0, size)free(ptr)}int main() {cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128*1024*1024);kernel<<<1, 128>>>();}Динамически из ядер malloc() из ядра выделяет память в куче Не освобождается между запусками ядер Освобождение по free() только с устройства Компилировать с –arch=sm_20 Доступны memcpy(), memset()Динамически из ядер Память под кучу выделяется на устройстве приинициализации CUDA runtime и освобождается призавершении программы После создания размер кучи не может быть изменен По-умолчанию 8мб Можно задать до первого вызова ядра c malloc черезcudaDeviceSetLimit(cudaLimitMallocHeapSize, N)Режимы работы кеша L1 Кеш может работать в двух режимах: 48KB и 16KB Переключение режимов: cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig)Устанавливает режим работы кеша cacheConfig для всего устройства cudaFuncSetCacheConfig ( const void* func,cudaFuncCache cacheConfig )Устанавливает режим работы кеша cacheConfig для всего отдельногоядраРежимы работы кеша L1 cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig) Возможные режимы:cudaFuncCachePreferNone – без предпочтений(по умолчанию).Выбирается последняя использованная конфигурация.

Начальнаяконфигурация – 16KB L1cudaFuncCachePreferShared: 16КB L1cudaFuncCachePreferL1: 48KB L1 cudaFuncSetCacheConfig ( const void* func,cudaFuncCache cacheConfig ) По умолчанию - cudaFuncCachePreferNone - запускать с режимомустройстваТранзакции Глобальная память оптимизирована с целью увеличенияполосы пропускания Отдать максимум данных за одно обращениеТранзакции Транзакция – выполнение загрузки из глобальной памятисплошного отрезка в 128 байт, с началом кратным 128 (naturallyaligned) Инструкция обращения в память выполняется одновременнодля всех нитей варпа (SIMT) Выполняется столько транзакций, сколько нужно дляпокрытия обращений всех нитей варпа Если нужен один байт – все равно загрузится 128Шаблоны доступаОбращения нитей варпа….512640768Обращения нитей варпа….512640Все обращенияумещаются в однутранзакциюПорядок не важен,главное, чтобы попадалив одну кеш-линию768Шаблоны доступаОбращения нитей варпа….512640Обращения нитей варпаЗапрашивается 128 байт, но с невыровненного адреса2 транзакции - 256 байт768Запрашивается 128 байт, нообращения разбросаны впределах трёх кеш-линий 3 транзакции – 384 байта….512640768Кеш-линии Ядра взаимодействуют не с памятью напрямую, а с кешами Транзакция – выполнение загрузки кеш-линии У кеша L1 кеш-линии 128 байт, у L2 - 32 байта, naturally aligned Кеш грузит из памяти всю кеш-линию, даже если нужен одинбайт Можно обращаться в память минуя кеш L1 Транзакции будут по 32 байтаТранзакции: L1 включенВарпL1512640768L2512576640768Транзакции: L1 выключенВарпL1512640768L2512576640768Включение \ отключение L1 Кеширование в L1 можно отключить при компиляции nvcc -Xptxas -dlcm=caс кешированием в L1 (по умолчанию) nvcc -Xptxas -dlcm=cgВ бинарном коде обращения в глобальную память будуттранслированы в инструкции, не использующие кеш L1 привыполненииРазличия именно на уровне бинарного кода – другие инструкцииассемблераШаблоны доступа: L1 выключенОбращения нитей варпа….512576640Обращения нитей варпаЗапрашивается 128 байт, но сневыровненного адреса.

Умещаются вчетыре кеш-линии L24 транзакции по 32 байта – 128 байт768Запрашивается 128 байт, нообращения разбросаны по пятикеш-линиям L25 транзакций по 32 – 160 байт….512608640768Вывод Если в ядре не используется общая память (см. далее), тозаведомо стоит включить cudaFuncCachePreferL1 Если разреженный доступ – кеширование в L1 отключаем В общем случае, стоит проверить производительность работывсех 4-х вариантов: (-dlcm=ca, -clcm=cg)x(16KB, 48KB)Матрицы и глобальная памятьПусть транзакция – 8*4=32 байта, адрес транзакции выровнен по 32 байтаЕсли ширина матрицы не кратна 32 байтам – большая часть строк не выровнена0713Матрицы и глобальная памятьОбращения варпа072 транзакции!13Матрицы и глобальная память Матрицы хранятся в линейном виде, по строкам Пусть длина строки матрицы – 480 байт (120 float) обращение – matrix[idy*120 + idx]Адрес начала каждой строки, кромепервой, не выровнен по 128 –обращения варпов в память будутвыполняться за 2 транзакции512Строка 0992Строка 11472Строка 21952Строка 3Матрицы и глобальная память Дополним каждую строку до размера, кратного 128 байтам – в нашемслучае, 480 + 32 = 512, это наш pitch – фактическая ширина в байтах Эти байты никак не будут использоваться, т.е.

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)Выделение памяти с «паддингом» cudaError_t cudaMallocPitch (void ** devPtr, size_t * pitch,size_t width, size_t height) Адрес элемента (Row, Column) матрицы, выделенной при помощиcudaMallocPitch:T* pElement = (T*)((char*) devPtr + Row * pitch) + ColumnКопирование в матрицу с padding-ом cudaError_t cudaMemcpy2D ( void* dst, size_t dpitch, constvoid* src, size_t spitch,size_t width, size_t height,cudaMemcpyKind kind )dst - указатель на матрицу, в которую нужно копировать ,dpitch – фактическая ширина её строк в байтахsrc - указатель на матрицу из которой нужно копировать,spitch – фактическая ширина её строк в байтахwidth – сколько байтов каждой строки нужно копироватьheight – число строкkind – направление копирования (как в обычном cudaMemcpy)Копирование в матрицу с padding-ом cudaError_t cudaMemcpy2D ( void* dst, size_t dpitch, constvoid* src, size_t spitch,size_t width, size_t height,cudaMemcpyKind kind ) Из начала каждой строки исходной матрицы копируется по 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] При первом чтении варпу нужно всего 4 байта, а скачается 128float **A;A[i][j] = 1;Выводы Обращения нитей варпа в память должны быть пространственнолокальными Начала строк матрицы должны быть выровнены Массивы структур -> структура с массивами 16КB vs 48KB L1 Избегаем косвенной адресации Избегаем обращений нитей варпа к столбцу матрицы В случае сильно разреженного доступа проверяем работу сотключенным кешемend.

Характеристики

Список файлов лекций