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

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

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

завершениявыполнения всех команд, в т.ч. запуска ядраcudaEventElapsedTime( &time, start, stop ); // Время между двумя событиямиcudaEventDestroy( start );cudaEventDestroy( stop );Особое отношение к .cuПри работе с CUDA используются расширения Си++: Конструкция запуска ядра <<< …. >>> Встроенные переменные threadIdx, blockIdx Квалификаторы __global__ __device__ и т.д. …. Эти расширения могут быть обработаны только в *.cu файлах! cudafe не запускается для файлов с другим расширением В этих файлах можно не делать #include <cuda_runtime.h> Вызовы библиотечных функций вида cuda* можно располагать в*.cpp файлах Они будут слинкованы обычным линковщиком из библиотекиlibcudart.soКомпиляция хост-кодаВ файле test.cpp : Основной хост-код.

Т.к. конструкцию запуска ядра в *.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 Объём до 4Gb Параметр устройства totalGlobalMem Кешируется в кешах L1 и L2:SMSMSMCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCoreCore L1 – на каждом мультипроцессоремаксимальный размер 48KBминимальный размер 16KB L2 – на устройствемаксимальный размер 768 KBПараметр устройства l2CacheSizeL1 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 Переключение режимов: cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig)Устанавливает режим работы кеша cacheConfig для всего устройства Возможные режимы: cudaFuncCachePreferNone – без предпочтений(по умолчанию).

Выбираетсяпоследняя использованная конфигурация. Начальная конфигурация – 16KB L1 cudaFuncCachePreferShared: 16КB L1 cudaFuncCachePreferL1: 48KB L1 cudaError_t cudaFuncSetCacheConfig ( const void* func,cudaFuncCachecacheConfig )Устанавливает режим работы кеша cacheConfig для заданной функции func По умолчанию - 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 Кеширование в 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)Транзакции: L1 включенВарпL1512640768L2512576640768Транзакции: L1 выключенВарпL1512640768L2512576640768Шаблоны доступа: L1 включенОбращения нитей варпа….512640768Обращения нитей варпа….512640Все обращенияумещаются в однутранзакциюПорядок не важен,главное, чтобы попадалив одну кеш-линию768Шаблоны доступа: L1 включенОбращения нитей варпа….512640Обращения нитей варпаЗапрашивается 128 байт, но с невыровненного адреса2 транзакции - 256 байт768Запрашивается 128 байт, нообращения разбросаны впределах трёх кеш-линий 3 транзакции – 384 байта….512640768Шаблоны доступа: L1 выключенОбращения нитей варпа….512576640Обращения нитей варпаЗапрашивается 128 байт, но сневыровненного адреса.

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

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

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