Параллельное программирование с CUDA. Часть 1: Введение
Еще одна статья о CUDA — зачем?
На Хабре было уже немало хороших статей по CUDA — раз, два и другие. Однако, поиск комбинации «CUDA scan» выдал всего 2 статьи никак не связанные с, собственно, алгоритмом scan на GPU —, а это один из самых базовых алгоритмов. Поэтому, вдохновившись только что просмотренным курсом на Udacity — Intro to Parallel Programming, я и решился написать более полную серию статей о CUDA. Сразу скажу, что серия будет основываться именно на этом курсе, и если у вас есть время — намного полезнее будет пройти его.Содержание
На данный момент планируются следующие статьи: Часть 1: Введение.Часть 2: Аппаратное обеспечение GPU и шаблоны параллельной коммуникации.Часть 3: Фундаментальные алгоритмы GPU: свертка (reduce), сканирование (scan) и гистограмма (histogram).Часть 4: Фундаментальные алгоритмы GPU: уплотнение (compact), сегментированное сканирование (segmented scan), сортировка. Практическое применение некоторых алгоритмов.Часть 5: Оптимизация GPU программ.Часть 6: Примеры параллелизации последовательных алгоритмов.Часть 7: Дополнительные темы параллельного программирования, динамический параллелизм.Задержка vs пропускная способность
Первый вопрос, который должен задать каждый перед применением GPU для решения своих задач —, а для каких целей хорош GPU, когда стоит его применять? Для ответа нужно определить 2 понятия: Задержка (latency) — время, затрачиваемое на выполнение одной инструкции/операции.Пропускная способность — количество инструкций/операций, выполняемых за единицу времени.Простой пример: имеем легковой автомобиль со скоростью 90 км/ч и вместимостью 4 человека, и автобус со скоростью 60 км/ч и вместимостью 20 человек. Если за операцию принять перемещение 1 человека на 1 километр, то задержка легкового автомобиля — 3600/90=40с — за столько секунд 1 человек преодолеет расстояние в 1 километр, пропускная способность автомобиля — 4/40=0.1 операций/секунду; задержка автобуса — 3600/60=60с, пропускная способность автобуса — 20/60=0.3(3) операций/секунду.Так вот, CPU — это автомобиль, GPU — автобус: он имеет большую задержку, но также и большую пропускную способность. Если для вашей задачи задержка каждой конкретной операции не настолько важна как количество этих операций в секунду — стоит рассмотреть применение GPU.Базовые понятия и термины CUDA
Итак, разберемся с терминологией CUDA: 
Устройство (device) — GPU. Выполняет роль «подчиненного» — делает только то, что ему говорит CPU.
Хост (host) — CPU. Выполняет управляющую роль — запускает задачи на устройстве, выделяет память на устройстве, перемещает память на/с устройства. И да, использование CUDA предполагает, что как устройство так и хост имеют свою отдельную память.
Ядро (kernel) — задача, запускаемая хостом на устройстве.
При использовании CUDA вы просто пишете код на своем любимом языке программирования (список поддерживаемых языков, не учитывая С и С++), после чего компилятор CUDA сгенерирует код отдельно для хоста и отдельно для устройства. Небольшая оговорка: код для устройства должен быть написан только на языке C с некоторыми 'CUDA-расширениями'.Основные этапы CUDA-программы
Хост выделяет нужное количество памяти на устройстве.
Хост копирует данные из своей памяти в память устройства.
Хост стартует выполнение определенных ядер на устройстве.
Устройство выполняет ядра.
Хост копирует результаты из памяти устройства в свою память.
Естественно, для наибольшей эффективности использования GPU нужно чтобы соотношение времени, потраченного на работу ядер, к времени, потраченного на выделение памяти и перемещение данных, было как можно больше.Ядра
Рассмотрим более детально процесс написания кода для ядер и их запуска. Важный принцип — ядра пишутся как (практически) обычные последовательные программы — то-есть вы не увидите создания и запуска потоков в коде самих ядер. Вместо этого, для организации параллельных вычислений GPU запустит большое количество копий одного и того же ядра в разных потоках —, а точнее, вы сами говорите сколько потоков запустить. И да, возвращаясь к вопросу эффективности использования GPU — чем больше потоков вы запускаете (при условии что все они будут выполнять полезную работу) — тем лучше.Код для ядер отличается от обычного последовательного кода в таких моментах: Внутри ядер вы имеете возможность узнать «идентификатор» или, проще говоря, позицию потока, который сейчас выполняется — используя эту позицию мы добиваемся того, что одно и то же ядро будет работать с разными данными в зависимости от потока, в котором оно запущено. Кстати, такая организация параллельных вычислений называется SIMD (Single Instruction Multiple Data) — когда несколько процессоров выполняют одновременно одну и ту же операцию, но на разных данных.
В некоторых случаях в коде ядра необходимо использовать различные способы синхронизации.
Каким же образом мы задаем количество потоков, в которых будет запущено ядро? Поскольку GPU это все таки Graphics Processing Unit, то это, естественно, повлияло на модель CUDA, а именно на способ задания количества потоков: Сначала задаются размеры так называемой сетки (grid), в 3D координатах: grid_x, grid_y, grid_z. В результате, сетка будет состоять из grid_x*grid_y*grid_z блоков.
Потом задаются размеры блока в 3D координатах: block_x, block_y, block_z. В результате, блок будет состоять из block_x*block_y*block_z потоков. Итого, имеем grid_x*grid_y*grid_z*block_x*block_y*block_z потоков. Важное замечание — максимальное количество потоков в одном блоке ограничено и зависит от модели GPU — типичны значения 512 (более старые модели) и 1024 (более новые модели).
Внутри ядра доступны переменные threadIdx и blockIdx с полями x, y, z — они содержат 3D координаты потока в блоке и блока в сетке соответственно. Также доступны переменные blockDim и gridDim с теми же полями — размеры блока и сетки соответственно.
Как видите, данный способ запуска потоков действительно подходит для обработки 2D и 3D изображений: например, если нужно определенным образом обработать каждый пиксел 2D либо 3D изображения, то после выбора размеров блока (в зависимости от размеров картинки, способа обработки и модели GPU) размеры сетки выбираются такими, чтобы было покрыто все изображение, возможно, с избытком — если размеры изображения не делятся нацело на размеры блока.Пишем первую программу на CUDA
Довольно теории, время писать код. Инструкции по установке и конфигурации CUDA для разных ОС — docs.nvidia.com/cuda/index.html. Также, для простоты работы с файлами изображений будем использовать OpenCV, а для сравнения производительности CPU и GPU — OpenMP.Задачу поставим довольно простую: конвертация цветного изображения в оттенки серого. Для этого, яркость пиксела pix в серой шкале считается по формуле: Y = 0.299*pix.R + 0.587*pix.G + 0.114*pix.B.Сначала напишем скелет программы: main.cpp
#include
#include
#include «openMP.hpp» #include «CUDA_wrappers.hpp» #include «common/image_helpers.hpp»
using namespace cv; using namespace std;
int main (int argc, char** argv) { using namespace std: chrono;
if (argc!= 2) { cout <<" Usage: convert_to_grayscale imagefile" << endl; return -1; }
Mat image, imageGray; uchar4 *imageArray; unsigned char *imageGrayArray;
prepareImagePointers (argv[1], image, &imageArray, imageGray, &imageGrayArray, CV_8UC1);
int numRows = image.rows, numCols = image.cols;
auto start = system_clock: now ();
RGBtoGrayscaleOpenMP (imageArray, imageGrayArray, numRows, numCols);
auto duration = duration_cast
memset (imageGrayArray, 0, sizeof (unsigned char)*numRows*numCols);
RGBtoGrayscaleCUDA (imageArray, imageGrayArray, numRows, numCols);
return 0;
}
Тут все довольно очевидно — читаем файл с изображением, подготавливаем указатели на цветное и в оттенках серого изображение, запускаем вариантс OpenMP и вариант с CUDA, замеряем время. Функция prepareImagePointers имеет следующий вид: prepareImagePointers
template
inputImage = imread (inputImageFileName, IMREAD_COLOR);
if (inputImage.empty ()) { cerr << "Couldn't open input file." << endl; exit(1); }
//allocate memory for the output outputImage.create (inputImage.rows, inputImage.cols, outputImageType);
cvtColor (inputImage, inputImage, cv: COLOR_BGR2BGRA);
*inputImageArray = (T1*)inputImage.ptr
#include
void RGBtoGrayscaleOpenMP (uchar4 *imageArray, unsigned char *imageGrayArray, int numRows, int numCols) { #pragma omp parallel for collapse (2) for (int i = 0; i < numRows; ++i) { for (int j = 0; j < numCols; ++j) { const uchar4 pixel = imageArray[i*numCols+j]; imageGrayArray[i*numCols+j] = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z; } } } Все довольно прямолинейно — мы всего лишь добавили директиву omp parallel for к однопоточному коду — в этом вся красота и мощь OpenMP. Я пробовал поиграться с параметром schedule, но получалось только хуже, чем без него.Наконец, переходим к CUDA. Тут распишем более детально. Сначала нужно выделить память под входные данные, переместить их с CPU на GPU и выделить память под выходные данные:Скрытый текст void RGBtoGrayscaleCUDA(const uchar4 * const h_imageRGBA, unsigned char* const h_imageGray, size_t numRows, size_t numCols) { uchar4 *d_imageRGBA; unsigned char *d_imageGray; const size_t numPixels = numRows * numCols; cudaSetDevice(0); checkCudaErrors(cudaGetLastError()); //allocate memory on the device for both input and output checkCudaErrors(cudaMalloc(&d_imageRGBA, sizeof(uchar4) * numPixels)); checkCudaErrors(cudaMalloc(&d_imageGray, sizeof(unsigned char) * numPixels));
//copy input array to the GPU
checkCudaErrors (cudaMemcpy (d_imageRGBA, h_imageRGBA, sizeof (uchar4) * numPixels, cudaMemcpyHostToDevice));
Стоит обратить внимание на стандарт именования переменных в CUDA — данные на CPU начинаются с h_ (host), данные да GPU — с d_ (device). checkCudaErrors — макрос, взят с github-репозитория Udacity курса. Имеет следующий вид: Скрытый текст
#include
#define checkCudaErrors (val) check ((val), #val, __FILE__, __LINE__)
template
cudaEvent_t start, stop; cudaEventCreate (&start); cudaEventCreate (&stop);
threadNum = 1024;
blockSize = dim3(threadNum, 1, 1);
gridSize = dim3(numCols/threadNum+1, numRows, 1);
cudaEventRecord (start);
rgba_to_grayscale_simple<<
__global__
void rgba_to_grayscale_optimized (const uchar4* const d_imageRGBA,
unsigned char* const d_imageGray,
int numRows, int numCols,
int elemsPerThread)
{
int y = blockDim.y*blockIdx.y + threadIdx.y;
int x = blockDim.x*blockIdx.x + threadIdx.x;
const int loop_start = (x/WARP_SIZE * WARP_SIZE)*(elemsPerThread-1)+x;
for (int i=loop_start, j=0; j
