Структура ПО CUDA Приложения, использующие возможности CUDA для параллельной обработки данных, взаимодействуют с GPU через несколько разных программных интерфейсов (обычно интерфейсы CUDA-runtime и CUDA-driver одновременно не используются, т.е. используется либо CUDA-runtime, либо CUDA-driver): CPU GPU CUDA- библиотеки (CuFFT, CuBLAS) CUDA- runtime cudart.dll CUDA- driver П Р И Л О Ж Е Н И Е

CUDA-программы Программы для CUDA пишутся на "расширенном" С, при этом их параллельная часть (так называемые «ядра») выполняется на GPU, а обычная часть – на CPU. Компоненты CUDA, принимающие участие в подготовке и исполнении приложения, автоматически осуществляют разделение частей и управление их запуском. Ядро есть исполнение сетки (grid) блоков потоков (thread blocks) Блок потоков – это набор потоков, выполняемых одним мультипроцессором. Потоки одного блока могут взаимодействовать через разделяемую память и синхронизироваться между собой. Единовременно выполняется в точности одна сетка (одно ядро), поэтому сетка не имеет координат. Блоки и потоки имеют координаты и могут рассматриваться как линия (1D), поверхность (2D) и объем (3D).

Block(i) Thread (0,0) Thread (1,0) Thread (m,0) Thread (0,1) Thread (1,1) Thread (m,1) Thread (0,p) Thread (1,p) Thread (m,p) Thread (0,0) Thread (1,0) Thread (m,0) Thread (0,1) Thread (1,1) Thread (m,1) Thread (0,p) Thread (1,p) Thread (m,p) GPU Grid2 Возможные структуры сетки и блоков потоков CPU Поток: Ядро 1 Ядро 2 Grid1 Block (0,0) Block (1,0) Block (0,1) Block (1,1) Block (k,0) Block (k,1) Block (0,n) Block (1,n) Block (k,n) Block (0) Block (1) Block (n) Thread (n,0,0) Thread (n,1,0) Thread (n.m,0) Thread (n,0,1) Thread (n,1,1) Thread (n,m,1) Thread (n,0,p) Thread (n.1,p) Thread (n,m,p) Block(s,t) Thread (0) Thread (1) Thread (m)

Пример: программа сложения двух векторов //Попарное сложение элементов вектора: //один поток выполняет сложение одной пары элементов Код программы ядра для GPU: __global__ void vectorAdd( float *resultV, float *srcA, float *srcB ){ //вычисление индекса складываемых элементов: int idx = blockDim.x * blockIdx.x + threadIdx.x;//!!! //вычисление результата: resultV[idx] = srcA[idx] + srcB[idx]; }

Пример: программа сложения двух векторов Код программы CPU для сложения векторов: const int N = ; //Размер вектора в элементах = 256 * 4096 const int dataSize = N * sizeof(float);//размер вектора в байтах //Выделение памяти CPU: float *h_A = (float *)malloc( dataSize ); float *h_B = (float *)malloc( dataSize ); _ float *h_C = (float *)malloc( dataSize ); //Выделение памяти GPU float *d_A, *d_B, *d_C; cudaMalloc((void **)&d_A, dataSize)); cudaMalloc((void **)&d_B dataSize)); cudaMalloc((void **)&d_C, dataSize)); //Инициализировать векторы h_A[], h_B[] … Продолжение на следующем слайде

Пример: программа сложения двух векторов //Скопировать входные данные в GPU для обработки cudaMemcpy(d_A, h_A, dataSize, cudaMemcpyHostToDevice) ); cudaMemcpy(d_B, h_B, dataSize, cudaMemcpyHostToDevice) ); //Запустить ядро из N / 256 блоков по 256 потоков //Предполагается, что N кратно 256 //Инициализация ядра в GPU: vectorAdd >>(d_C, d_A, d_B); //Переписать результаты работы GPU в память CPU: cudaMemcpy(h_C, d_C, dataSize, cudaMemcpyDeviceToHost) );

