Лекторы: Боресков А.В. (ВМиК МГУ) Харламов А.А. (NVidia) CUDA Stream. Иерархия памяти в CUDA. Глобальная память. - презентация

1 Лекторы: Боресков А.В. (ВМиК МГУ) Харламов А.А. (NVidia) CUDA Stream. Иерархия памяти в CUDA. Глобальная память.

2 2011 Примеры многоядерных систем На первой лекции мы рассмотрели – Intel Core 2 Duo – SMP – Cell – BlueGene/L – G80 / Tesla / Fermi

3 2011 Подход CUDA Исходная задача разбивается на подзадачи, которые можно решать независимо друг от друга. Каждая из этих подзадач решается набором взаимодействующих между собой нитей

4 2011 SIMT (Single Instruction, Multiple Threads) z Параллельно на каждом SM выполняется большое число отдельных нитей (threads) z Нити подряд разбиваются на warpы (по 32 нити) и SM управляет выполнением warpов z Нити в пределах одного warpа выполняются физически параллельно z Большое число warpов покрывает латентность

5 2011 Технические детали RTM CUDA Programming Guide Run CUDAHelloWorld – Печатает аппаратно зависимые параметры Размер shared памяти Кол-во SM Размер warpа Кол-во регистров на SM т.д.

6 2011 Программная модель CUDA z Код состоит как из последовательных, так и из параллельных частей z Последовательные части кода выполняются на CPU z Массивно-параллельные части кода выполняются на GPU как ядра

7 2011 Программная модель CUDA GPU (device) это вычислительное устройство, которое: – Является сопроцессором к CPU (host) – Имеет собственную память (DRAM) – Выполняет одновременно очень много нитей

8 2011 Программная модель CUDA Последовательные части кода выполняются на CPU Массивно-параллельные части кода выполняются на GPU как ядра Отличия нитей между CPU и GPU – Нити на GPU очень «легкие» – HW планировщик задач – Для полноценной загрузки GPU нужны тысячи нитей Для покрытия латентностей операций чтения / записи Для покрытия латентностей sfu инструкций

9 2011 Программная модель CUDA Параллельная часть кода выполняется как большое количество нитей (threads) Нити группируются в блоки (blocks) фиксированного размера Блоки объединяются в сеть блоков (grid) Ядро выполняется на сетке из блоков Каждая нить и блок имеют свой уникальный идентификатор

10 2011 Программная модель CUDA Десятки тысяч потоков for ( int ix = 0; ix < nx; ix++ ) { pData[ix] = f(ix); } for ( int ix = 0; ix < nx; ix++ ) for ( int iy = 0; iy < ny; iy++ ) { pData[ix + iy * nx] = f(ix) * g(iy); } for ( int ix = 0; ix < nx; ix++ ) for ( int iy = 0; iy < ny; iy++ ) for ( int iz = 0; iz < nz; iz++ ) { pData[ix + (iy + iz * ny) * nx] = f(ix) * g(iy) * h(iz); }

11 2011 Программная модель CUDA Потоки в CUDA объединяются в блоки: – Возможна 1D, 2D, 3D топология блока Общее кол-во потоков в блоке ограничено В текущем HW это 512 потоков

12 2011 Программная модель CUDA Потоки в блоке могут разделять ресурсы со своими соседями float data[N]; for ( int ix = 0; ix < nx; ix++ ) data[ix] = f(ix, data[ix / n]);

13 2011 Программная модель CUDA Блоки могут использовать shared память – Т.к. блок целиком выполняется на одном SM – Объем shared памяти ограничен и зависит от HW Внутри Блока потоки могут синхронизоваться – Т.к. блок целиком выполняется на одном SM

14 2011 Программная модель CUDA Блоки потоков объединяются в сетку (grid) потоков – Возможна 1D, 2D топология сетки блоков потоков

15 2011 Синтаксис CUDA CUDA – это расширение языка C/C++ – [+] спецификаторы для функций и переменных – [+] новые встроенные типы – [+] встроенные переменные (внутри ядра) – [+] директива для запуска ядра из C кода Как скомпилировать CUDA код – [+] nvcc компилятор – [+].cu расширение файла

16 2011 Синтаксис CUDA Спецификаторы СпецификаторВыполняется наМожет вызываться из __device__device __global__devicehost __host__host zСпецификатор функций zСпецификатор переменных СпецификаторНаходитсяДоступнаВид доступа __device__device R __constant__devicedevice / hostR / W __shared__deviceblock RW / __syncthreads()

17 2011 Расширения языка C z Спецификатор __global__ соответствует ядру z Может возвращать только void z Спецификаторы __host__ и __device__ могут использоваться одновременно z Компилятор сам создаст версии для CPU и GPU z Спецификаторы __global__ и __host__ не могут быть использованы одновременно

