Разделы презентаций


Powerpoint Templates Основы технологии CUDA. Работа с памятью

Содержание

Сравнение архитектур CPU и GPUПараллельная обработка данных распределяет элементы данных на параллельно обрабатываемых потоках. GPU особенно хорошо подходит для решения проблем, которые могут быть выражены как вычислений данными параллельно - та

Слайды и текст этой презентации

Слайд 1Powerpoint Templates
Основы технологии CUDA.
Работа с памятью

Powerpoint TemplatesОсновы технологии CUDA. Работа с памятью

Слайд 2Сравнение архитектур CPU и GPU
Параллельная обработка данных распределяет элементы данных

на параллельно обрабатываемых потоках. GPU особенно хорошо подходит для решения

проблем, которые могут быть выражены как вычислений данными параллельно - та же программа выполняется на многих элементов данных параллельно - с высокой интенсивностью - арифметическое отношение арифметических операций к операциям с памятью.
Сравнение архитектур CPU и GPUПараллельная обработка данных распределяет элементы данных на параллельно обрабатываемых потоках. GPU особенно хорошо

Слайд 3Вычислительная модель GPU
Двумерная блочная структура

Вычислительная модель GPU Двумерная блочная структура

Слайд 4Структура блоков
Трехмерная структура блоков

Структура блоковТрехмерная структура блоков

Слайд 5Подключаемые библиотеки
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include
#include

Подключаемые библиотеки#include

Слайд 6Оценка затраченного на вычисления времени
cudaEvent_t start, stop;
float gpuTime;
cudaEventCreate(

&start );
cudaEventCreate( &stop );
cudaEventRecord( start, 0 );
...


cudaEventRecord( stop, 0 ) ;
cudaEventSynchronize( stop ) );
cudaEventElapsedTime( &gpuTime, start, stop ) ;
printf("time spent executing by the GPU: %.2f millseconds\n", gpuTime );
cudaEventDestroy( start) ;
cudaEventDestroy( stop) ;

cudaEventCreate() - создание временных меток start, stop ;
cudaEventRecord() - фиксация времени старта ;
cudaEventRecord( ) - фиксация времени завершения.
cudaEventSynchronize() - синхронизация асинхронных процессов;
cudaEventElapsedTime() - вычисление разницы во времени.