Структура памяти GPU CPUGrid Block (0,0)Block (1,0) Global Memory Constant Memory Texture Memory Thread (0,0) Thread (1,0) Local Memory Registers Shared Memory Thread (0,0) Thread (1,0) Local Memory Registers Shared Memory …… … Registers: R/W, on chip, per thread Shared Memory: R/W, on chip, per block Local Memory: R/W, DRAM, uncached, per thread Global Memory: R/W, DRAM, uncached, per context Constant Memory: R, DRAM, cached, per context Texture Memory: R, cached, per context

Аппаратная реализация GPU Мультипроцессор N Мультипроцессор 2 Мультипроцессор 1 Кэш констант Кэш текстур Устройство управления АЛУ1 Регистры АЛУ2 Регистры АЛУk Регистры Разделяемая память (Shared Memory) Глобальная память устройства … CPU

Версии CUDA Поддерживаемые возможности (не указанные здесь возможности поддерживаются во всех версиях) Версия CUDA x Целочисленные atomic-функции, работающие над 32- разрядными словами в глобальной памяти НетДа Целочисленные atomic-функции, работающие над 64- разрядными словами в глобальной памяти НетДа Целочисленные atomic-функции, работающие над 32- разрядными словами в разделяемой памяти Warp-функции голосования Операции с плавающей точкой двойной точностиНетДа Дополнительные atomic-функции, работающие над 32- разрядными числами с плавающей точкой в глобальной и разделяемой памяти НетДа Дополнительные функции синхронизации

Технические характеристики Версия CUDA x Максимальные x- или y- размерности сетки блоков65535 Максимальное количество нитей в блоке Максимальные x- или y- размерности блока Максимальная z- размерность блока64 Размер Warp в блоках32 Максимальное количество резидентных блоков в мультипроцессоре 8 Максимальное количество резидентных warp-ов в мультипроцессоре Максимальное количество резидентных нитей в мультипроцессоре Количество 32-битных регистров в мультипроцессоре 8 K16 K32 K

Технические характеристики Версия CUDA x Максимальный объем разделяемой памяти в мультипроцессоре 16 KB48 KB Количество банков разделяемой памяти1632 Объем локальной памяти одной нити16 KB512 KB Размер памяти констант64 KB Размер кэша памяти констант одного мультипроцессора 8 KB Размер кэша памяти текстур одного мультипроцессора От 6 KB до 8 KB Максимальная ширина ссылки на 1D текстуру, связанную с CUDA массивом Максимальная ширина ссылки на 1D текстуру, связанную с линейной памятью 2 27

Технические характеристики Версия CUDA x Максимальная ширина и высота ссылки на 2D текстуру, связанную с линейной памятью или CUDA массивом x x Максимальная ширина, высота и глубина ссылки на 3D текстуру, связанную с линейной памятью или CUDA массивом 2048 x 2048 x 2048 Максимальное количество текстур, которые могут быть связаны с ядром 128 Максимальная ширина ссылки на 1D поверхность, связанную с CUDA массивом – 8192 Максимальная ширина и высота ссылки на 1D поверхность, связанную с CUDA массивом 8192 x 8192 Максимальное количество поверхностей, которые могут быть связаны с ядром 8 Максимальное количество инструкций на ядро2 х 10 6

Модель исполнения программ в GPU: Блок потоков всегда исполняется только на одном мультипроцессоре Каждый блок потоков состоит из warpов. Warpом называется группа потоков фиксированного размера, состоящая из скалярных потоков с последовательными координатами, исполняемая в режиме ОКМД. Локальные регистры динамически делятся между всеми АЛУ, исполняющими блок потоков. Общие переменные блока потоков хранятся в разделяемой памяти мультипроцессора