18 2011 Расширения языка C Ограничения на функции, выполняемые на GPU: z Нельзя брать адрес (за исключением __global__) z Не поддерживается рекурсия z Не поддерживаются static-переменные внутри функции z Не поддерживается переменное число входных аргументов

19 2011 Расширения языка C Ограничения на спецификаторы переменных: z Нельзя применять к полям структуры или union z Не могут быть extern z Запись в __constant__ может выполнять только CPU через специальные функции z __shared__ - переменные не могут инициализироваться при объявлении

20 2011 Расширения языка C Новые типы данных: z 1/2/3/4-мерные вектора из базовых типов z (u)char, (u)int, (u)short, (u)long, longlong z float, double z dim3 – uint3 с нормальным конструкторов, позволяющим задавать не все компоненты z Не заданные инициализируются единицей

21 2011 Расширения языка С int2 a = make_int2 ( 1, 7 ); float4 b = make_float4 ( a.x, a.y, 1.0f, 7 ); float2 x = make_float2 ( b.z, b.w ); dim3 grid = dim3 ( 10 ); dim3 blocks = dim3 ( 16, 16 ); Для векторов не определены покомпонентные операции Для double и longlong возможны только вектора размера 1 и 2.

22 2011 Синтаксис CUDA Встроенные переменные Сравним CPU код vs CUDA kernel: __global__ void incKernel ( float * data ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; data [idx] = data [idx] + 1.0f; } float * data; for ( int i = 0; i < n; i++ ) { data [x] = data[i] + 1.0f; } Пусть nx = 2048 Пусть в блоке 256 потоков кол-во блоков = 2048 / 256 = 8 [ ][ == 256][ ]

23 2011 Синтаксис CUDA Встроенные переменные В любом CUDA kernele доступны: – dim3 gridDim; – uint3 blockIdx; – dim3 blockDim; – uint3 threadIdx; – int warpSize; dim3 – встроенный тип, который используется для задания размеров kernelа По сути – это uint3.

24 2011 Синтаксис CUDA Директивы запуска ядра Как запустить ядро с общим кол-во тредов равным nx? incKernel ( data ); dim3 threads ( 256 ); dim3 blocks ( nx / 256 ); float * data; >> угловые скобки, внутри которых задаются параметры запуска ядра: Кол-во блоке в сетке Кол-во потоков в блоке … Неявно предпологаем, что nx кратно 256

25 2011 Расширения языка С Общий вид команды для запуска ядра incKernel >> ( data ); z bl – число блоков в сетке z th – число нитей в сетке z ns – количество дополнительной shared- памяти, выделяемое блоку z st – поток, в котором нужно запустить ядро

26 2011 Как скомпилировать CUDA код NVCC – компилятор для CUDA – Основными опциями команды nvcc являются: – -deviceemu - компиляция в режиме эмуляции, весь код будет выполняться в многонитевом режиме на CPU и можно использовать обычный отладчик (хотя не все ошибки могут проявится в таком режиме) – --use_fast_math - заменить все вызовы стандартных математических функций на их быстрые (но менее точные) аналоги – -o - задать имя выходного файла CUDA файлы обычно носят расширение.cu

27 2011 Основы CUDA host API Два API z Низкоуровневый driver API (cu*) z Высокоуровневый runtime API (cuda*) z Реализован через driver API z Не требуют явной инициализации z Все функции возвращают значение типа cudaError_t z cudaSuccess в случае успеха

28 2011 Основы CUDA API Многие функции API асинхронны: z Запуск ядра z Копирование при помощи функций *Async z Копирование device device z Инициализация памяти

29 2011 Основы CUDA API char * cudaGetErrorString ( cudaError_t ); cudaError_t cudaGetLastError (); cudaError_t cudaThreadSynchronize (); cudaError_t cudaEventCreate ( cudaEvent_t * ); cudaError_t cudaEventRecord ( cudaEvent_t * ); cudaStream_t ); cudaError_t cudaEventQuery ( cudaEvent_t ); cudaError_t cudaEventSynchronize ( cudaEvent_t ); cudaError_t cudeEventElapsedTime ( float * time, cudaEvent_t start, cudaEvent_t stop ); cudaError_t cudaEventDestroy ( cudaEvent_t ); cudaError_t cudaGetDeviceCount ( int * ); cudaError_t cudaGetDevicePropertis ( cudaDeviceProp * props, int deviceNo );

30 2011 CUDA Compute Capability Возможности GPU обозначаются при помощи Compute Capability, например 1.1 z Старшая цифра соответствует архитектуре z Младшая – небольшим архитектурным изменениям z Можно получить из полей major и minor структуры cudaDeviceProp

