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

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

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

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

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

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

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

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

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

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); }

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

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

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

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

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

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

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

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

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

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

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.

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][ ]

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

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

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

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

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

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

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 );

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

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; }

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

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 Влияет на правила работы с глобальной памятью

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

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)

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

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

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 ); Пример работы с глобальной памятью

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 ); Функции для работы с глобальной памятью

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

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; }

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 );

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

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

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

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

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

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

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

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

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

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)

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

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

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

2011