Иерархия памяти CUDA. Глобальная память. Параллельные решения задач умножения матриц и решения СЛАУ. zЛекторы: yБоресков А.В. (ВМиК МГУ)Боресков А.В. (ВМиК. - презентация

1 Иерархия памяти CUDA. Глобальная память. Параллельные решения задач умножения матриц и решения СЛАУ. z Лекторы: y Боресков А.В. (ВМиК МГУ)Боресков А.В. (ВМиК МГУ) y Харламов А.А. (NVidia)Харламов А.А. (NVidia)

2 Типы памяти в 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)

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

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

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

6 Умножение матриц. Простейшая реализация. #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; }

7 Умножение матриц. Простейшая реализация. 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 cudaMemcpy ( adev, a, numBytes, cudaMemcpyHostToDevice ); // from CPU to DRAM cudaMemcpy ( bdev, b, numBytes, cudaMemcpyHostToDevice ); // from CPU to DRAM matMult >> ( adev, bdev, N, cdev ); cudaThreadSynchronize(); cudaMemcpy ( c, cdev, numBytes, cudaMemcpyDeviceToHost ); cudaFree ( adev ); cudaFree ( bdev ); cudaFree ( cdev );

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

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

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

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

12 Device 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

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

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

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

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

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

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

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

20 Использование отдельных массивов 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

21 Решение системы линейных алгебраических уравнений z Традиционные методы ориентированы на последовательное вычисление элементов и нам не подходят z Есть еще итеративные методы Ax=f, A – матрица размера N*N, f – вектор размера N

22 Итеративные методы z Эффективны когда z Матрица А сильна разрежена z Параллельные вычисления zВ обоих случаях цена (по времени) одной итерации O(N)

23 Сходимость z Если есть сходимость, то только к решению системы z Записав уравнения для погрешности получаем достаточное условие сходимости z За счет выбора достаточно малого значения параметра получаем сходимость

24 Код на CUDA // // one iteration // __global__ void kernel ( float * a, float * f, float alpha, float * x0, float * x1, int n ) { int idx = blockIdx.x * blockDim.x + threadId.x; int ia = n * idx; float sum = 0.0f; for ( int I = 0; i < n; i++ ) sum += a [ia + I] * x0 [I]; x1 [idx] = x0 [idx] + alpha * (sum – f [idx] ); }

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

26 Вопросы