31 2011 Получение информации о GPU int main ( int argc, char * argv [] ) { intdeviceCount; cudaDevicePropdevProp; cudaGetDeviceCount ( &deviceCount ); printf ( "Found %d devices\n", deviceCount ); for ( int device = 0; device < deviceCount; device++ ) { cudaGetDeviceProperties ( &devProp, device ); printf ( "Device %d\n", device ); printf ( "Compute capability : %d.%d\n", devProp.major, devProp.minor ); printf ( "Name : %s\n", devProp.name ); printf ( "Total Global Memory : %d\n", devProp.totalGlobalMem ); printf ( "Shared memory per block: %d\n", devProp.sharedMemPerBlock ); printf ( "Registers per block : %d\n", devProp.regsPerBlock ); printf ( "Warp size : %d\n", devProp.warpSize ); printf ( "Max threads per block : %d\n", devProp.maxThreadsPerBlock ); printf ( "Total constant memory : %d\n", devProp.totalConstMem ); } return 0; }

32 2011 Compute Capability GPUCompute Capability Tesla S GeForce GTX GeForce 9800 GX2 1.1 GeForce 9800 GTX 1.1 GeForce 8800 GT 1.1 GeForce 8800 GTX 1.0 RTM Appendix A.1 CUDA Programming Guide

33 2011 Compute Capability z Compute Caps. – доступная версия CUDA y Разные возможности HW y Пример: x В 1.1 добавлены атомарные операции в global memory x В 1.2 добавлены атомарные операции в shared memory x В 1.3 добавлены вычисления в double z Узнать доступный Compute Caps. можно через cudaGetDeviceProperties() y См. CUDAHelloWorld z Сегодня Compute Caps: y Влияет на правила работы с глобальной памятью

34 2011 Компиляция программ z Используем утилиту make/nmake, явно вызывающую nvcc z Используем MS Visual Studio z Подключаем cuda.rules z Используем CUDA Wizard ( ard)

35 2011 Типы памяти в CUDA Тип памятиДоступУровень выделения Скорость работы Регистры R/WPer-threadВысокая(on-chip) Локальная R/WPer- thread Низкая (DRAM) Shared R/WPer-blockВысокая(on-chip) Глобальная R/WPer-gridНизкая (DRAM) Constant R/OPer-gridВысокая(L1 cache) Texture R/OPer-gridВысокая(L1 cache)

36 2011 Типы памяти в CUDA z Самая быстрая – shared (on-chip) и регистры z Самая медленная – глобальная (DRAM) z Для ряда случаев можно использовать кэшируемую константную и текстурную память z Доступ к памяти в CUDA идет отдельно для каждой половины warpа (half-warp)

37 2011 Работа с памятью в CUDA Основа оптимизации – оптимизация работы с памятью: z Максимальное использование shared- памяти z Использование специальных паттернов доступа к памяти, гарантирующих эффективный доступ y Паттерны работают независимо в пределах каждого half-warpа

38 2011 Работа с глобальной памятью в CUDA float * devPtr; // pointer device memory // allocate device memory cudaMalloc ( (void **) &devPtr, 256*sizeof ( float ); // copy data from host to device memory cudaMemcpy ( devPtr, hostPtr, 256*sizeof ( float ), cudaMemcpyHostToDevice ); // process data // copy results from device to host cudaMemcpy ( hostPtr, devPtr, 256*sizeof( float ), cudaMemcpyDeviceToHost ); // free device memory cudaFree ( devPtr ); Пример работы с глобальной памятью

39 2011 Работа с глобальной памятью в CUDA cudaError_t cudaMalloc ( void ** devPtr, size_t size ); cudaError_t cudaMallocPitch ( void ** devPtr, size_t * pitch, size_t width, size_t height ); cudaError_t cudaFree ( void * devPtr ); cudaError_t cudaMemcpy ( void * dst, const void * src, size_t count, enum cudaMemcpyKind kind ); cudaError_t cudaMemcpyAsync ( void * dst, const void * src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream ); cudaError_t cudaMemset ( void * devPtr, int value, size_t count ); Функции для работы с глобальной памятью

40 2011 Пример: умножение матриц z Произведение двух квадратных матриц A и B размера N*N, N кратно 16 z Матрицы расположены в глобальной памяти z По одной нити на каждый элемент произведения z 2D блок – 16*16 z 2D grid

41 2011 Умножение матриц. Простейшая реализация. #define BLOCK_SIZE 16 __global__ void matMult ( float * a, float * b, int n, float * c ) { int bx = blockIdx.x; int by = blockIdx.y; int tx = threadIdx.x; int ty = threadIdx.y; float sum = 0.0f; int ia = n * BLOCK_SIZE * by + n * ty; int ib = BLOCK_SIZE * bx + tx; int ic = n * BLOCK_SIZE * by + BLOCK_SIZE * bx; for ( int k = 0; k < n; k++ ) sum += a [ia + k] * b [ib + k*n]; c [ic + n * ty + tx] = sum; }

