ТЕХНОЛОГИЯ ПРОГРАММИРОВАНИЯ CUD. Учебное пособие казань 2017 2 удк 004. 4, 681. 3 Ббк 32. 81 Печатается по постановлению Редакционноиздательского совета
 Скачать 1.28 Mb.
|
|
Глава 2. Свойства устройства В этом разделе опишем, какие существуют видеокарты, и приведем характеристики некоторых из них. Также рассмотрим функцию, благодаря которой можно узнать параметры и характеристики видеокарты, установленной непосредственно на устройстве. Архитектура видеокарт NVidia Рассмотрим подробнее физическое устройство видеокарт. На рис. 4 приведена схема архитектуры видеокарты GeForce 8800 GT. Центральный процессор обменивается данными с GPU через мост (Bridge). Вычислительная часть видеокарты состоит из семи текстурных процессорных кластеров (Texture Processor Cluster), каждый из которых содержит два мультипроцессора (SM), блок текстур (Texture Unit) и совмещенный текстурный кэш и кэш первого уровня (Tex/L1). Рис. 4. Архитектура видеокарты GeForce 8800 GT 17 Блоки текстур осуществляют выборку и фильтрацию текстурных данных, необходимых для построения сцены. Число текстурных блоков в видеочипе определяет текстурную производительность и скорость выборки из текстур. Текстурная производительность является одним из важнейших параметров видеочипов. Особое влияние этот параметр оказывает на скорость при использовании трилинейной и анизотропной фильтраций, требующих дополнительных текстурных выборок. Мультипроцессор, в свою очередь, содержит 8 потоковых процессоров (Streaming Processors) или ядер CUDA (CUDA cores). Таким образом, общее число SP = 14 * 8 = 112. Также каждый SM содержит по два блока для вычисления специальных (трансцендентных) функций (SFU), таких как синус, косинус и т.д. В состав SM входит разделяемая память (Shared Memory), блок инструкций нитей (Multithreaded Instruction Unit), константный кэш (C-Cache) и кэш инструкций (L-Cache). Таблица 2. Сравнительные характеристики видеокарт GF 8800 GeForce 8800 Ultra GeForce 8800 GTX GeForce 8800 GTS GeForce 8800 GT GeForce 8800 GS Stream процессоры 128 128 96 112 96 Частота ядра (МГц) 612 575 500 600 550 Частота шейдерного блока (МГц) 1500 1350 1200 1500 1375 Частота памяти (МГц) 1080 900 800 900 800 Объем памяти 768MB 768MB 640MB и 320MB 512MB 384MB Интерфейс памяти 384-bit 384-bit 320-bit 256-bit 192-bit Полоса пропускания памяти (ГБ/с) 103.7 86.4 64 57.6 38.4 18 Вычислительная часть устройства соединяется с видеопамятью (DRAM) посредством кэша второго уровня (L2) и блоков операций растеризации (ROP). Блоки опреаций растеризации осуществляют операции записи рассчитанных видеокартой пикселей в буферы и операции их смешивания (блендинга). Производительность блоков ROP влияет на филлрейт (cкорость заполнения пикселами), а это является одной из основных характеристик видеокарт. Рис. 5. Архитектура нового мультипроцессора семейства Volta GV100 19 Надо иметь в виду, что видеокарты одной линейки имеют одинаковую производительность, и, следовательно, одинаковое число SP в MP. Однако остальные характеристики могут отличаться. В таблице 2 приведены характеристики для GeForce 8800. Рассмотрим также новый, недавно анонсированный, мультипроцессор семейства Volta (Рис. 5). Каждый SM содержит четыре блока, в которых есть помимо ядер выполняющих целочисленные вычисления, ядра для работы с вещественными 32 и 64-битными числами, а также новые, появившиеся в данной линейке процессоров, тензорные ядра (Tensor Cores), разработанные специально для обучения нейросетей в задачах глубокого обучения. На Рис. 5 таже представлены планировщик варпов (Warp Scheduler), блок выборки (Dispatch Unit), регистры мультипроцессора (Register File) и блоки загрузки и сохранения (LD/ST – Load/Store Units). Типы видеокарт, поддерживающих CUDA Приведем в таблице 1 характеристики некоторых видеокарт, поддерживающих технологию CUDA. Первая колонка таблицы – название видеокарты: указывается модель и объем памяти, в названии GT обозначает обычные видеокарты низкого уровня производительности, GTX обозначает видеоадаптеры среднего и высокого уровня, Ti или Titan обозначает десктопные карты. Вторая колонка обозначает маркировку кристаллов: GT – Tesla, GF – Fermi, GK – Kepler, GM – Maxwell, GP – Pascal. Третья и четверая колонки фактически показывают производительность видеокарты. Цифры в последней колонке, разрядности шины памяти, указывают на пропускную способность памяти. Отметим, что частота ядра, как правило, немного ниже частот шейдерных блоков. Частота же памяти в современных видеокартах существенно выше частоты ядра. 20 Таблица 1. Характеристики некоторых видеокарт GeForce GPU / кристалл Кол-во процессоров- ядер / текстурных блоков / ROP – блоков растеризации Частота ядра GPU / шейдерного блока / памяти в MHz. Разрядность шины пямяти Titan X 12gb GP102 3584/224/96 1417/1530/10000 384 bit GTX 1080 8gb GP104 2560/160/64 1607/1733/10000 256 bit GTX 1060 6gb GP106 1280/80/48 1506/1709/8000 192 bit GTX 980 Ti 4gb GM200 2816/176/96 1000/1075/7000 384 bit GTX 980 4 gb GM204 2048/104/64 1126/1216/7000 256 bit GTX 960 2 gb GM206 1024/64/32 1126/1178/7000 128 bit Titan Z 12gb 2xGK110 5760/448/96 705/876/7000 2×384 GTX 780 Ti 3gb GK110 2880/240/48 875/928/7000 384 bit GTX 760 2 gb GK104 1152/96/32 980/1033/6008 256 bit GTX 750 1 gb GM107 512/32/16 1020/1085/5000 128 bit GTX 680 2 gb GK104 1538/128/32 1006/1058/6008 256 bit GTX 660 2 gb GK106 960/80/24 980/1033/6008 192 bit GT 640 2 gb GK107 384/32/16 900/900/1784 128 bit GT 630 1 gb GF108 96/16/4 810/1620/3200 128 bit GT 610 1 gb GF119 48/8/4 810/1620/1800 64 bit GTX 580 1.5 gb GF110 512/64/48 772/1544/4008 384 bit GTX 480 1.5 gb GF100 480/60/48 701/1402/3700 384 bit GTX 560 1 gb GF114 336/56/32 675/1620/4004 256 bit GTX 460 1 gb GF104 336/56/32 675/1350/3600 256 bit GTX 550 Ti 1gb GF116 192/32/24 900/1800/4100 192 bit GTS 450 1 gb GF106 192/32/16 783/1566/3600 128 bit GT 440 1 gb GF108 96/16/4 810/1620/3200 128 bit GT 520 1 gb GF119 48/8/4 810/1620/1800 64 bit GTX 280 1 gb GT200 240/80/32 602/1296/2215 512 bit GTS 250 1 gb G92b 128/64/16 738/1836/2200 256 bit GT 240 1 gb GT215 96/32/8 550/1340/3400 128 bit GT 220 1 gb GT216 48/16/8 625/1360/1580 128 bit GT 210 1 gb GT218 16/8/4 589/1402/1000 64 bit 9800 GT 1 gb G92 112/56/16 600/1500/1800 256 bit 9600 GT 1 gb G94 64/32/16 650/1625/1800 128 bit 9500 GT 1 gb G96 32/16/8 550/1400/1600 128 bit 8400 GS 512 mb G86 16\8\4 450/900/800 64 bit 21 Функция cudaGetDeviceProperties() Получим от видеокарты информацию о ее возможностях и разберемся, что они дают. Для определения свойств видеокарты используется функция cudaGetDeviceProperties(), принимающая в качестве параметра номер видеокарты и возвращающая в структуре cudaDeviceProp интересующие нас значения. Выведем на экран полученные переменные. Приведем основные переменные структуры cudaDeviceProp: struct cudaDeviceProp { char name[256]; size_t totalGlobalMem; size_t sharedMemPerBlock; int regsPerBlock; int warpSize; size_t memPitch; int maxThreadsPerBlock; int maxThreadsDim[3]; int maxGridSize[3]; size_t totalConstMem; int clockRate; int deviceOverlap; int multiProcessorCount; int kernelExecTimeoutEnabled; int integrated; int canMapHostMemory; int concurrentKernels; … } Здесь: • name[256] представляет собой ASCII строку идентификации устройства; • totalGlobalMem – общий объем глобальной памяти, доступной на устройстве в байтах; 22 • sharedMemPerBlock – максимальный объем общей памяти, доступной для нити блока в байтах; • regsPerBlock – максимальное количество 32-битных регистров, доступных для нити блока; • warpSize – размер Варпа в потоках; • maxThreadsPerBlock – максимальное количество потоков на блок; • maxThreadsDim[3] содержит максимальный размер каждого измерения блока; • maxGridSize[3] содержит максимальный размер каждого измерения сетки; • totalConstMem – общий объем константной памяти, доступной на устройстве в байтах; • clockRate – тактовая частота в килогерцах; • deviceOverlap равно 1, если устройство может одновременно копировать память между хостом и устройством во время выполнения ядра, или 0, если нет; • multiProcessorCount – число мультипроцессоров на устройстве; • kernelExecTimeoutEnabled равно 1, если существует ограничение времени выполнения для ядер, выполненных на устройстве, или 0, если нет. • integrated равен 1, если устройство является интегрированным (материнская плата) GPU и 0, если оно является дискретным. • canMapHostMemory равно 1, если устройство может отображать память хоста в адресное пространство CUDA для использования с cudaHostAlloc() / cudaHostGetDevicePointer(), или 0, если нет; • concurrentKernels равно 1, если устройство поддерживает выполнение нескольких ядер в том же контексте одновременно, или 0, если нет. 23 Приведем пример использования функции cudaGetDeviceProperties: cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp, 0);//определение параметров GPU с номером 0 printf("Device name : %s\n", deviceProp.name); printf("Total global memory : %d MB\n", deviceProp.totalGlobalMem / 1024 / 1024); printf("Shared memory per block : %d\n", deviceProp.sharedMemPerBlock); printf("Registers per block : %d\n", deviceProp.regsPerBlock); printf("Warp size : %d\n", deviceProp.warpSize); printf("Memory pitch : %d\n", deviceProp.memPitch); printf("Max threads per block : %d\n", deviceProp.maxThreadsPerBlock); printf("Max threads dimensions : x = %d, y = %d, z = %d\n", deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]); printf("Max grid size: x = %d, y = %d, z = %d\n", deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]); printf("Clock rate: %d\n", deviceProp.clockRate); printf("Total constant memory: %d\n", deviceProp.totalConstMem); printf("Compute capability: %d.%d\n", deviceProp.major, deviceProp.minor); printf("Texture alignment: %d\n", deviceProp.textureAlignment); printf("Device overlap: %d\n", deviceProp.deviceOverlap); printf("Multiprocessor count: %d\n", deviceProp.multiProcessorCount); 24 printf("Kernel execution timeout enabled: %s\n", deviceProp.kernelExecTimeoutEnabled ? "true" : "false"); scanf(""); Посмотрим, что выдала программа: Device name : GeForce 940M Total global memory : 2048 MB Shared memory per block : 49152 Registers per block : 65536 Warp size : 32 Memory pitch : 2147483647 Max threads per block : 1024 Max threads dimensions : x = 1024, y = 1024, z = 64 Max grid size: x = 2147483647, y = 65535, z = 65535 Clock rate: 1176000 Total constant memory: 65536 Compute capability: 5.0 Texture alignment: 512 Device overlap: 1 Multiprocessor count: 3 Kernel execution timeout enabled: true Также важным моментом является и определение версии видеокарты. Следующий код выводит производительность видеокарты cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp, 0); printf("Device has Compute Capability %d.%d\n", deviceProp.major, deviceProp.minor); 25 Вопросы к разделу 1. Опишите примерную архитектуру современных видеокарт NVidia. 2. Какие параметры видеокарты характеризуют её производительность? 3. Какие основные характеристики видеокарты можно узнать с помощью функции cudaGetDeviceProperties()? Лабораторная работа 1. Определите следующие параметры видеокарты с поддержкой технологии CUDA: • наименование; • общий объем графической памяти; • объем разделяемой памяти в пределах блока; • число регистров в пределах блока; • размер варпа; • максимально допустимое число потоков в блоке; • версию вычислительных возможностей; • число потоковых мультипроцессоров; • тактовую частоту ядра; • объем кэша второго уровня; • ширину шины памяти видеокарты; • максимальную размерность при конфигурации потоков в блоке и блоков в сетке. 2. Сравните результаты, полученные программно, с информацией от производителя. Домашняя работа 26 1. Доработайте лабораторное задание 1, добавив в него вывод объема константной памяти и пиковую частоту видеокарты в МГц. 27 Глава 3. Простейшая программа на CUDA Каждая программа, использующая технологию CUDA, делится на две части: последовательная часть (код, выполняющийся на CPU) и параллельная часть (код, выполняющийся на GPU). Напишем программу «Hello, world!», которая будет вызвана на двух блоках и пяти нитях. #include "cuda_runtime.h" #include "device_launch_parameters.h" #include __global__ void HelloWorld() { printf("Hello world, %d, %d\n", blockIdx.x, threadIdx.x); } int main() { HelloWorld << <2, 5 >> >(); // хост ожидает завершения работы девайса cudaDeviceSynchronize(); // ожидаем нажатия любой клавиши getchar(); return 0; } Ядро, так называется функция, выполняющаяся на GPU. Эта функция имеет спецификатор __global__, который сообщает компилятору, что функция вызывается с CPU и выполняется на GPU. На ядро налагается ряд некоторых ограничений: не содержать рекурсию, не иметь переменное число аргументов, не иметь переменные static внутри себя. Вызов функции осуществляются следующим образом: имя_функции << 28 Тройные угловые скобки перед функцией, выполняющейся на GPU, – это синтаксис CUDA. Параметры в угловых скобках определяют количество и конфигурацию параллельных процессов GPU: размер сетки и размер блока. Функция копирования Рассмотрим более сложный пример использования технологии CUDA. Для написания программы, приведенной ниже, необходимо изучить принципы работы с глобальной памятью GPU. Глобальная память является самой медленной, и если необходима высокая производительность, то количество операций с глобальной памятью следует свести к минимуму. Назначение функции Синтаксис функции Выделение памяти cudaMalloc( void **devPtr, size_t size); Освобождение памяти cudaFree( void **devPtr); Копирование данных cudaMemcpy( void *dst, const void *src, size_t size, enum cudaMemcpyKind kind) В качестве параметра kind, задающее направление копирования, могут выступать: cudamemcpyHostToDevice, cudamemcpyDeviceToHost, cudamemcpyDeviceToDevice, cudamemcpyHostToHost. Итак, рассмотрим немного более сложную программу. Данная программа складывает два целых числа на GPU. //подключение библиотек #include "cuda_runtime.h" 29 #include "device_launch_parameters.h" #include #include // ядро __global__ void add( int *a, int *b, int *c ) { *c = *a + *b; } //главная функция int main() { // переменные на CPU int a, b, c; // переменные на GPU int *dev_a, *dev_b, *dev_c; int size = sizeof( int ); //размерность // выделяем память на GPU cudaMalloc( (void**)&dev_a, size ); cudaMalloc( (void**)&dev_b, size ); cudaMalloc( (void**)&dev_c, size ); // инициализация переменных a = 2; b = 7; // копирование информации с CPU на GPU cudaMemcpy( dev_a, &a, size, cudaMemcpyHostToDevice ); cudaMemcpy( dev_b, &b, size, cudaMemcpyHostToDevice ); // вызов ядра add<<< 1, 1 >>>( dev_a, dev_b, dev_c ); // копирование результата работы ядра с GPU на CPU 30 cudaMemcpy( &c, dev_c, size, cudaMemcpyDeviceToHost ); // вывод информации printf("%d + %d = %d\n", a, b, c); // очищение памяти на GPU cudaFree( dev_a ); cudaFree( dev_b ); cudaFree( dev_c ); return 0; } Замер времени работы части кода программы Одним из важных параметров при написании программы является время, затраченное на ее выполнение. Для замера времени выполнения операций на CPU есть множество возможностей. Для того, чтобы замерить время, которое было затрачено на выполнение вычислений на GPU, можно воспользоваться механизмом событий (CUDA events). Событие – это объект типа cudaEvent_t, используемый для обозначения «точки» среды вызовов CUDA. Приведем пример фрагмента кода, замеряющий время выполнения ядра на GPU: // инициализируем события cudaEvent_t start, stop; float elapsedTime; // создаем события cudaEventCreate(&start); cudaEventCtreate(&stop); // запись события cudaEventRecord(start, 0); // вызов ядра 31 cudaEventRecord(stop,0); // ожидание завершения работы ядра cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsedTime, start, stop); // вывод информации printf("Time spent executing by the GPU: %.2f millseconds\n", elapsedTime); // уничтожение события cudaEventDestroy(start); cudaEventDestroy(stop); Обработка ошибок После выполнения функций на GPU, они возвращают значения типа cudaError_t. Функции, отработавшие на девайсе, возвращают значение cudaSuccess, в противном случае возвращается код ошибки. Назначение функции Синтаксис функции Возвращает описание ошибки в виде строки cudaGetErrorString(cudaError_t error); Возвращает код последней ошибки cudaGetLastError(); Основные выводы ответа при обработке ошибок: 1. cudaSuccess – при удачном выполнении. 2. cudaErrorMemoryAllocation – при ошибке выделения памяти. 3. cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен). 4. cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте. 5. cudaErrorInvalidMemcpyDirection – неверное направление (например, перепутан источник и место-назначение копирования). 6. cudaErrorInitializationError – ошибка инициализации. 7. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном 32 запуске функции. 8. cudaErrorInvalidResourceHandle – неверный обработчик события. Приведем пример использования функции cudaGetLastError(). cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) printf("%s ", cudaGetErrorString(err)); В данном примере проверяется, есть ли ошибка у последней команды, выполненной на девайсе. Если ошибка есть, то выводится ее описание. Вопросы к разделу 1. На какие условно две части можно разделить CUDA-программу? 2. Для чего нужна глобальная память? Какие имеются функции в CUDA для работы с глобальной памятью? 3. Опишите механизм, использующийся в CUDA-программах, для замера времени работы ядра. 4. Какие механизмы в CUDA-программах используются для обработки ошибок? Лабораторная работа 1. Измерьте скорость копирования данных (ГБ/сек) между CPU и GPU. 2. Напишите программу, вычисляющую число пи методом интегрирования четверти окружности единичного радиуса (можно использовать формулу для площади четверти окружности: 4 / 1 1 0 2 π = − = ∫ dx x S ). В программе предусмотрите проверку на ошибку выполнения функции. 33 Домашняя работа 1. Напишите программу вычисления дзета-функции Римана путем суммирования ряда из обратных степеней. |