Параллельное программирование с CUDA. Часть 1: Введение

  • Tutorial

Еще одна статья о 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-программы


  1. Хост выделяет нужное количество памяти на устройстве.
  2. Хост копирует данные из своей памяти в память устройства.
  3. Хост стартует выполнение определенных ядер на устройстве.
  4. Устройство выполняет ядра.
  5. Хост копирует результаты из памяти устройства в свою память.

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

Ядра


Рассмотрим более детально процесс написания кода для ядер и их запуска. Важный принцип — ядра пишутся как (практически) обычные последовательные программы — то-есть вы не увидите создания и запуска потоков в коде самих ядер. Вместо этого, для организации параллельных вычислений GPU запустит большое количество копий одного и того же ядра в разных потоках — а точнее, вы сами говорите сколько потоков запустить. И да, возвращаясь к вопросу эффективности использования GPU — чем больше потоков вы запускаете (при условии что все они будут выполнять полезную работу) — тем лучше.
Код для ядер отличается от обычного последовательного кода в таких моментах:
  1. Внутри ядер вы имеете возможность узнать «идентификатор» или, проще говоря, позицию потока, который сейчас выполняется — используя эту позицию мы добиваемся того, что одно и то же ядро будет работать с разными данными в зависимости от потока, в котором оно запущено. Кстати, такая организация параллельных вычислений называется SIMD (Single Instruction Multiple Data) — когда несколько процессоров выполняют одновременно одну и ту же операцию но на разных данных.
  2. В некоторых случаях в коде ядра необходимо использовать различные способы синхронизации.

Каким же образом мы задаем количество потоков, в которых будет запущено ядро? Поскольку 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 <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"

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<milliseconds>(system_clock::now() - start);
  cout<<"OpenMP time (ms):" << duration.count() << endl;

  memset(imageGrayArray, 0, sizeof(unsigned char)*numRows*numCols);  

  RGBtoGrayscaleCUDA(imageArray, imageGrayArray, numRows, numCols);

  return 0;
}


Тут все довольно очевидно — читаем файл с изображением, подготавливаем указатели на цветное и в оттенках серого изображение, запускаем вариант
с OpenMP и вариант с CUDA, замеряем время. Функция prepareImagePointers имеет следующий вид:
prepareImagePointers
template <class T1, class T2>
void prepareImagePointers(const char * const inputImageFileName,
                          cv::Mat& inputImage, 
                          T1** inputImageArray, 
                          cv::Mat& outputImage,
                          T2** outputImageArray, 
                          const int outputImageType)
{
  using namespace std;
  using namespace cv;

  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<char>(0);
  *outputImageArray  = (T2*)outputImage.ptr<char>(0); 
}


Я пошел на небольшую хитрость: дело в том, что мы выполняем очень мало работы на каждый пиксел изображения — то-есть при варианте с CUDA встает упомянутая выше проблема соотношения времени выполнения полезных операций к времени выделения памяти и копирования данных, и в результате общее время CUDA варианта будет больше OpenMP варианта, а мы же хотим показать что CUDA быстрее:) Поэтому для CUDA будет измеряться только время, потраченное на выполнение собственно конвертации изображения — без учета операций с памятью. В свое оправдание скажу, что для большого класса задач время полезной работы будет все-таки доминировать, и CUDA будет быстрее даже с учетом операций с памятью.
Далее напишем код для OpenMP варианта:
openMP.hpp
#include <stdio.h>

#include <omp.h>
#include <vector_types.h>

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 <cuda.h>

#define checkCudaErrors(val) check( (val), #val, __FILE__, __LINE__)

template<typename T>
void check(T err, const char* const func, const char* const file, const int line) {
  if (err != cudaSuccess) {
    std::cerr << "CUDA error at: " << file << ":" << line << std::endl;
    std::cerr << cudaGetErrorString(err) << " " << func << std::endl;
    exit(1);
  }
}


cudaMalloc — аналог malloc для GPU, cudaMemcpy — аналог memcpy, имеет дополнительный параметр в виде enum-а, который указывает тип копирования: cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice.
Далее необходимо задать размеры сетки и блока и вызвать ядро, не забыв измерить время:
Скрытый текст
 dim3 blockSize;
  dim3 gridSize;
  int threadNum;

  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<<<gridSize, blockSize>>>(d_imageRGBA, d_imageGray, numRows, numCols);
  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
  float milliseconds = 0;
  cudaEventElapsedTime(&milliseconds, start, stop);
  std::cout << "CUDA time simple (ms): " << milliseconds << std::endl;


