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

Сентябрь 1, 2022

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

Содержание

2. Типы памяти в CUDA
3. Типы памяти в CUDA Самая быстрая – shared (on-chip) Самая медленная – глобальная (DRAM) Для ряда
4. Работа с памятью в CUDA Основа оптимизации – оптимизация работы с памятью Максимальное использование shared-памяти Использование
5. Умножение матриц Произведение двух квадратных матриц A и B размера N*N, N кратно 16 Матрицы расположены
6. Умножение матриц. Простейшая реализация. #define BLOCK_SIZE 16 __global__ void matMult ( float * a, float *
7. Умножение матриц. Простейшая реализация. int numBytes = N * N * sizeof ( float ); float
8. Простейшая реализация. На каждый элемент 2*N арифметических операций 2*N обращений к глобальной памяти Memory bound (тормозит
9. Оптимизация работы с глобальной памятью. Обращения идут через 32/64/128-битовые слова При обращении к t[i] sizeof( t
10. Использование выравнивания. struct vec3 { float x, y, z; }; struct __align__(16) vec3 { float x,
11. Device Compute Capability Compute Caps. – доступная версия CUDA Разные возможности HW Пример: В 1.1 добавлены
12. Device Compute Capability RTM Appendix A.1 CUDA Programming Guide
13. Объединение запросов к глобальной памяти. GPU умеет объединять рад запросов к глобальной памяти в один блок
14. Объединение (coalescing) для GPU с CC 1.0/1.1 Нити обращаются к 32-битовым словам, давая 64-байтовый блок 64-битовым
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 Нити обращаются к 8-битовым словам, дающим один 32-байтовы сегмент
18. Объединение (coalescing) Если хотя бы одно условие не выполнено 1.0/1.1 – 16 отдельных транзаций 1.2/1.3 –
19. Объединение (coalescing) Можно добиться заметного увеличения скорости работы с памятью Лучше использовать не массив структур, а
20. Использование отдельных массивов struct vec3 { float x, y, z; }; vec3 * a; float x
21. Решение системы линейных алгебраических уравнений Традиционные методы ориентированы на последовательное вычисление элементов и нам не подходят
22. Итеративные методы Эффективны когда Матрица А сильна разрежена Параллельные вычисления В обоих случаях цена (по времени)
23. Сходимость Если есть сходимость, то только к решению системы Записав уравнения для погрешности получаем достаточное условие
24. Код на CUDA // // one iteration // __global__ void kernel ( float * a, float
25. Ресуры нашего курса CUDA.CS.MSU.SU Место для вопросов и дискуссий Место для материалов нашего курса Место для
27. Скачать презентацию

Слайд 2

Типы памяти в CUDA

Слайд 3

Типы памяти в CUDA
Самая быстрая – shared (on-chip)
Самая медленная – глобальная

(DRAM)
Для ряда случаев можно использовать кэшируемую константную и текстурную память
Доступ к памяти в CUDA идет отдельно для каждой половины warp’а (half-warp)

Слайд 4

Работа с памятью в CUDA
Основа оптимизации – оптимизация работы с памятью
Максимальное

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

Слайд 5

Умножение матриц
Произведение двух квадратных матриц A и B размера N*N, N

кратно 16
Матрицы расположены в глобальной памяти
По одной нити на каждый элемент произведения
2D блок – 16*16
2D 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

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

bound (тормозит именно доступ к памяти)

Слайд 9

Оптимизация работы с глобальной памятью.
Обращения идут через 32/64/128-битовые слова
При обращении к

t[i]
sizeof( t [0] ) равен 4/8/16 байтам
t [i] выровнен по sizeof ( t [0] )
Вся выделяемая память всегда выровнена по 256 байт

Слайд 10

Использование выравнивания.
struct vec3
{
float x, y, z;
};
struct align(16) vec3
{
float x,

y, z;
};

Размер равен 12 байт
Элементы массива не будут выровнены в памяти

Размер равен 16 байт
Элементы массива всегда будут выровнены в памяти

Слайд 11

Device Compute Capability
Compute Caps. – доступная версия CUDA
Разные возможности HW
Пример:
В 1.1

добавлены атомарные операции в global memory
В 1.2 добавлены атомарные операции в shared memory
В 1.3 добавлены вычисления в double
Узнать доступный Compute Caps. можно через cudaGetDeviceProperties()
См. CUDAHelloWorld
Сегодня Compute Caps:
Влияет на правила работы с глобальной памятью

Слайд 12

Device Compute Capability
RTM Appendix A.1 CUDA Programming Guide

Слайд 13

Объединение запросов к глобальной памяти.
GPU умеет объединять рад запросов к глобальной

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

Слайд 14

Объединение (coalescing) для GPU с CC 1.0/1.1
Нити обращаются к
32-битовым словам,

давая 64-байтовый блок
64-битовым словам, давая 128-байтовый блок
Все 16 слов лежат в пределах блока
k-ая нить 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
Нити обращаются к
8-битовым словам, дающим

один 32-байтовы сегмент
16-битовым словам, дающим один 64-байтовый сегмент
32-битовым словам, дающим один 128-байтовый сегмент
Получающийся сегмент выровнен по своему размеру

Слайд 18

Объединение (coalescing)
Если хотя бы одно условие не выполнено
1.0/1.1 – 16 отдельных

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

Слайд 19

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

массив структур, а набор массивов отдельных компонент – это позволяет использовать 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

Решение системы линейных алгебраических уравнений
Традиционные методы ориентированы на последовательное вычисление элементов

и нам не подходят
Есть еще итеративные методы

Ax=f,
A – матрица размера N*N,
f – вектор размера N

Слайд 22

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

цена (по времени) одной итерации O(N)

Слайд 23

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

получаем достаточное условие сходимости
За счет выбора достаточно малого значения параметра получаем сходимость

Слайд 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

Ресуры нашего курса
CUDA.CS.MSU.SU
Место для вопросов и дискуссий
Место для материалов нашего курса
Место

для ваших статей!
Если вы нашли какой-то интересный подход!
Или исследовали производительность разных подходов и знаете, какой из них самый быстрый!
Или знаете способы сделать работу с CUDA проще!
www.steps3d.narod.ru
www.nvidia.ru

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

Содержание

Типы памяти в CUDA

Типы памяти в CUDAСамая быстрая – shared (on-chip)Самая медленная – глобальная

Работа с памятью в CUDAОснова оптимизации – оптимизация работы с памятьюМаксимальное

Умножение матрицПроизведение двух квадратных матриц A и B размера N*N, N

Умножение матриц. Простейшая реализация.#define BLOCK_SIZE 16__global__ void matMult ( float *

Умножение матриц. Простейшая реализация.int numBytes = N * N * sizeof

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

Оптимизация работы с глобальной памятью.Обращения идут через 32/64/128-битовые словаПри обращении к

Использование выравнивания.struct vec3{ float x, y, z;};struct __align__(16) vec3{ float x,

Device Compute CapabilityCompute Caps. – доступная версия CUDAРазные возможности HWПример:В 1.1

Device Compute CapabilityRTM Appendix A.1 CUDA Programming Guide

Объединение запросов к глобальной памяти.GPU умеет объединять рад запросов к глобальной

Объединение (coalescing) для GPU с CC 1.0/1.1Нити обращаются к 32-битовым словам,

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

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

Объединение (coalescing) для GPU с CC 1.2/1.3Нити обращаются к8-битовым словам, дающим

Объединение (coalescing)Если хотя бы одно условие не выполнено1.0/1.1 – 16 отдельных

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

Использование отдельных массивовstruct vec3{ float x, y, z;};vec3 * a;float x

Решение системы линейных алгебраических уравненийТрадиционные методы ориентированы на последовательное вычисление элементов

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

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

Код на CUDA//// one iteration//__global__ void kernel ( float * a,

Ресуры нашего курса CUDA.CS.MSU.SUМесто для вопросов и дискуссийМесто для материалов нашего курсаМесто

Похожие презентации