Программирование для GPU с использованием NVidia CUDA.

1. Программирование для GPU с использованием NVidia CUDA. Половинкин А.Н.

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

3. Почему GPU?!

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

5. Архитектура GPU

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

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

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

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

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

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

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

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

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

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

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

17. Дополнительные типы данных • [u]char[1..4] • [u]int[1..4] • [u]long[1..4] • float[1..4] • double2

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

19. Управление устройствами • перечисление устройств: • cudaError_tcudaGetDeviceCount(int* count) – возвращает число доступных устройств; • cudaError_tcudaGetDevice(int* dev) – возвращает используемое устройство • cudaError_tcudaGetDeviceProperties(structcudaDeviceProp* prop, int dev) – возвращает структуру, содержащую свойства устройства • выбор устройства: • cudaError_tcudaChooseDevice(int* dev, const structcudaDeviceProp* prop) – устанавливает устройство, на котором выполняется device код, в наибольшей степени соответствующее конфигурации • cudaError_tcudaSetDevice(int dev) – устанавливает устройство, на котором выполняется device код; Замечание: Nvidia Tesla D870 представляется в виде двух устройств

20. Управление памятью • выделение и освобождение памяти на устройстве: • cudaError_tcudaMalloc(void** devPtr, size_t count) – выделяет память на устройстве и возвращает указатель на нее • cudaError_tcudaFree(void* devPtr) – освобождает память на устройстве • копирование данных между хостом и устройством: • cudaError_tcudaMemcpy(void* dst, const void* src, size_t count, enumcudaMemcpyKind kind) – копирует данные между хостом и устройством

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

22. Вызов ядер (2) __global__ void KernelFunc() …. dim3 DimGrid(100, 50); dim3 DimBlock (8, 8, 8); size_tSMSize = 64; KernelFunc<<<DimGrid, DimBlock, SMSize>>>();

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

24. Пример: сложение 2 векторов (1) #include <stdio.h> #include <stdlib.h> #include <cutil.h> 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) { inti = threadIdx.x; C[i] = A[i] + B[i]; }

25. Пример: сложение 2 векторов (2) intmain(intargc, char **argv) { float *h_A, *h_B, *h_C; float *d_A, *d_B, *d_C; inti; 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); }

26. Пример: сложение 2 векторов (3) // инициализация 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<<<1, N>>>(d_A, d_B, d_C); // копирование вектора, содержащего сумму A и B с GPU на хост CUDA_SAFE_CALL( cudaMemcpy(h_C, d_C, DATA_SZ, cudaMemcpyDeviceToHost) );

27. Пример: сложение 2 векторов (4) // освобождение памяти на 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); }

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

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

30. Производительность арифметических инструкций (2) • Целочисленное деление и взятие остатка по модулю следует заменять битовыми операциями везде, где это возможно: • если 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 версии арифметических функций

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

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

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

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

35. Доступ к глобальной памяти (3) • Структуры размера больше 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; }; 2 128-bit load instructions 5 32-bit load instructions

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

37. Доступ к глобальной памяти (5)

38. Доступ к глобальной памяти (6)

39. Доступ к глобальной памяти (7)

40. Доступ к shared памяти (1) ... ... ... ... ... ... ... ... ... ... ... ... ... ... ... ... • shared memory состоит из блоков памяти равного размера, доступ к которым может быть осуществлен одновременно, - банков памяти • банки в shared memory организованы таким образом, что последовательно идущие 32-bit words относятся к последовательно идущим банкам памяти Bank 8 Bank 5 Bank 4 Bank 2 Bank 15 Bank 7 Bank 1 Bank 6 Bank 10 Bank 11 Bank 0 Bank 12 Bank 13 Bank 14 Bank 9 Bank 3 Address 8 Address 2 Address 9 Address 6 Address 4 Address 3 Address 11 Address 10 Address 12 Address 5 Address 13 Address 14 Address 0 Address 15 Address 7 Address 1 Address 25 Address 21 Address 28 Address 22 Address 24 Address 17 Address 29 Address 20 Address 30 Address 18 Address 23 Address 31 Address 26 Address 16 Address 19 Address 27

41. Доступ к shared memory (2) __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), тогда для того, чтобы избежать bankconflicts, необходимо, чтобы d = 1

42. Доступ к shared memory (3)

43. Доступ к shared memory (4)

44. Доступ к shared memory (5)

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

46. Библиотеки CUBLAS и CUFFT (2) Имя любой процедуры BLAS имеет следующую структуру: • <character code><name><mod>() • <character code>:символ, описывающий тип данных, с которым работает процедура.

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

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

49. Библиотеки CUBLAS и CUFFT (5) #include <stdio.h> #include <mkl.h> 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 }

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