Половинкин А.Н.

Вычисления общего назначения на GPU Архитектура GPU Программная модель выполнения на CUDA Программирование с использованием CUDA Настольная вычислительная суперкомпьютерая система Nvidia Tesla D870

CPU GPU предназначен для вычислений, параллельных по данным: одна и та же операция выполняется над многими данными параллельно (SIMD) в которых отношение вычислительных операций к числу операций по доступу к памяти велико Вместо кэша и сложных элементов управления на кристалле размещено большее число вычислительных элементов GPU

8 SP (Streaming Processor) - потоковые скалярные процессоры 2 SFU (Super Functions Unit) – предназначен для вычисления сложных математических функций RFn (Register File) Shared Memory – разделяемая память

ядро (kernel) – функция, выполняемая решеткой (grid) блоков потоков (threads block) блок потоков (threads block) – набор потоков, выполняющих одну функцию (kernel) на одном мультипроцессоре, способных общаться между собой посредством: разделяемой памяти (shared memory) точек синхронизации два потока из двух различных блоков не взаимодействуют между собой

каждый поток и блок потоков имеют идентификаторы каждый поток может определить, с какими данными он должен работать Block ID (1D или 2D) Thread ID (1D, 2D или 3D) данный подход упрощает адресацию памяти при обработке многомерных данных

registers (чтение/запись, одним SP*) local (чтение/запись, одним SP) shared (чтение/запись, всеми SP, входящими в состав MP**) constant cache (только чтение, всеми SP, входящими в состав MP) texture cache (только чтение, всеми SP, входящими в состав MP) device (global) (чтение/запись, всеми SP, входящими в состав всех MP) *SP – scalar processor **MP – multiprocessor

данные, расположенные в глобальной памяти, реально располагаются в памяти устройства (доступ к device memory много медленнее доступа к shared memory) общий подход к ускорению вычислений заключается в следующем: разбить множество обрабатываемых данных на подмножества, убирающиеся в shared memory обрабатывать каждое подмножество данных одним блоком потоков: загрузить подмножество данных из global memory в shared memory выполнить вычисления над элементами данных из подмножества скопировать результаты из shared memory в global memory

host = CPU device = GPU = набор мультипроцессоров device memory = собственная память GPU kernel (ядро) – подпрограмма, выполняемая на GPU grid (решетка) – массив блоков потоков, которые выполняют одно и то же ядро thread block (блок потоков) – набор потоков, которые выполняют ядро и могут взаимодействовать, используя общую память (shared memory)

Стандартный язык C для разработки параллельных приложений на GPU Библиотеки FFT (Fast Fourier Transform) и BLAS (Basic Linear Algebra Subroutine) Специализированный драйвер для вычислений, обеспечивающий быструю передачу данных между CPU и GPU Драйвер CUDA, обеспечивающий взаимодействие с OpenGL и DirectX

Поддержка видеокарт NVidia >= G80, Tesla, Quadro Поддержка Windows XP 32/64bit, Windows Vista 32/64 bit, Linux 32/64bit, Mac OS Комплект поставки CUDA driver CUDA toolkit CUDA SDK

API представляет собой расширение языка C Состав CUDA API: расширения языка C библиотека времени выполнения (runtime library): общий компонент, обеспечивающий встроенные векторные типы и подмножество C runtime library поддерживающее как host, так и device код host component, обеспечивающий управление и доступ к одному или нескольким устройствам с хоста device component, обеспечивающий функции, специфичные для устройства

КвалификаторВыполняется на:Вызывается с: __device__device __host__host __global__devicehost __global__ - определяет функцию ядро (kernel) должна возвращать результат типа void __device__ и __host__ могут использоваться совместно невозможно взять адрес __device__ функции функции, выполняемые на устройстве, не допускают: рекурсию объявление статических переменных внутри функции переменное число аргументов Пример: __global__ void KernelFunc(float arg);

__device__ объявляет переменную, размещаемую на GPU размещается в глобальном пространстве памяти; время жизни переменной совпадает со временем жизни приложения; доступ к переменной может быть осуществлен из всех потоков, выполняемых на устройстве, а также с хоста через библиотеки времени выполнения. __constant__ объявляет переменную, которая размещается в константном пространстве памяти; время жизни переменной совпадает со временем жизни приложения; доступ к переменной может быть осуществлен из всех потоков, выполняемых на устройстве, а также с хоста через библиотеки времени выполнения. __shared__ объявляет переменную, которая размещается в пространстве общей памяти блока потоков; время жизни переменной совпадает со временем жизни блока потоков; доступ к переменной может быть осуществлен из потоков, принадлежащих блоку потоков.