Обратите внимание на формат вызова ядра — kernel_name<<<gridSize, blockSize>>>. Код самого ядра также не очень сложный:
rgba_to_grayscale_simple
__global__
void rgba_to_grayscale_simple(const uchar4* const d_imageRGBA,
                              unsigned char* const d_imageGray,
                              int numRows, int numCols)
{
    int y = blockDim.y*blockIdx.y + threadIdx.y;
    int x = blockDim.x*blockIdx.x + threadIdx.x;
    if (x>=numCols || y>=numRows)
      return;
    const int offset = y*numCols+x;
    const uchar4 pixel = d_imageRGBA[offset];
    d_imageGray[offset] = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z;
}


Здесь мы вычисляем координаты y и x обрабатываемого пиксела, используя ранее описанные переменные threadIdx, blockIdx и blockDim, ну и выполняем конвертацию. Обратите внимание на проверку if (x>=numCols || y>=numRows) — так как размеры изображения не обязательно будут делится нацело на размеры блоков, некоторые блоки могут «выходить за рамки» изображения — поэтому необходима эта проверка. Также, функция ядра должна помечаться спецификатором __global__ .
Последний шаг — cкопировать результат назад с GPU на CPU и освободить выделенную память:
Скрытый текст
  checkCudaErrors(cudaMemcpy(h_imageGray, d_imageGray, sizeof(unsigned char) * numPixels, cudaMemcpyDeviceToHost));
  cudaFree(d_imageGray);
  cudaFree(d_imageRGBA);


Кстати, CUDA позволяет использовать C++ компилятор для host-кода — так что запросто можно написать обертки для автоматического освобождения памяти.
Итак, запускаем, измеряем (размер входного изображения — 10,109 × 4,542):
OpenMP time (ms):45
CUDA time simple (ms): 43.1941

Конфигурация машины, на которой проводились тесты:
Скрытый текст
Процессор: Intel® Core(TM) i7-3615QM CPU @ 2.30GHz.
GPU: NVIDIA GeForce GT 650M, 1024 MB, 900 MHz.
RAM: DD3, 2x4GB, 1600 MHz.
OS: OS X 10.9.5.
Компилятор: g++ (GCC) 4.9.2 20141029.
CUDA компилятор: Cuda compilation tools, release 6.0, V6.0.1.
Поддерживаемая версия OpenMP: OpenMP 4.0.

Получилось как-то не очень впечатляюще:) А проблема все та же — слишком мало работы выполняется над каждым пикселом — мы запускаем тысячи потоков, каждый из которых отрабатывает практически моментально. В случае с CPU такой проблемы не возникает — OpenMP запустит сравнительно малое количество потоков (8 в моем случае) и разделит работу между ними поровну — таким образом процессоры будет занят практически на все 100%, в то время как с GPU мы, по сути, не используем всю его мощь. Решение довольно очевидное — обрабатывать несколько пикселов в ядре. Новое, оптимизированное, ядро будет выглядеть следующим образом:
rgba_to_grayscale_optimized
#define WARP_SIZE 32

__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<elemsPerThread && i<numCols; i+=WARP_SIZE, ++j)
    {
      const int offset = y*numCols+i;
      const uchar4 pixel = d_imageRGBA[offset];
      d_imageGray[offset] = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z;
    }
}


Здесь не все так просто как с предыдущим ядром. Если разобраться, теперь каждый поток будет обрабатывать elemsPerThread пикселов, причем не подряд, а с расстоянием в WARP_SIZE между ними. Что такое WARP_SIZE, почему оно равно 32, и зачем обрабатывать пиксели пободным образом, будет более детально рассказано в следующих частях, сейчас только скажу что этим мы добиваемся более эффективной работы с памятью. Каждый поток теперь обрабатывает elemsPerThread пикселов с расстоянием в WARP_SIZE между ними, поэтому x-координата первого пиксела для этого потока исходя из его позиции в блоке теперь рассчитывается по несколько более сложной формуле чем раньше.
Запускается это ядро следующим образом:
Скрытый текст
  threadNum=128;
  const int elemsPerThread = 16;
  blockSize = dim3(threadNum, 1, 1);
  gridSize = dim3(numCols / (threadNum*elemsPerThread) + 1, numRows, 1);
  cudaEventRecord(start);
  rgba_to_grayscale_optimized<<<gridSize, blockSize>>>(d_imageRGBA, d_imageGray, numRows, numCols, elemsPerThread);
  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
  milliseconds = 0;
  cudaEventElapsedTime(&milliseconds, start, stop);
  std::cout << "CUDA time optimized (ms): " << milliseconds << std::endl;