42 2011 Умножение матриц. Простейшая реализация. int numBytes = N * N * sizeof ( float ); float * adev, * bdev, * cdev ; dim3 threads ( BLOCK_SIZE, BLOCK_SIZE ); dim3 blocks ( N / threads.x, N / threads.y); cudaMalloc ( (void**)&adev, numBytes );// allocate DRAM cudaMalloc ( (void**)&bdev, numBytes ); // allocate DRAM cudaMalloc ( (void**)&cdev, numBytes ); // allocate DRAM // copy from CPU to DRAM cudaMemcpy ( adev, a, numBytes, cudaMemcpyHostToDevice ); cudaMemcpy ( bdev, b, numBytes, cudaMemcpyHostToDevice ); matMult >> ( adev, bdev, N, cdev ); cudaThreadSynchronize(); cudaMemcpy ( c, cdev, numBytes, cudaMemcpyDeviceToHost ); // free GPU memory cudaFree ( adev ); cudaFree ( bdev ); cudaFree ( cdev );

43 2011 Простейшая реализация. z На каждый элемент z 2*N арифметических операций z 2*N обращений к глобальной памяти z Memory bound (тормозит именно доступ к памяти)

44 2011 Используем CUDA Profiler z Легко видно, что основное время (84.15%) ушло на чтение из глобальной памяти z Непосредственно вычисления заняли всего около 10%

45 2011 Оптимизация работы с глобальной памятью. z Обращения идут через 32/64/128- битовые слова z При обращении к t[i] y sizeof(t [0]) равен 4/8/16 байтам y t [i] выровнен по sizeof ( t [0] ) z Вся выделяемая память всегда выровнена по 256 байт

46 2011 Использование выравнивания. struct vec3 { float x, y, z; }; struct __align__(16) vec3 { float x, y, z; }; z Размер равен 12 байт z Элементы массива не будут выровнены в памяти z Размер равен 16 байт z Элементы массива всегда будут выровнены в памяти

47 2011 Объединение запросов к глобальной памяти. z GPU умеет объединять ряд запросов к глобальной памяти в один блок (транзакцию) z Независимо происходит для каждого half-warpа z Длина блока должна быть 32/64/128 байт z Блок должен быть выровнен по своему размеру

48 2011 Объединение (coalescing) для GPU с CC 1.0/1.1 z Нити обращаются к y 32-битовым словам, давая 64-байтовый блок y 64-битовым словам, давая 128-байтовый блок z Все 16 слов лежат в пределах блока z k-ая нить half-warpа обращается к k-му слову блока

49 2011 Объединение (coalescing) для GPU с CC 1.0/1.1 Coalescing

50 2011 Объединение (coalescing) для GPU с CC 1.0/1.1 Not Coalescing

51 2011 Объединение (coalescing) для GPU с CC 1.2/1.3 z Нити обращаются к y 8-битовым словам, дающим один 32- байтовый сегмент y 16-битовым словам, дающим один 64- байтовый сегмент y 32-битовым словам, дающим один 128- байтовый сегмент z Получающийся сегмент выровнен по своему размеру

52 2011 Объединение (coalescing) z Если хотя бы одно условие не выполнено y 1.0/1.1 – 16 отдельных транзакций y 1.2/1.3 – объединяет их в блоки (2,3,…) и для каждого блока проводится отдельная транзакция z Для 1.2/1.3 порядок в котором нити обращаются к словам внутри блока не имеет значения (в отличии от 1.0/1.1)

53 2011 Объединение (coalescing) z Можно добиться заметного увеличения скорости работы с памятью z Лучше использовать не массив структур, а набор массивов отдельных компонент – это позволяет использовать coalescing

54 2011 Использование отдельных массивов struct vec3 { float x, y, z; }; vec3 * a; float x = a [threadIdx.x].x; float y = a [threadIdx.x].y; float z = a [threadIdx.x].z; float * ax, * ay, * az; float x = ax [threadIdx]; float y = ay [threadIdx]; float z = az [threadIdx]; Не можем использовать coalescing при чтении данных Поскольку нити одновременно обращаются к последовательно лежащим словам памяти, то будет происходить coalescing

55 2011 Ресурсы нашего курса CUDA.CS.MSU.SU – Место для вопросов и дискуссий – Место для материалов нашего курса – Место для ваших статей! Если вы нашли какой-то интересный подход! Или исследовали производительность разных подходов и знаете, какой из них самый быстрый! Или знаете способы сделать работу с CUDA проще!

56 2011