Оценка затраченного на вычисления времениcudaEvent_t start, stop; float gpuTime; cudaEventCreate( &start ); cudaEventCreate( &stop ); cudaEventRecord( start,

Слайд 7О компоновке нитей и блоков
#define DGX 8
#define DGY 32


#define DBX 8
#define DBY 8
#define DBZ 8
#define

N (DBX*DBY*DBZ*DGX*DGY)
__global__ void kern( float *a )
{ int bs = blockDim.x*blockDim.y*blockDim.z;
int idx = threadIdx.x + threadIdx.y*blockDim.x
+ threadIdx.z* (blockDim.x*blockDim.y)
+ blockIdx.x*bs + blockIdx.y*bs*gridDim.x ;
if(idx > N-1) return;
a[idx] -= 0.5f;
}
Размер блока 8x8x8, что как раз равно 512 тредов на один блок, размер грида 8х32 блока, таким образом общее количество параллельных процессов 131072=8x8x8x8x32.
При этом, адресация выделенной памяти -- линейная и сложный номер треда пересчитывается в индекс ячейки масива памяти.

О компоновке нитей и блоков#define DGX 8 	#define DGY 32 	#define DBX 8 	#define DBY 8 	#define

Слайд 8Отладка программ
Функции из CUDA runtime API могут возвращать различные коды

ошибок. Можно использовать следующий макрос для отлова ошибок:
#define CUDA_DEBUG
#ifdef CUDA_DEBUG
#define

CUDA_CHECK_ERROR(err)
if (err != cudaSuccess) {
printf("Cuda error: %s\n", cudaGetErrorString(err));
printf("Error in file: %s, line: %i\n", __FILE__, __LINE__);
} \
#else
#define CUDA_CHECK_ERROR(err)
#endif

Если определена переменная среды CUDA_DEBUG, происходит проверка кода ошибки и выводиться информация о файле и строке, где она произошла. Эту переменную можно включить при компиляции под отладку и отключить при компиляции под релиз.

Отладка программФункции из CUDA runtime API могут возвращать различные коды ошибок. Можно использовать следующий макрос для отлова

Слайд 9Типы памяти

Типы памяти

Слайд 10Типы памяти
Регистры.
Локальная память.
Глобальная память.
Разделяемая память.
Константная память.
Текстурная память

Типы памятиРегистры. Локальная память. Глобальная память. Разделяемая память. Константная память.Текстурная память

Слайд 11Организация памяти устройства

Организация памяти устройства

Слайд 12Регистровая память (register)
Является самой быстрой из всех видов. Определить количество регистров

доступных GPU можно с помощью уже функции cudaGetDeviceProperties.
Рассчитать количество

регистров, доступных одной нити GPU - для этого необходимо разделить общее число регистров на произведение количества нитей в блоке и количества блоков в гриде.
Все регистры GPU 32 разрядные.
В CUDA нет явных способов использования регистровой памяти. Это определяет компилятор.
Регистровая память (register)Является самой быстрой из всех видов. Определить количество регистров доступных GPU можно с помощью уже функции

Слайд 13Расчет количества регистров, доступных одной нити GPU
 
 
При вызове функций ядра
myKernelFunc

gridSize, blockSize, sharedMemSize,  cudaStream >>>(float* param1,float * param2),
Вызов функции
cudaGetDeviceProperties

Расчет количества регистров, доступных одной нити GPU  При вызове функций ядраmyKernelFunc>(float* param1,float * param2), Вызов функцииcudaGetDeviceProperties

Слайд 14Локальная память
Локальная память (local memory) может быть использована компилятором при большом количестве

локальных переменных в какой-либо функции. По скоростным характеристикам локальная память

значительно медленнее, чем регистровая. В документации от nVidia рекомендуется использовать локальную память только в самых необходимых случаях.
Явных средств, позволяющих блокировать использование локальной памяти, не предусмотрено, поэтому при падении производительности стоит тщательно проанализировать код и исключить лишние локальные переменные.
Локальная памятьЛокальная память (local memory) может быть использована компилятором при большом количестве локальных переменных в какой-либо функции. По скоростным

Слайд 15Глобальная память
Глобальная память (global memory) – самый медленный тип памяти, из

доступных GPU. Глобальные переменные можно выделить с помощью спецификатора __global__,

а так же динамически, с помощью функций из семейства cudMallocXXX. Глобальная память в основном служит для хранения больших объемов данных, поступивших на device с host’а, данное перемещение осуществляется с использованием функций cudaMemcpyXXX.
В алгоритмах, требующих высокой производительности, количество операций с глобальной памятью необходимо свести к минимуму.

Глобальная памятьГлобальная память (global memory) – самый медленный тип памяти, из доступных GPU. Глобальные переменные можно выделить с

Слайд 16Разделяемая память
Разделяемая память (shared memory) относиться к быстрому типу памяти. Разделяемую

память рекомендуется использовать для минимизации обращение к глобальной памяти, а

так же для хранения локальных переменных функций.
Адресация разделяемой памяти между нитями потока одинакова в пределах одного блока, что может быть использовано для обмена данными между потоками в пределах одного блока.
Для размещения данных в разделяемой памяти используется спецификатор __shared__.
Разделяемая памятьРазделяемая память (shared memory) относиться к быстрому типу памяти. Разделяемую память рекомендуется использовать для минимизации обращение к

Слайд 17Константная память
Константная память (constant memory) является достаточно быстрой из доступных GPU.

Отличительной особенностью константной памяти является возможность записи данных с хоста,

но при этом в пределах GPU возможно лишь чтение из этой памяти, что и обуславливает её название.
Для размещения данных в константной памяти предусмотрен спецификатор __constant__. Если необходимо использовать массив в константной памяти, то его размер необходимо указать заранее, так как динамическое выделение в отличие от глобальной памяти в константной не поддерживается.
Для записи с хоста в константную память используется функция cudaMemcpyToSymbol, и для копирования с device’а на хост cudaMemcpyFromSymbol, как видно этот подход несколько отличается от подхода при работе с глобальной памятью.
Константная памятьКонстантная память (constant memory) является достаточно быстрой из доступных GPU. Отличительной особенностью константной памяти является возможность записи

Слайд 18Текстурная память
Предназначена главным образом для работы с текстурами. Она оптимизирована

под выборку 2D данных и имеет следующие возможности:
быстрая выборка значений

фиксированного размера из одномерного или двухмерного массива;
нормализованная адресация числами типа float в интервале [0,1).
аппаратная линейная или билинейная интерполяция соседних значений в случае нормализованной адресации;
аппаратная обработка выхода за границу массива с использованием двух режимов: clamp и wrap.
Размер текстурной памяти ограничивается только максимальным размером памяти, которую может выделить устройство. Но так же из текстурной памяти можно читать данные только встроенных в nvcc типов, имеющих размер 1, 2, 4, 8 или 16 байт, и только с помощью специальных функций — tex1D, tex2D или tex1Dfetch, tex2Dfetch. Другими словами, нельзя сделать указатель на текстурную память и переименовать его произвольным образом.

Текстурная памятьПредназначена главным образом для работы с текстурами. Она оптимизирована под выборку 2D данных и имеет следующие

Слайд 19Пример использования различных типов памяти
При операции транспонирования матрицы

Пример использования различных типов памятиПри операции транспонирования матрицы

Слайд 20Транспонирование матрицы на CPU
// inputMatrix - указатель на исходную матрицу


// outputMatrix - указатель на матрицу результат
// width - ширина

исходной матрицы (она же высота матрицы-результата)
// height - высота исходной матрицы (она же ширина матрицы-результата)
__host__ void transposeMatrixCPU(float* inputMatrix, float* outputMatrix, int width, int height) {   for (int y = 0; y < height; y++)   {     for (int x = 0; x < width; x++)     {      outputMatrix[x * height + y] = inputMatrix[y * width + x];     }   } }
Транспонирование матрицы на CPU// inputMatrix - указатель на исходную матрицу // outputMatrix - указатель на матрицу результат//

Слайд 21Использование только глобальной памяти.
__global__ void transposeMatrixSlow(float* inputMatrix, float* outputMatrix, int

width, int height)
{
int xIndex = blockDim.x * blockIdx.x +

threadIdx.x;
int yIndex = blockDim.y * blockIdx.y + threadIdx.y;
if ((xIndex < width) && (yIndex < height))
{
int inputIdx = xIndex + width * yIndex;
int outputIdx = yIndex + height * xIndex;
outputMatrix[outputIdx] = inputMatrix[inputIdx];
}
}

Использование только глобальной памяти.__global__ void transposeMatrixSlow(float* inputMatrix, float* outputMatrix, int width, int height){ int xIndex = blockDim.x

Слайд 22Использование константной памяти.
#define N 100
__constant__ float devInputMatrix[N];
__global__ void transposeMatrixSlow(float* inputMatrix, float*

outputMatrix, int width, int height)
{ int xIndex = blockDim.x

* blockIdx.x + threadIdx.x;
int yIndex = blockDim.y * blockIdx.y + threadIdx.y;
if ((xIndex < width) && (yIndex < height))
{ int inputIdx = xIndex + width * yIndex;
int outputIdx = yIndex + height * xIndex;
outputMatrix[outputIdx] = inputMatrix[inputIdx]; } }
void host_function() {
   float devInputMatrix[N];
    ...
cudaMemcpy(devInputMatrix, inputMatrix, byteSize,
cudaMemcpyHostToDevice);
}
Использование константной памяти.#define N 100__constant__ float devInputMatrix[N];__global__ void transposeMatrixSlow(float* inputMatrix, float* outputMatrix, int width, int height){  int

Слайд 23Использование разделяемой памяти
#define BLOCK_DIM 16
__global__ void transposeMatrixFast(float* inputMatrix, float* outputMatrix,

int width, int height)
{ __shared__ float temp[BLOCK_DIM][BLOCK_DIM];
int xIndex =

blockIdx.x * blockDim.x + threadIdx.x;
int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
if ((xIndex < width) && (yIndex < height))
{
int idx = yIndex * width + xIndex;
temp[threadIdx.y][threadIdx.x] = inputMatrix[idx];
}
__syncthreads();
xIndex = blockIdx.y * blockDim.y + threadIdx.x;
yIndex = blockIdx.x * blockDim.x + threadIdx.y;
if ((xIndex < height) && (yIndex < width))
{
int idx = yIndex * height + xIndex;
outputMatrix[idx] = temp[threadIdx.x][threadIdx.y];
}
}
Использование разделяемой памяти#define BLOCK_DIM 16__global__ void transposeMatrixFast(float* inputMatrix, float* outputMatrix, int width, int height){ __shared__ float temp[BLOCK_DIM][BLOCK_DIM];

Слайд 24Результаты вычислений
матрица размерностью 2048 * 1536 = 3145728 элементов и

20 итераций в нагрузочных циклах

Результаты вычисленийматрица размерностью 2048 * 1536 = 3145728 элементов и 20 итераций в нагрузочных циклах

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

Если не удалось найти и скачать доклад-презентацию, Вы можете заказать его на нашем сайте. Мы постараемся найти нужный Вам материал и отправим по электронной почте. Не стесняйтесь обращаться к нам, если у вас возникли вопросы или пожелания:

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

Что такое TheSlide.ru?

Это сайт презентации, докладов, проектов в PowerPoint. Здесь удобно  хранить и делиться своими презентациями с другими пользователями.


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

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