Содержание Часть 1: Введение.
Часть 2. Аппаратное обеспечение графического процессора и шаблоны параллельной связи.
Часть 3. Основные алгоритмы графического процессора: сокращение, сканирование и гистограмма.
Часть 4. Основные алгоритмы графического процессора: компактное, сегментированное сканирование, сортировка.
Практическое применение некоторых алгоритмов.
Часть 5. Оптимизация программ на графическом процессоре.
Часть 6: Примеры распараллеливания последовательных алгоритмов.
Часть 7: Дополнительные темы по параллельному программированию, динамический параллелизм.
Шаблоны параллельного общения
Что такое параллельные вычисления? Не что иное, как набор потоков, решающих конкретную задачу совместно .
Ключевое слово здесь «кооператив» — для достижения сотрудничества необходимо использовать определенные механизмы связи между потоками.
При использовании CUDA связь происходит через память: потоки могут читать входные данные, изменять выходные данные или обмениваться «промежуточными» результатами.
В зависимости от того, как именно потоки взаимодействуют через память, существуют разные параллельные модели общения .
В предыдущей части в качестве простого примера использования CUDA рассматривалась задача преобразования цветного изображения в оттенки серого.
Для этого интенсивность каждого пикселя выходного полутонового изображения рассчитывалась по формуле I=A*пиксель.
R+B*пиксель.
G+C*пиксель.
B , Где А, Б, С - константы, пикс — соответствующий пиксель исходного изображения.
Графически этот процесс выглядит так: 
Если абстрагироваться от самого метода расчета выходного значения, мы получим первый шаблон параллельной связи — карта : по каждому элементу входных данных с индексом я , выполняется та же функция, а результат сохраняется в массиве выходных данных под тем же индексом я .
Образец карта очень эффективен на GPU, а также просто выражается в терминах CUDA — достаточно запустить один поток для каждого входного элемента (как это было сделано в предыдущей части для задачи преобразования изображений).
Однако лишь малая часть проблем может быть решена с использованием только этого шаблона.
Если для расчета выходного значения с индексом я Если используется несколько входных элементов, такой шаблон называется собирать , и это может выглядеть так: 
Или вот так: 
Эффективность реализации этого паттерна на CUDA зависит от того, какие входные значения используются при вычислении выходных и их количества — лучше всего, когда используется небольшое количество последовательных элементов.
Обратный узор, разбрасывать - каждый входной элемент влияет на несколько (или один) выходных элементов, графически выглядит так же, как собирать Однако смысл меняется: теперь мы «стартуем» не с выходных элементов, для которых вычисляем их значение, а с входных, влияющих на значения тех или иных выходных элементов.
Довольно часто одну и ту же задачу можно решить в рамках шаблона.
собирать , так разбрасывать .
Например, если мы хотим усреднить 3 соседних входных элемента и записать их в выходной массив, мы можем:
- запустить поток на каждом выходном элементе, где каждый поток будет усреднять значения 3 соседних входных элементов — собирать ;
- или запустить поток для каждого входного элемента, где каждый поток будет добавлять 1/3 значения своего входного элемента к значению соответствующего выходного элемента — разбрасывать .
В случае 2D/3D изображений можно использовать различные типы этого шаблона, например двумерный трафарет фон Неймана: 
или 2D-трафарет Мура: 
Из-за этого ограничения шаблон трафарет обычно реализуется достаточно эффективно в рамках CUDA: достаточно запустить один поток для каждого выходного элемента, и поток сам прочитает нужные ему входные элементы.
При такой организации вычислений эффективность обеспечивается двумя факторами:
- Все данные, необходимые для одного потока, группируются в памяти (в случае одномерного массива — сплошной «кусок» памяти, в 2D случае — несколько участков памяти, расположенных на одинаковом расстоянии друг от друга).
- Значение некоторого входного элемента считывается несколько раз из соседних потоков (конкретное количество чтений зависит от выбранной маски) — появляется возможность «повторно использовать» данные, предоставленные нам CUDA — об этом речь пойдет далее в статье.
Аппаратное обеспечение графического процессора
Давайте посмотрим на высокоуровневую структуру аппаратного обеспечения графического процессора:
- CUDA-совместимый графический процессор состоит из нескольких (обычно десятков) потоковые мультипроцессоры (потоковые мультипроцессоры), далее С.
М.
.
- Каждый С.
М.
, в свою очередь, состоит из нескольких десятков простые/потоковые процессоры (SP) (обычные/потоковые процессоры), а точнее, Ядра CUDA (ядра CUDA).
Эти ребята уже больше похожи на обычный CPU — у них свои регистры, кэш и т. д. Каждый С.
М.
тоже есть свой Общая память (разделяемая память) — это своего рода дополнительный кеш, который доступен всем SP и может использоваться как как кеш для часто используемых данных, так и для «общения» между потоками одного и того же блока CUDA.
- Графический процессор также имеет собственную память, называемую память устройства , общий для всех потоков CUDA — с этим работают функции cudaMalloc И cudaMemcpy (его размеры любят измерять и школьники, и производители графических процессоров).
Соответствие модели CUDA и аппаратного обеспечения графического процессора.
Гарантии CUDA Согласно модели CUDA, программист разбивает задачу на блоки, а блоки на потоки.
Как эти программные объекты сопоставляются с аппаратными блоками графического процессора, описанными выше? 
- Каждый блок будет полностью выполнен в выделенном для него месте.
С.
М.
.
- Распределение блоков по С.
М.
Он занимается графическим процессором, а не программистом.
- Все потоки блокировки Икс будут разделены на группы, называемые коробит (так обычно говорят - перекосы), и сделаны на С.
М.
.
Размер этих групп зависит от модели графического процессора, например для моделей с микроархитектурой.
Ферми оно равно 32. Все потоки из одного варпа выполняются одновременно, занимая определенную часть ресурсов С.
М.
.
Причём они либо выполняют одну и ту же инструкцию (но на разных данных), либо простаивают.
- Все потоки в определенном блоке будут выполняться за один раз.
С.
М.
.
- Все темы определенного ядра будет выполнен до того, как начнет выполняться следующее ядро.
- Какой-то блок Икс будет выполнено до/после/одновременно с каким-либо блоком Да .
- Какой-то блок Икс будет выполнено на каком-то конкретном С.
М.
З .
Синхронизация
Итак, перечислим основные механизмы синхронизации, предоставляемые CUDA:- Барьер — точка в коде ядра, при достижении которой поток может «пройти» дальше, только если все потоки из его квартала достигли этой точки.
Еще раз: барьер позволяет синхронизировать только потоки один блок , и не все темы в принципе! Ограничение вполне естественное, ведь количество блоков, заданное программистом, может существенно превышать количество доступных.
С.
М.
.
- Атомарные операции - аналогично атомарным операциям ЦП, можно просмотреть полный список доступных операций здесь .
- __threadfence — не совсем примитив синхронизации: по достижении этой инструкции поток может продолжить выполнение только после того, как все его манипуляции с памятью станут видны другим потокам — по сути, он заставляет поток сбросить кеш.
Основные принципы эффективного использования CUDA
- Принцип увеличения соотношения (полезное время работы)/(время работы памяти) был рассмотрен в предыдущей статье.
Значение дроби можно увеличить двумя способами — увеличить числитель, уменьшить знаменатель: то есть либо нужно выполнить больше работы, либо потратить меньше времени на операции с памятью.
Помимо очевидного решения — чтобы по возможности сократить количество обращений к памяти, используются следующие принципы эффективного управления памятью:
- Переместите часто используемые данные в более быструю память: поточную локальную память > общая блочная память > > общая память устройства > > память хоста.
Таким образом, если несколько потоков в одном блоке используют одни и те же данные, скорее всего, имеет смысл переместить их в общую память блока.
- Последовательный доступ к памяти: поскольку потоки в блоках фактически выполняются группами деформаций, если потоки в одной и той же деформации работают с данными, расположенными последовательно в памяти, CUDA сможет читать один большой фрагмент памяти для каждой инструкции.
В противном случае, если потоки в варп-доступе к данным разбросаны по памяти, количество обращений к памяти увеличивается.
- Переместите часто используемые данные в более быструю память: поточную локальную память > общая блочная память > > общая память устройства > > память хоста.
- Уменьшение расхождения потоков: особенность CUDA заключается в том, что потоки в одной и той же деформации Всегда либо они выполняют одну и ту же инструкцию, либо простаивают. Таким образом, если код ядра содержит код типа
, то половина потоков деформации (с нечетным индексом) будет ждать, пока вторая половина не выполнит код внутри if. Поэтому подобных ситуаций следует избегать.if (threadIdx.x % 2 == 0) { .} .
Пишем вторую программу на CUDA
Перейдем к практике.В качестве примера, иллюстрирующего изложенную теорию, напишем программу, размывающую изображение.
Принцип работы следующий: значение каналов Р, Г, Б пикселей в выходном размытом изображении рассчитывается как взвешенная сумма значений каналов.
Р, Г, Б (соответственно) все пиксели исходного изображения в определенном трафарете: 
Веса рассчитываются с использованием двумерного распределения Гаусса, но как именно это делается, для нашей задачи не очень важно.
Как видно из описания задачи, для реализации данного алгоритма вполне естественно выбрать шаблон трафарет - ведь каждый пиксель выходного изображения рассчитывается на основе соответствующих соседних пикселей исходного изображения.
Начнем со скелета программы: main.cpp #include <chrono>
#include <iostream>
#include <cstring>
#include <string>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/opencv.hpp>
#include <vector_types.h>
#include "openMP.hpp"
#include "CUDA_wrappers.hpp"
#include "common/image_helpers.hpp"
void prepareFilter(float **filter, int *filterWidth, float *filterSigma)
{
static const int blurFilterWidth = 9;
static const float blurFilterSigma = 2.;
*filter = new float[blurFilterWidth * blurFilterWidth];
*filterWidth = blurFilterWidth;
*filterSigma = blurFilterSigma;
float filterSum = 0.f;
const int halfWidth = blurFilterWidth/2;
for (int r = -halfWidth; r <= halfWidth; ++r)
{
for (int c = -halfWidth; c <= halfWidth; ++c)
{
float filterValue = expf( -(float)(c * c + r * r) / (2.f * blurFilterSigma * blurFilterSigma));
(*filter)[(r + halfWidth) * blurFilterWidth + c + halfWidth] = filterValue;
filterSum += filterValue;
}
}
float normalizationFactor = 1.f / filterSum;
for (int r = -halfWidth; r <= halfWidth; ++r)
{
for (int c = -halfWidth; c <= halfWidth; ++c)
{
(*filter)[(r + halfWidth) * blurFilterWidth + c + halfWidth] *= normalizationFactor;
}
}
}
void freeFilter(float *filter)
{
delete[] filter;
}
int main( int argc, char** argv )
{
using namespace cv;
using namespace std;
using namespace std::chrono;
if( argc != 2)
{
cout <<" Usage: blur_image imagefile" << endl;
return -1;
}
Mat image, blurredImage, referenceBlurredImage;
uchar4 *imageArray, *blurredImageArray;
prepareImagePointers(argv[1], image, &imageArray, blurredImage, &blurredImageArray, CV_8UC4);
int numRows = image.rows, numCols = image.cols;
float *filter, filterSigma;
int filterWidth;
prepareFilter(&filter, &filterWidth, &filterSigma);
cv::Size filterSize(filterWidth, filterWidth);
auto start = system_clock::now();
cv::GaussianBlur(image, referenceBlurredImage, filterSize, filterSigma, filterSigma, BORDER_REPLICATE);
auto duration = duration_cast<milliseconds>(system_clock::now() - start);
cout<<"OpenCV time (ms):" << duration.count() << endl;
start = system_clock::now();
BlurImageOpenMP(imageArray, blurredImageArray, numRows, numCols, filter, filterWidth);
duration = duration_cast<milliseconds>(system_clock::now() - start);
cout<<"OpenMP time (ms):" << duration.count() << endl;
cout<<"OpenMP similarity:" << getEuclidianSimilarity(referenceBlurredImage, blurredImage) << endl;
for (int i=0; i<4; ++i)
{
memset(blurredImageArray, 0, sizeof(uchar4)*numRows*numCols);
start = system_clock::now();
BlurImageCUDA(imageArray, blurredImageArray, numRows, numCols, filter, filterWidth);
duration = duration_cast<milliseconds>(system_clock::now() - start);
cout<<"CUDA time full (ms):" << duration.count() << endl;
cout<<"CUDA similarity:" << getEuclidianSimilarity(referenceBlurredImage, blurredImage) << endl;
}
freeFilter(filter);
return 0;
}
Точки:
- Читаем файл изображения, готовим указатели на исходное изображение и полученное размытое изображение.
Функция подготовитьImagePointers остается прежним, при необходимости вы можете просмотреть его исходный код на bitbucket.
- Готовим фильтр Гаусса — то есть набор наших весов.
Мы также запоминаем используемые параметры фильтра, чтобы затем передать их в OpenCV и получить образец размытого изображения для проверки корректности работы наших алгоритмов.
- Вызываем функцию размытия по Гауссу из OpenCV, сохраняем полученный образец и измеряем затраченное время.
- Вызываем функцию размытия по Гауссу, написанную с использованием OpenMP, измеряем затраченное время и сравниваем результат с образцом.
Функция расчета сходства изображений getEuclidianSimilarity следующее: getEuclidianSimilarity
double getEuclidianSimilarity(const cv::Mat& a, const cv::Mat& b) { double errorL2 = cv::norm(a, b, cv::NORM_L2); double similarity = errorL2 / (double) (a.rows * a.cols); return similarity; }
По сути, он находит среднюю сумму квадратов разностей значений всех каналов всех пикселей в двух изображениях. - Мы вызываем CUDA-версию размытия по Гауссу 4 раза, каждый раз измеряя затраченное время и сверяя результат с образцом.
Зачем звонить 4 раза? Дело в том, что при самом первом вызове на инициализацию будет потрачено определенное количество времени — поэтому лучше запустить ее несколько раз и измерить время, затраченное на последующие вызовы.
#include <stdio.h>
#include <omp.h>
#include <vector_types.h>
#include <vector_functions.h>
void BlurImageOpenMP(const uchar4 * const imageArray,
Теги: #Параллельное программирование #GPGPU #cuda #Udacity
-
Джеффрис, Гарольд
19 Oct, 24 -
Яндексу Исполнилось Десять Лет
19 Oct, 24 -
Простой И Аккуратный Короб Для Проводов.
19 Oct, 24 -
Почему Внедрение Crm Снова Терпит Неудачу
19 Oct, 24 -
Работа Со Спрайтами (Unity3D). Часть Ii
19 Oct, 24