[u]char[1..4] [u]int[1..4] [u]long[1..4] float[1..4] double2

gridDim – переменная типа dim3, содержит текущую размерность решетки; blockIdx – переменная типа uint3, содержит индекс блока потоков внутри решетки; blockDim – переменная типа dim3, содержит размерность блока потоков; threadIdx – переменная типа uint3, содержит индекс потока внутри блока потоков; warpSize – переменная типа int, содержит размер «свёртки» (warp) в потоках. Замечание: данные переменные предназначены только для чтения и не могут быть изменены из вызывающий программы

перечисление устройств: cudaError_t cudaGetDeviceCount(int* count) – возвращает число доступных устройств; cudaError_t cudaGetDevice (int* dev) – возвращает используемое устройство cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp* prop, int dev) – возвращает структуру, содержащую свойства устройства выбор устройства: cudaError_t cudaChooseDevice(int* dev, const struct cudaDeviceProp* prop) – устанавливает устройство, на котором выполняется device код, в наибольшей степени соответствующее конфигурации cudaError_t cudaSetDevice(int dev) – устанавливает устройство, на котором выполняется device код; Замечание: Nvidia Tesla D870 представляется в виде двух устройств

выделение и освобождение памяти на устройстве: cudaError_t cudaMalloc(void** devPtr, size_t count) – выделяет память на устройстве и возвращает указатель на нее cudaError_t cudaFree(void* devPtr) – освобождает память на устройстве копирование данных между хостом и устройством: cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind) – копирует данные между хостом и устройством

Функция ядра должна быть вызвана с указанием конфигурации исполнения Конфигурация определяется использованием выражения специального вида >> между именем функции и списком ее аргументов, где: Dg – определяет размерность и размер сетки, так что Dg.x * Dg.y равно числу блоков потоков, которые будут запущены, Dg.z не используется. Db – определяет размерность и размер каждого блока потоков, Db.x * Db.y * Db.z равно числу потоков на блок. Ns – переменная типа size_t, определяет число байт в разделяемой памяти, которое дополнительно выделяется на блок добавление к автоматически выделенной компилятором памяти.

__global__ void KernelFunc() …. dim3 DimGrid(100, 50); dim3 DimBlock (8, 8, 8); size_t SMSize = 64; KernelFunc >>();

void __syncthreads() – синхронизирует все потоки внутри блока потоков; как только все потоки достигли данной точки, они продолжают свое выполнение используется, чтобы избежать RAW/WAR/ WAW конфликтов при доступе к shared или global памяти

#include const int N = 256; const int DATA_SZ = N * sizeof(float); float RandFloat(float low, float high) { float t = (float)rand() / (float)RAND_MAX; return (1.0f - t) * low + t * high; } // ядро, каждый поток вычисляет сумму элементов массивов A и B с индексами, // соответствующими индексу потока __global__ void vecAdd(float* A, float* B, float* C) { int i = threadIdx.x; C[i] = A[i] + B[i]; }