Модель исполнения программ в GPU: Мультипроцессор может обрабатывать несколько блоков потоков одновременно. При этом: Локальные регистры и разделяемая память динамически делятся между всеми потоками (блоками потоков), работающими одновременно Уменьшение числа используемых локальных регистров на поток и размера используемой разделяемой памяти на блок позволяет увеличивать число потоков / блоков потоков, которые могут одновременно исполняться мультипроцессором Слишком большой размер блока потоков, использующих слишком много локальных регистров, может привести к сбою процесса исполнения ядра. На логическом уровне блоки потоков всегда изолированы друг от друга. Один поток CPU управляет только одним GPU, каждый поток CPU управляет своим GPU

Рекомендуемая разработчиками CUDA технология программирования: 1.Разбить обрабатываемые данные на фрагменты, целиком помещающиеся в разделяемую память (Shared memory). 2.Обеспечить загрузку данных в глобальную память. 3.Запустить ядро, для каждого блока: 3.1. Обеспечить перемещение фрагментов данных в локальную память блока 3.2. Выполнить нужные вычисления 3.3. Скопировать сформированные результаты обработки назад в глобальную память 4.Обработать полученные результаты.

Среда программирования CUDA Расширение среды программирования языка C: –Расширения синтаксиса для написания кода для GPU, реализуемые дополнительным препроцессором nvcc. –Run-time библиотеки: Общая часть – встроенные векторные типы данных и набор функций, исполняемых и на CPU, и на GPU CPU-компонента, для доступа и управления одним несколькими GPU (драйвер) GPU-компонента, предоставляющая функции, специфические только для GPU

Расширения языка: Дополнительные виды функций: Исполняется наВызывается из __host__ float HostFunc()CPUCPU __global__ void KernelFunc()GPUCPU __device__ float DeviceFunc()GPUGPU __global__ является функцией-ядром, она не может возвращать никакого значения (всегда void) Нельзя получить указатель на функцию вида __device__ В функциях, исполняемых на GPU: Недопустима рекурсия Нет статических переменных внутри функций Количество аргументов не может быть переменным

Виды переменных Тип ОбластьВремя памяти видимостижизни __device__ __shared__ int shared blockblock __device__ int global gridcontext __device__ __constant__ intconstant gridcontext Ключевое слово __device__ необязательно, если используется __shared__ или __constant__ Локальные переменные без идентификатора вида хранятся в регистрах, за исключением больших структур или массивов, располагающихся в локальной памяти Тип памяти указателя адаптируется к типу присвоенного ему выражения

Встроенные векторные типы [u]char[1..4], [u]short[1..4], [u]int[1..4], [u]long[1..4], float[1..4] uint4 эквивалентно Struct uint4 { unsigned int x; unsigned int y; unsigned int z; unsigned int w; }; Доступ к полям по именам: x, y, z, w, например: uint4 param; int y = param.y; int abc = param.w; dim3 - cиноним uint3. Обычно тип dim3 используется для задания параметров сетки (grid)

Запуск ядра (программы GPU) __global__ void CSolve(int nDim, float * pfMatr, float * pfVect, float * fAcc, int * pG) {//тело функции, являющейся ядром … } dim3 dimBlock = dim3(nDim, 1); dim3 dimGrid = dim3(nDim, 1); float *d_pfMatr; float *d_pfVect; float *d_fAcc; int *d_pG; //Выделение памяти, ввод исходных данных, перенос в память GPU: … //При запуске программному обеспечению CUDA передаются // обязательные параметры конфигурации сетки (grid): //Собственно запуск ядра: CSolve >>(nDim, d_pfMatr, d_pfVect, d_fAcc, d_pG);

Параметры запуска ядра KernelFunc (...); DimGrid: определяет размерность сетки (в блоках). DimBlock: определяет размерность блока (в потоках). Опциональные SharedMemBytes байт: Выделяются в дополнение к статически объявленным разделяемым переменным в ядре. Отображаются на любую переменную вида: extern __shared__ float DynamicSharedMem[]. Запуск ядра асинхронен, сразу после вызова функции управление немедленно возвращается к следующему оператору программы CPU.