Количество блоков по x-координате теперь рассчитывается как numCols / (threadNum*elemsPerThread) + 1 вместо numCols / threadNum + 1. В остальном все осталось так же.
Запускаем:
OpenMP time (ms):44
CUDA time simple (ms): 53.1625
CUDA time optimized (ms): 15.9273

Получили прирост по скорости в 2.76 раза (опять же, не учитывая время на операции с памятью) — для такой простой проблемы это довольно неплохо. Да-да, эта задача слишком простая — с ней достаточно хорошо справляется и CPU. Как видно из второго теста, простая реализация на GPU может даже проигрывать по скорости реализации на CPU.
На сегодня все, в следующей части рассмотрим аппаратное обеспечение GPU и основные шаблоны параллельной коммуникации.
Весь исходный код доступен на bitbucket.
EPAM
137.51
Компания для карьерного и профессионального роста
Share post

Comments 36

    –7
    Интересный подход, вместо того, чтобы выбрать задачу, на которой хорошо показать мощь GPU, вы просто взяли, и сделали на CPU больше работы под таймером. Интересно, производители нового железа тоже так показывают, что новое железо лучше?
      +6
      Ну целью первой статьи было скорее показать, как написать простую программу на CUDA, чем доказать, что вычисления на GPU можно выполнить быстрее чем на CPU. В следующих статьях уже будут примеры, где CUDA вариант действительно быстрее чем CPU вариант — просто они будут несколько сложнее.
        –7
        Ну знаете, такие вот публично написанные вещи навевают соответствующие мысли о компании. Например, что это нормальная практика с заказчиками :)
          +2
          Не совсем Вас понимаю. Что именно имеется ввиду под «нормальная практика с заказчиками»?
            –8
            Написать одно, померять другое, обнаружить фэйковое улучшение, разумеется. Или еще, найти дурацкое решение задачи, полениться найти пригодное, сделать дурацкое и сказать, мол, ну, нам лень было, потому мы вот тут подкрутили показатели.
            Короче, негативные впечатления.
              +6
              А где было сказано про фэйковые улучшения и подкрутку показателей? Я довольно четко написал, что для данной задачи реальный выигрыш мы не получили — однако сами вычисления выполнились быстрее. В примере для следующей статьи уже будет получено реальное ускорение, с учетом всех операций. Ну а про компанию это вообще не понятно к чему.
              Повторю еще раз: цель данной конкретной статьи не состояла в написании супер-пупер быстрого кода на GPU — но в объяснении основ CUDA и написании простого примера для иллюстрации этих основ. Вы даже могли бы обратить внимание на плашку «tutorial».
        +3
        Не сердитесь, то что вычисления на GPU – это быстро и круто все и так понимают, а вот пояснения принципов работы CUDA и простые примеры для тех кто делает первые шаги в этом деле лишними никогда не будут.
        0
        ну хоть что-то реальное делают на куде, а то кругом одни решатели СЛАУ
        может, напишите mp3 енкодер? мне кажется, эту задачу можно легко распараллелить. вот чем реально было бы сравнить производительность CPU и GPU
          +1
          Обработка изображений, много чего реализовано на CUDA в OpenCV, кстати недавно сталкивался с обучением сверточных нейронок на CUDA, прирост впечатляет
            +1
            В примере к второй статье будет размытие изображения по Гауссу. К четвертой — поразрядная сортировка. К шестой — смешивание изображений по Пуассону. Насчет mp3 енкодера — я посмотрю, может действительно будет примером к одной из статей, если получится разобраться:)
              +1
              Я по-быстрому погуглил — так как одним из шагов mp3 енкодинга является FFT или какой-то его аналог, а разные бенчмарки говорят, что для большого количества данных GPU реализация FFT дает определенный прирост в скорости — можно потенциально ускорить этот процесс. Однако, думаю что это нетривиальная задача. — Вот тут например пишут, что в 2008 nVidia проводила конкурс на реализацию mp3 енкодера на CUDA.
              +1
              Из того, что реально работает:
              У нас в разработке программа, вычисляющая специфические интегралы теперь и с использованием GPU. Пришлось изрядно повозиться, чтобы настроить одновременное вычисление рациональных функций в разных точках, но результат был действительно хорош.
              К сожалению, проект пока закрытый, так что никаких точных данных дать не могу, но даже игровая десктопная GPU в несколько раз перебивает четырехядерный процессор.
                0
                Ну, с каждым годом доля GPU в суперкомпьютерах стремительно растет. Поверьте, там действительно работают реальные программы и производят реальные, полезные вычисления :)
                0
                А нет ли какого-либо способа обобщенно писать ядра для OpenCL и CUDA? Или это из области фантастики? Я бы видел это так: пишешь на обобщенном языке, применяешь тулзу с параметром --nvidia-kernel или --amd-kernel и получаешь код требующий компиляции или ядро, которое уже можно применять.
                  0
                  Я ни о чем таком пока не слышал и в гугле не нашел, хотя тут и говорят, что конвертация CUDA кода в OpenCL код довольно прямолинейна. Все что нашел — cu2cl — автоматический транслятор CUDA в OpenCL, естественно, с некоторыми ограничениями.
                  Как понимаете, все дело в определенных различиях между технологиями (которые, кстати, хорошо описаны по первой ссылке).
                    –1
                    Спасибо! // Кармы нет чтоб +1 поставить, приходится не-по-программерски )
                    0
                    А зачем примешивать CUDA к OpenCL, если последняя работает и так на всех устройствах без изменения кода?
                    Другое дело, что обычно для оптимального кода под CPU/AVX нужно писать немного разные ядра, чем на GPU. Но все ядра работают везде, с разной оптимальностью конечно (проверено от titan и firepro, до xeon phi и CPU).
                      +1
                      Не всё так гладко. Nvidia видеокарты не держат OpenCL 2.0. (и даже с предыдущими версиями драйвера выходят очень поздно). т.е. по факту Nvidia довольно давно забивает на поддержку OpenCL.

                      только начиная с OpenCL 2.0 введены atomic операции над float'ами (в CUDA они есть начиная с CC 2.0, т.е. грубо говоря начиная с GTX 4xx серии, довольно древних карт)

                      без atomic операций над float'ами очень многие алгоритмы (особенно различные физические симуляции и агрегации реальных вещественных данных) — становятся нежизнеспособны.

                      да, существует код для более ранних версий OpenCL — позволяющий реализовать костыль, имитирующий atomic операции над float'ами, но производительность эти трюки убивают в разы.

                      поэтому на практике — приходится писать CUDA код для Nvidia и OpenCL код для AMD как минимум ради atomic float'ов.
                        0
                        Никогда не заморачивался с атомик операциями, наверное в силу специфики или простоты задач. Можете привести примеры или ссылку, где атомик спасают ситуацию?

                        Да и устройства с opencl 2.0 я увидел реально сосвсем недавно, кажись как вышел amd драйвер омега или как его там. До этого всегда было 1.2 или даже 1.1 на нвидиа.

                        У нас проблема другая, девайсы с двойной точностью от нвидиа становятся с совсем невменяемой ценой(хорошо что есть пока titan). А карты amd с их новой GCN1.1 работают медленнее старых 7970 и 280Х, не смотря на большее кол-во ядер.
                          0
                          Доольно простой пример, где нужен атомик над float — построение гистограммы для любых float значений. Скажем, имея большое количество записей о спортсменах, вы хотите найти средний вес в зависимости от роста — тогда вам сначала нужно построить гистограмму 'рост'->'суммарный вес всех спортсменов с таким ростом', и тут уже без атомик не обойтись, а вес вполне может быть float.
                            0
                            да, в большинстве случаев это — построение гистограмм. могу привести пару примеров из своих симуляций:

                            1. есть множество силовых точек, связанных «резинками» (например, как в таких задачах en.wikipedia.org/wiki/Force-directed_graph_drawing или в задачах на симуляцию физики мягкого тела (мой случай))

                            при условии, что количество резинок, связанных с одной силовой точкой, может быть произвольным, и на каждом шаге симуляции нужно «повлиять» на каждую силовую точку вектором силы каждой растянутой «резинки», связанной с силовой точкой — стандартный способ избежать atomic float'ов, когда kernel совершает проход по множеству точек, и каждая точка неконкурентно аккумулирует вектора сил всех связанных с ней резинок — не подходит (количество резинок у каждой силовой точки значительно отличается, что будет запирать синхронные WARPы. к тому же, необходимо будет использовать динамические списки резинок — т.е. гораздо более сложный код менеджмента модели).

                            при наличии atomic float'ов — мы просто совершаем пробег по множеству резинок, каждая из которых добавляет свой вектор силы к обоем силовым точкам, к которым привязана (через atomicAdd float). т.к. две резинки могут одновременно попытаться добавить свой вектор силы к одной и той же силовой точке, без atomic float'ов такой подход не работает.

                            приятно то, что atomic float'ы в современных nvidia картах настолько производительно реализованы, что в большинстве случаев попытки обойти их использование более хитрым и производительным алгоритмом приведут к уменьшению производительности.

                            т.е. при желании решать реальные задачи, а не фигурно извращаться с архитектурой GPGPU — atomic float'ы + CUDA + современные nvidia GPU (и чем дальше, тем лучше с этим) позволяют просто писать SIMD код в лоб, при этом имея производительность на грани теоретического максимума.

                            — это наиболее простой и понятный пример применения atomic float'ов.

                            в моей практике была куча ситуаций, когда происходило такое распараллеленное аккумулирование векторов сил на силовые точки. во всех случаях сначала писался код в лоб с atomicAdd (т.к. это по сути мнемонический эквивалент предполагаемой операции всегда)

                            Затем в некоторых случаях производились попытки оптимизировать это, «вывернув» структуры данных под то, чтобы каждая силовая точка становилась аккумулятором для списка внешних воздействий. это всегда усложняло код (денормализация данных под каждый конкретный kernel), но, как ни странно — в большинстве случаев не только не увеличивало производительность, а, наоборот, уменьшало её (я для себя пришёл к выводу, что денормализованные данные занимают большое количество памяти, что приводит к резкому уменьшению эффективности работы кешей).

                            Таким образом, здравый смысл и добро восторжествовали, и сегодня в CUDA + NVidia окружении простой код в большинстве случаев работает наиболее эффективно, и «камень с души» о непроизведённой тонкой оптимизации оказывается снят в общем случае. Atomic float'ы рулят!

                            p.s. субьективно — начиная с GTX 4 серии (т.е. с того момента, как в CC 2.0 практически сошла на нет необходимость возни c global memory coalescing) смысла в попытках обеспечить когерентность чтений, использовать shared memory в кернелах, использовать хитрые алгоритмы с синхронизациями и прочий ловлевел hardcore — канули в лета, и я очень рад, что это так.

                            для примера: я очень плотно занимался поиском оптимального алгоритма для SPH симуляций. в CUDA samples с давних времён был включен классический пример «particles», в котором ещё остались оптимизации под СС 1.2 (GTX 2 серии и раньше, когда CUDA только зарождался). но все эти ухищрения только мешали раскрыться потенциалу GTX 7 серии (я экспериментировал на GTX Titan).

                            в итоге, оптимальной оказалась практика, предложенная Rama Hoetzlein (он же провёл очень хороший обзор эволюции решений этой задачи здесь: on-demand.gputechconf.com/gtc/2014/presentations/S4117-fast-fixed-radius-nearest-neighbor-gpu.pdf). К сожалению, его собственное решение fluids3.com/ содержит несколько эммм досадных недоразумений в коде, в результате чего его код из коробки выполняется медленнее, чем «particle» демка из cuda samples. Но сделанный мной гибрид обоих решений сумел быть существенно более производительным, чем оба исходных варианта.

                            Это я всё к чему. Итоговый код не содержит никаких низкоуровневых фокусов с __syncthreads, sharedMem, разделений структур на отдельные вектора с денормализацией и сортировкой по ключам, хитрых fastRadixSort; и при этом выполняется в разы быстрее на современных Nvidia GPU.

                            Поэтому, если вы видите применение низкоуровневых фич CUDA, настоятельно рекомендую убедиться в том, что автор — не писал это «по книжкам» под допотопные устройства году эдак в 2009-2010 и/или не является нанотехнологически-академическим теоретиком из серии произвели исследование на миллион и выяснили «на GTX 9800 мой код даёт x10 ускорение». Велика вероятность, что на GTX 780 он будет выполняться в лучшем случае с той же скоростью, а то и медленнее решения в лоб на тех же пресловутых atomic float'ах. И здесь по опыту — очень легко попасть в ловушку устаревших мировоззрений сумрачного гения, потратить кучу времени, заморочить голову и в итоге так и не получить оптимального кода под свою задачу.

                            На этом всё. Извините, накипело. Надеюсь, мой опыт кому-то будет полезен.
                              0
                              Вы, судя по всему, «в теме»:) Хотел спросить: то-есть если из этого примера убрать перемещение весов фильтра в общую память блока — то на GPU с Compute Capability >= 2.0 производительность останеться прежней?
                                0
                                Прошу прощения, не из этого примера, а из примера ко второй статье, и не «останеться» а «останется»:)
                                  0
                                  не берусь оценивать, я бы посоветовал провести эксперимент и сравнить результаты.

                                  здесь слишком простой код, и то, что в моих «реальных» задачах даёт размытую на общем фоне оптимизацию в 2-3%, здесь легко может давать 20-30% и больше.

                                  это, кстати, одна из ловушек: вы пишите примитивный тестовый код, он с применением тонкой оптимизации даёт вам выигрыш в 30-50%, вы усложняете код и продолжаете поддерживать общую идею тонкой оптимизации, удерживая в голове её «стоимость» в 30-50%. но эффект уже давно мог быть размыт и в реальной ситуации составлять менее 2-3%, а для вас любое изменение алгоритма — всё ещё будет обозначать тотальную концентрацию и кучу оптимизационных экспериментов.

                                  у меня есть сложная симуляция (развитие вот этого начинания: habrahabr.ru/post/153169/), код которой расширялся и менялся десятки раз, и я множество раз проводил повторные оптимизации.

                                  за это время я выделил для себя такие принципы:

                                  — не стоит переоценивать теоретическую предсказуемость изменений производительности при изменениях кода. если есть 2-3 решения и хочется хороший результат — лучше реализовать их все и промерять, результаты часто оказываются очень неожиданными в силу того, что архитектура современных GPGPU очень нелинейной стала, как и поведение nvcc.

                                  например, добавил float4 в структуру в которой их уже 8, и получил -30% производительности, хотя до этого пару раз добавлял без всякого оверхеда. значит, вылез за пределы какого-то кеша на своих данных, или количества регистров

                                  — в итоге, ни один реальный алгоритм мне не удалось оптимизировать с использованием sharedMemory или __syncthreads. эти фокусы работают только на очень простых данных, либо на очень старых устройствах.

                                  — разделять данные { float3 xyz; float3 fxfyfz } на два отдельных вектора — бывает полезно попробовтаь

                                  — не стоит бояться atomic float'ов, часто они работают быстрее варианта с денормализованными данными (сгенерировать много данных для более простого пробега по ним бывает менее оптимально, несмотря на WARP блокировки — видимо, в силу того, что более компактные структуры данных лучше кешируются).

                                  — эксперименты с параметрами CUDA компилятора часто приводят к очень интересным результатам: CC 1.x на простых симуляциях был сильно быстрее чем CC 2.0, но CC 1.x не умел atomicFloat'ы. зато CC 3.x можно заставить одновременно использовать atomicFloat'ы и генерировать менее прецезионный код ключиками "--use_fast_math –Xptxas –v,–abi=no". обещают эту фичу выпилить (там не по самым строгим стандартам считаются float'ы), но 180-250% прозиводительности заставляют меня до последнего генерировать не до конца ГОСТ, зато быстрый код.
                                    0
                                    Ясно, спасибо. Спросил потому-что практические задания для udacity-курса выполнялись на GPU с архитектурой Fermi — то-есть СС=2.x, но инструкторы все-таки советовали перемещать данные в shared memory, да я и сам замечал прирост в производительности. Я ведь правильно понимаю, дело в том, что начиная с СС 2.0 CUDA и сама перемещает часто используемые данные в L1 кэш — именно поэтому использование shared memory может и не дать прироста?
                                      0
                                      Честно говоря, не знаю. Я делюсь своим практическим опытом. В версию с автоперемещением данных в L1 кеш готов верить.

                                      >но инструкторы все-таки советовали перемещать данные в shared memory

                                      Думаю, доля смысла в этом есть. но очень важно не спугнуть новичков всеми этими наворотами!

                                      Я до сих пор не видел хорошей статьи, в которой главным образом утверждалась бы аксиома «не бойтесь и пробуйте решать ваши реальный задачи в лоб без вникания в моменты тонкой оптимизации! на современных устройствах оптимизации перестали быть критически значимыми». сразу начинают нагонять страху кучей моментов, которые якобы должен удерживать в голове GPGPU программист. А в результате — community чуть менее чем никакое, приток новичков плохой, GPGPU развивается в разы медленнее, чем мог бы.

                                      когда я 3 года назад входил «в тему», мне понадобилось много самурайского духа, чтобы сначала решиться и вкурить все эти особенности архитектуры GPGPU, а затем ещё столько же, чтобы осознать, что все эти фичи в 95% случаев не критичны и в первую очередь нужно просто решиться писать GPGPU код, простой и прямой как молоток. и всё будет хорошо работать.
                              0
                              >У нас проблема другая, девайсы с двойной точностью от нвидиа становятся с совсем невменяемой ценой(хорошо что есть пока titan). А карты amd с их новой GCN1.1 работают медленнее старых 7970 и 280Х, не смотря на большее кол-во ядер.

                              я тоже запасся двумя GTX Titan, однако, до пользы от применения двойной точности у меня ещё ни разу не дошло. греет душу, что если понадобится — есть на чём посчитать :-)

                              >Да и устройства с opencl 2.0 я увидел реально сосвсем недавно, кажись как вышел amd драйвер омега или как его там. До этого всегда было 1.2 или даже 1.1 на нвидиа.

                              Жесть конечно. Начинаю понимать, почему convnet и много других интересных вещей сходу пишутся под nivida. жить без atomic float'ов в XXI веке — это знатное извращенство.
                        +3
                        Статья интересная, но ваш стиль корпоративного блога (боковые полоски а-ля ковер) отвлекают взгляд от чтения. Читать реально сложно.
                          0
                          Эх, я бы с удовольствием ковёр с таким узором на стену повесил! С компьютерными разъёмами-то :)
                            0
                            Я спрошу, можно ли что-то сделать:)
                              0
                              Ну может, чуть менее ярко =) Так-то идея классная, но немного мешается =)
                            0
                            Вопрос: на картинке какого разрешения проводились тестирование? А то я не нашел нигде в коде и статье.
                              0
                              Немаленького: 10,109 × 4,542. Для изображений стандартных размеров смысла использовать CUDA вообще бы не было — подозреваю, что даже вычислительная часть выполнялась бы медленнее, чем на CPU, не учитывая операций с памятью. Думаю, в следующих статьях (кроме 2 — она уже написана) буду проводить тестирование на входах разной размерности.
                                0
                                Первое стоит отметить, что в реальности при конвертации BGRA -> Gray обычно ведут расчет в целых числах.
                                Так же ради интереса запустил функцию преобразования BGRA -> Gray, которая работает на одном ядре i7-4770. Получил следующие результаты:
                                Scalar version — 54.7 ms.
                                SSE2 version — 20.9 ms.
                                AVX2 version — 19.0 ms.
                                Вывод — задача явно упирается в пропускную способность памяти даже для CPU. Иначе говоря, слишком мало вычислений на каждый пиксель, что бы было оправдано использование GPU, хотя наверное и подойдет как учебный пример.
                                  0
                                  Вы правы, я почти так и написал в статье:
                                  слишком мало работы выполняется над каждым пикселом

                                  эта задача слишком простая — с ней достаточно хорошо справляется и CPU

                                  Но это действительно учебный пример, на котором довольно просто показать основы CUDA.
                                  0
                                  Добавил размер изображения в статью.

                              Only users with full accounts can post comments. Log in, please.