Иерархия памяти CUDA Глобальная память Параллельные решения задач

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

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

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

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

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

Умножение матриц. Простейшая реализация. #define BLOCK_SIZE 16 __global__ void mat. Mult ( float * a, float * b, int n, float * c ) { int bx = block. Idx. x; int by = block. Idx. y; int tx = thread. Idx. x; int ty = thread. Idx. y; float sum = 0. 0 f; 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; }

Умножение матриц. Простейшая реализация. int float dim 3 cuda. Malloc num. Bytes = N * sizeof ( float ); * adev, * bdev, * cdev ; threads ( BLOCK_SIZE, BLOCK_SIZE ); blocks ( N / threads. x, N / threads. y); ( (void**)&adev, num. Bytes ); ( (void**)&bdev, num. Bytes ); ( (void**)&cdev, num. Bytes ); // allocate DRAM cuda. Memcpy ( adev, a, num. Bytes, cuda. Memcpy. Host. To. Device ); // from CPU to DRAM cuda. Memcpy ( bdev, b, num. Bytes, cuda. Memcpy. Host. To. Device ); // from CPU to DRAM mat. Mult<<<blocks, threads>>> ( adev, bdev, N, cdev ); cuda. Thread. Synchronize(); cuda. Memcpy ( c, cdev, num. Bytes, cuda. Memcpy. Device. To. Host ); cuda. Free ( adev ); ( bdev ); ( cdev );

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

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

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

Device 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. можно через cuda. Get. Device. Properties() y См. CUDAHello. World z Сегодня Compute Caps: y Влияет на правила работы с глобальной памятью

Device Compute Capability GPU Tesla S 1070 Compute Capability 1. 3 Ge. Force GTX 260 1. 3 Ge. Force 9800 GX 2 1. 1 Ge. Force 9800 GTX 1. 1 Ge. Force 8800 GTX 1. 0 RTM Appendix A. 1 CUDA Programming Guide

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

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

Объединение (coalescing) для GPU с CC 1. 0/1. 1 Coalescing

Объединение (coalescing) для GPU с CC 1. 0/1. 1 Not Coalescing

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

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

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

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

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

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

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

Код на CUDA // // one iteration // __global__ void kernel ( float * a, float * f, float alpha, float * x 0, float * x 1, int n ) { int idx = block. Idx. x * block. Dim. x + thread. Id. x; int ia = n * idx; float sum = 0. 0 f; for ( int I = 0; i < n; i++ ) sum += a [ia + I] * x 0 [I]; x 1 [idx] = x 0 [idx] + alpha * (sum – f [idx] ); }

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

Вопросы