int main(int argc, char **argv) { float *h_A, *h_B, *h_C; float *d_A, *d_B, *d_C; int i; h_A = (float *)malloc(DATA_SZ); h_B = (float *)malloc(DATA_SZ); h_C = (float *)malloc(DATA_SZ); for(i = 0; i < N; i++) { h_A[i] = RandFloat(0.0f, 1.0f); h_B[i] = RandFloat(0.0f, 1.0f); }

// инициализация GPU CUT_DEVICE_INIT(argc, argv); // выделение памяти на GPU CUDA_SAFE_CALL( cudaMalloc((void **)&d_A, DATA_SZ) ); CUDA_SAFE_CALL( cudaMalloc((void **)&d_B, DATA_SZ) ); CUDA_SAFE_CALL( cudaMalloc((void **)&d_C, DATA_SZ) ); // копирование данных с хоста на GPU CUDA_SAFE_CALL( cudaMemcpy(d_A, h_A, DATA_SZ, cudaMemcpyHostToDevice) ); CUDA_SAFE_CALL( cudaMemcpy(d_B, h_B, DATA_SZ, cudaMemcpyHostToDevice) ); // вызов ядра vecAdd >>(d_A, d_B, d_C); // копирование вектора, содержащего сумму A и B с GPU на хост CUDA_SAFE_CALL( cudaMemcpy(h_C, d_C, DATA_SZ, cudaMemcpyDeviceToHost) );

// освобождение памяти на GPU CUDA_SAFE_CALL( cudaFree(d_C) ); CUDA_SAFE_CALL( cudaFree(d_B) ); CUDA_SAFE_CALL( cudaFree(d_A) ); free(h_C); free(h_B); free(h_A); // освобождение ресурсов устройства СUT_EXIT(argc, argv); }

Для каждого потока в свертке: чтение операндов выполнение инструкции запись результата Для повышения производительности необходимо: уменьшить число арифметических инструкций с низким throughput максимизировать использование доступной пропускной способности для каждого типа памяти

4 clock cycles: floating point сложение, умножение, умножение- сложение (multiply-add) integer сложение 24-bit integer умножение (__mul24) побитовые операции, сравнение, минимум, максимум, преобразование типов 16 clock cycles: вычисление обратного числа, 1 / sqrt(x), __logf(x) умножение 32-bit integers

Целочисленное деление и взятие остатка по модулю следует заменять битовыми операциями везде, где это возможно: если n=2^p, тогда i/n ~ i>>log2(n), i%n ~ i&(n-1) 32 clock cycles: __sinf(x), __cosf(x), __expf(x) (доступны только из кода, выполняемого на устройстве) Рекомендуется использовать везде, где это возможно, floating point данные и floating point версии арифметических функций

условные операторы и операторы циклов (if, switch, do, while, for) влияют на производительность приложений if (condition) { code1; } else { code2; }... потоки в свёртке condition == true condition == false... потоки в свёртке code 1 idle code2

Включают в себя операции чтения/записи глобальной, локальной и shared памяти. Выполнение одной инструкции доступа к памяти требует 4 clock cycles. Глобальная память обладает латентностью clock cycles __shared__ float shared[32]; __device__ float device[32]; shared[threadIdx.x] = device[threadIdx.x];

Существуют атомарные инструкции для чтения 32-bit, 64-bit и 128-bit машинных слов. __device__ type device[32]; type data = device[tid]; sizeof(type) должен быть равен 4, 8 или 16 данные должны быть выровнены по sizeof(type) Выравнивание обеспечивается компилятором автоматически для встроенных типов данных (float2, float4, …)

Размер и выравнивание структур обеспечивается директивой компилятора __align__ struct __align__(8) { float a; float b; }; struct __align__(16) { float a; float b; float c; }; 1 64-bit load instruction bit load instruction

Структуры размера больше 16 байт следует определять, используя __align__(16) struct { float a; float b; float c; float d; float e; }; struct __align__(16) { float a; float b; float c; float d; float e; }; 5 32-bit load instructions bit load instructions

доступ к глобальной памяти всеми потоками половины свертки (half warp), объединяется в 1 или 2 инструкции, при выполнении условий: потоки совершают доступ к 32-bit, 64-bit или 128-bit words. все 16 машинных слов должны лежать в одном и том же сегменте, размер которого равен размеру memory transaction size k-ый поток совершает доступ к k-му слову

shared memory состоит из блоков памяти равного размера, доступ к которым может быть осуществлен одновременно, - банков памяти банки в shared memory организованы таким образом, что последовательно идущие 32-bit words относятся к последовательно идущим банкам памяти Address 0Address Bank 0 Address 1Address Bank 1 Address 2Address Bank 2 Address 3Address Bank 3 Address 4Address Bank 4 Address 5Address Bank 5 Address 6Address Bank 6 Address 7Address Bank 7 Address 8Address Bank 8 Address 9Address Bank 9 Address 10Address Bank 10 Address 11Address Bank 11 Address 12Address Bank 12 Address 13Address Bank 13 Address 14Address Bank 14 Address 15Address Bank 15

__shared__ float shared[32]; float data = shared[BaseIndex + s * tid]; tid – thread ID, s – шаг доступа к элементам массива потоки с ID tid и tid+n вызовут bank conflict, если sn кратно числу банков m пусть d = НОД(m, s), тогда для того, чтобы избежать bank conflicts, необходимо, чтобы d = 1

BLAS (Basic Linear Algebra Subroutines) – набор базисных подпрограмм линейной алгебры. Данный набор является основой для функций из пакета LAPACK. Состоит из 3 уровней: 1. Операции над векторами (vector-vector) 2. Вектор-матричные операции (matrix-vector) 3. Операции над матрицами (matrix-matrix)

Имя любой процедуры BLAS имеет следующую структуру: () : символ, описывающий тип данных, с которым работает процедура. sвещественный, одинарной точности скомплексный, одинарной точности dвещественный, двойной точности zкомплексный, двойной точности

Для некоторых процедур и функций данные символы могут комбинироваться. Например, функция scasum принимает на вход массив комплексных чисел и возвращает вещественное значение. : для BLAS Level 1 определяет тип операции (например, dot – скалярное произведение, swap – перестановка элементов векторов местами), для BLAS Level 2 и 3 определяет тип матричного аргумента.

Некоторые функции BLAS возвращают индекс элемента массива. Независимо от того, какая версия библиотеки (Fortran или C) используется, элементы массива нумеруются с 1. Следовательно, при использовании C-версии из результата, который вернула функция, следует вычесть 1.

#include void main() { int n = 4; float x[4] = {1., 2., 3., -4.}; float y[4] = {2., 2., -1., 10}; int incx = 1, incy = 1; float alpha = 1.0; int imax = 0; saxpy(&n, &alpha, x, &incx, y, &incy); // y := alpha*x + y }

cublasInit() - инициализация CUBLAS (должна быть вызвана перед использованием любой другой CUBLAS функции) cublasShutdown() - освобождает ресурсы, используемые библиотекой CUBLAS на стороне хоста

cublasAlloc(int n, int elemSize, void **devicePtr) – создает объект в пространстве памяти GPU, содержащий массив из n элементов размера elemSize. Указатель на созданный объект размещается в devicePtr (данный указатель не должен впоследствии изменяться в коде, выполняемом на хосте) cublasFree(const void *devicePtr) – освобождает память, выделенную на GPU

cublasSetVector (int n, int elemSize, const void *x, int incx, void *y, int incy) – копирует n элементов вектора x (пространство памяти CPU) в вектор y (пространство памяти GPU) cublasGetVector (int n, int elemSize, const void *x, int incx, void *y, int incy) - копирует n элементов вектора x (пространство памяти GPU) в вектор y (пространство памяти CPU)

#include void main() { int n = 4; float x_H[4] = {1., 2., 3., -4.}; float y_H[4] = {2., 2., -1., 10}; float* x_D = 0; float* y_D = 0; int incx = 1, incy = 1; float alpha = 1.0; cublasInit(); cublasAlloc(n, sizeof(float), &x_D); cublasAlloc(n, sizeof(float), &y_D); cublasSetVector(n, sizeof(float), x_H, incx, x_D, incx); cublasSetVector(n, sizeof(float), y_H, incy, y_D, incy);

// y_D := alpha*x_D + y_D cublasSaxpy(n, alpha, x_D, &incx, y_D, &incy); cublasGetVector(n, sizeof(float), y_D, incy, y_H, incy); cublasFree(x_D); cublasFree(y_D); cublasShutdown(); }

Настольная вычислительная суперкомпьютерная система Tesla D870 = +

2 GPU Tesla C (2x16) мультипроцессора 256 (2x128) потоковых процессорных ядер Рабочая частота ядер – 1.35 ГГц Пропускная способность шины памяти (пиковая) – 76.8 Гб/с Пиковая производительность – 1 Tflop/s Максимальная потребляемая мощность – 520 Вт

Физическое моделирование Вычислительная биология (задачи молекулярной динамики) Вычислительная химия Томография Финансовая математика Компьютерное зрение …

Умножение квадратных матриц Аппаратное обеспечение NVidia Tesla C870 Intel Core2 Quad Q6600 Используемое ПО: Windows XP 32bit; Microsoft Visual Studio 2005 CUBLAS 2.0 (включен в состав CUDA Toolkit 2.0) (для Tesla C870) Intel MKL (для Core2 Quad 6600) Замечание: при измерении общего времени работы алгоритма на видеокарте учитывается время загрузки данных на карту и получения данных с карты

Nvidia CUDA Programming Guide Многочисленные курсы по CUDA: (на русском языке)

?