Что такое findslide.org?

FindSlide.org - это сайт презентаций, докладов, шаблонов в формате PowerPoint.


Для правообладателей

Обратная связь

Email: Нажмите что бы посмотреть 

Яндекс.Метрика

Презентация на тему Иерархия памяти CUDA

Содержание

Типы памяти в CUDA
Иерархия памяти CUDA. Глобальная память. Параллельные решения задач умножения матриц и решения Типы памяти в CUDA Типы памяти в CUDAСамая быстрая – shared (on-chip)Самая медленная – глобальная (DRAM)Для Работа с памятью в CUDAОснова оптимизации – оптимизация работы с памятьюМаксимальное использование Умножение матрицПроизведение двух квадратных матриц A и B размера N*N, N кратно Умножение матриц. Простейшая реализация.#define BLOCK_SIZE 16__global__ void matMult ( float * a, Умножение матриц. Простейшая реализация.int     numBytes = N * Простейшая реализация.На каждый элемент 2*N арифметических операций2*N обращений к глобальной памятиMemory bound Оптимизация работы с глобальной памятью.Обращения идут через 32/64/128-битовые словаПри обращении к t[i]sizeof( Использование выравнивания.struct vec3{ float x, y, z;};struct __align__(16) vec3{ float x, y, 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 отдельных транзаций1.2/1.3 Объединение (coalescing)Можно добиться заметного увеличения скорости работы с памятьюЛучше использовать не массив Использование отдельных массивовstruct vec3{ float x, y, z;};vec3 * a;float x = Решение системы линейных алгебраических уравненийТрадиционные методы ориентированы на последовательное вычисление элементов и Итеративные методы Эффективны когдаМатрица А сильна разрежена Параллельные вычисленияВ обоих случаях цена СходимостьЕсли есть сходимость, то только к решению системыЗаписав уравнения для погрешности получаем Код на CUDA//// one iteration//__global__ void kernel ( float * a, float Ресуры нашего курса	CUDA.CS.MSU.SUМесто для вопросов и дискуссийМесто для материалов нашего курсаМесто для Вопросы
Слайды презентации

Слайд 2 Типы памяти в CUDA

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

Слайд 3 Типы памяти в CUDA
Самая быстрая – shared (on-chip)
Самая

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

медленная – глобальная (DRAM)
Для ряда случаев можно использовать кэшируемую

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

Слайд 4 Работа с памятью в CUDA
Основа оптимизации – оптимизация

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

работы с памятью
Максимальное использование shared-памяти
Использование специальных паттернов доступа к

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

Слайд 5 Умножение матриц
Произведение двух квадратных матриц A и B

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

размера N*N, N кратно 16
Матрицы расположены в глобальной памяти
По

одной нити на каждый элемент произведения
2D блок – 16*16
2D grid


Слайд 6 Умножение матриц. Простейшая реализация.
#define BLOCK_SIZE 16

__global__ void matMult

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

( 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

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

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 Простейшая реализация.
На каждый элемент
2*N арифметических операций
2*N обращений

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

к глобальной памяти
Memory bound (тормозит именно доступ к памяти)


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

Оптимизация работы с глобальной памятью.Обращения идут через 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)

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

vec3
{
float x, y, z;
};
Размер равен 12 байт
Элементы массива

не будут выровнены в памяти

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


Слайд 11 Device Compute Capability
Compute Caps. – доступная версия CUDA
Разные

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

возможности 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

Device Compute CapabilityRTM Appendix A.1 CUDA Programming Guide

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

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

запросов к глобальной памяти в один блок (транзакцию)
Независимо происходит

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

Слайд 14 Объединение (coalescing) для GPU с CC 1.0/1.1
Нити обращаются

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

к
32-битовым словам, давая 64-байтовый блок
64-битовым словам, давая

128-байтовый блок
Все 16 слов лежат в пределах блока
k-ая нить half-warp’а обращается к k-му слову блока


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

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

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

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

Слайд 17 Объединение (coalescing) для GPU с CC 1.2/1.3
Нити обращаются

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

к
8-битовым словам, дающим один 32-байтовы сегмент
16-битовым словам, дающим один

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

Слайд 18 Объединение (coalescing)
Если хотя бы одно условие не выполнено
1.0/1.1

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

– 16 отдельных транзаций
1.2/1.3 – объединяет их в блоки

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

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

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

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

компонент – это позволяет использовать coalescing

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

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

* 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 (

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

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.CS.MSU.SUМесто для вопросов и дискуссийМесто для материалов нашего курсаМесто

материалов нашего курса
Место для ваших статей!
Если вы нашли какой-то

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


  • Имя файла: ierarhiya-pamyati-cuda.pptx
  • Количество просмотров: 141
  • Количество скачиваний: 1
- Предыдущая Гвинея
Следующая - История чая