Пользователь
0,0
рейтинг
17 марта 2009 в 16:00

Разработка → CUDA: Как работает GPU

Внутренняя модель nVidia GPU – ключевой момент в понимании GPGPU с использованием CUDA. В этот раз я постараюсь наиболее детально рассказать о программном устройстве GPUs. Я расскажу о ключевых моментах компилятора CUDA, интерфейсе CUDA runtime API, ну, и в заключение, приведу пример использования CUDA для несложных математических вычислений.

Приступим.

Вычислительная модель GPU:


Рассмотрим вычислительную модель GPU более подробно.
  1. Верхний уровень ядра GPU состоит из блоков, которые группируются в сетку или грид (grid) размерностью N1 * N2 * N3. Это можно изобразить следующим образом:

    Рис. 1. Вычислительное устройство GPU.

    Размерность сетки блоков можно узнать с помощь функции cudaGetDeviceProperties, в полученной структуре за это отвечает поле maxGridSize. К примеру, на моей GeForce 9600M GS размерность сетки блоков: 65535*65535*1, то есть сетка блоков у меня двумерная (полученные данные удовлетворяют Compute Capability v.1.1).
  2. Любой блок в свою очередь состоит из нитей (threads), которые являются непосредственными исполнителями вычислений. Нити в блоке сформированы в виде трехмерного массива (рис. 2), размерность которого так же можно узнать с помощью функции cudaGetDeviceProperties, за это отвечает поле maxThreadsDim.


Рис. 2. Устройство блока GPU.

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

CUDA и язык C:


Сама технология CUDA (компилятор nvcc.exe) вводит ряд дополнительных расширений для языка C, которые необходимы для написания кода для GPU:
  1. Спецификаторы функций, которые показывают, как и откуда буду выполняться функции.
  2. Спецификаторы переменных, которые служат для указания типа используемой памяти GPU.
  3. Спецификаторы запуска ядра GPU.
  4. Встроенные переменные для идентификации нитей, блоков и др. параметров при исполнении кода в ядре GPU .
  5. Дополнительные типы переменных.

Как было сказано, спецификаторы функций определяют, как и откуда буду вызываться функции. Всего в CUDA 3 таких спецификатора:
  • __host__ — выполнятся на CPU, вызывается с CPU (в принципе его можно и не указывать).
  • __global__ — выполняется на GPU, вызывается с CPU.
  • __device__ — выполняется на GPU, вызывается с GPU.

Спецификаторы запуска ядра служат для описания количества блоков, нитей и памяти, которые вы хотите выделить при расчете на GPU. Синтаксис запуска ядра имеет следующий вид:

myKernelFunc<<<gridSize, blockSize, sharedMemSize, cudaStream>>>(float* param1,float* param2), где
  • gridSize – размерность сетки блоков (dim3), выделенную для расчетов,
  • blockSize – размер блока (dim3), выделенного для расчетов,
  • sharedMemSize – размер дополнительной памяти, выделяемой при запуске ядра,
  • cudaStream – переменная cudaStream_t, задающая поток, в котором будет произведен вызов.

Ну и конечно сама myKernelFunc – функция ядра (спецификатор __global__). Некоторые переменные при вызове ядра можно опускать, например sharedMemSize и cudaStream.

Так же стоит упомянуть о встроенных переменных:
  • gridDim – размерность грида, имеет тип dim3. Позволяет узнать размер гридa, выделенного при текущем вызове ядра.
  • blockDim – размерность блока, так же имеет тип dim3. Позволяет узнать размер блока, выделенного при текущем вызове ядра.
  • blockIdx – индекс текущего блока в вычислении на GPU, имеет тип uint3.
  • threadIdx – индекс текущей нити в вычислении на GPU, имеет тип uint3.
  • warpSize – размер warp’а, имеет тип int (сам еще не пробовал использовать).

Кстати, gridDim и blockDim и есть те самые переменные, которые мы передаем при запуске ядра GPU, правда, в ядре они могут быть read only.

Дополнительные типы переменных и их спецификаторы будут рассмотрены непосредственно в примерах работы с памятью.

CUDA host API:


Перед тем, как приступить к непосредственному использованию CUDA для вычислений, необходимо ознакомиться с так называемым CUDA host API, который является связующим звеном между CPU и GPU. CUDA host API в свою очередь можно разделить на низкоуровневое API под названием CUDA driver API, который предоставляет доступ к драйверу пользовательского режима CUDA, и высокоуровневое API – CUDA runtime API. В своих примерах я буду использовать CUDA runtime API.

В CUDA runtime API входят следующие группы функций:
  • Device Management – включает функции для общего управления GPU (получение инфор-мации о возможностях GPU, переключение между GPU при работе SLI-режиме и т.д.).
  • Thread Management – управление нитями.
  • Stream Management – управление потоками.
  • Event Management – функция создания и управления event’ами.
  • Execution Control – функции запуска и исполнения ядра CUDA.
  • Memory Management – функции управлению памятью GPU.
  • Texture Reference Manager – работа с объектами текстур через CUDA.
  • OpenGL Interoperability – функции по взаимодействию с OpenGL API.
  • Direct3D 9 Interoperability – функции по взаимодействию с Direct3D 9 API.
  • Direct3D 10 Interoperability – функции по взаимодействию с Direct3D 10 API.
  • Error Handling – функции обработки ошибок.

Понимаем работу GPU:


Как было сказано, нить – непосредственный исполнитель вычислений. Каким же тогда образом происходит распараллеливание вычислений между нитями? Рассмотрим работу отдельно взятого блока.

Задача. Требуется вычислить сумму двух векторов размерностью N элементов.

Нам известна максимальные размеры нашего блока: 512*512*64 нитей. Так как вектор у нас одномерный, то пока ограничимся использованием x-измерения нашего блока, то есть задействуем только одну полосу нитей из блока (рис. 3).

Рис. 3. Наша полоса нитей из используемого блока.

Заметим, что x-размерность блока 512, то есть, мы можем сложить за один раз векторы, длина которых N <= 512 элементов. В прочем, при более массивных вычислениях, можно использовать большее число блоков и многомерные массивы. Так же я заметил одну интересную особенность, возможно, некоторые из вас подумали, что в одном блоке можно задействовать 512*512*64 = 16777216 нитей, естественно это не так, в целом, это произведение не может превышать 512 (по крайней мере, на моей видеокарте).

В самой программе необходимо выполнить следующие этапы:
  1. Получить данные для расчетов.
  2. Скопировать эти данные в GPU память.
  3. Произвести вычисление в GPU через функцию ядра.
  4. Скопировать вычисленные данные из GPU памяти в ОЗУ.
  5. Посмотреть результаты.
  6. Высвободить используемые ресурсы.

Переходим непосредственно к написанию кода:

Первым делом напишем функцию ядра, которая и будет осуществлять сложение векторов:
// Функция сложения двух векторов
__global__ void addVector(float* left, float* right, float* result)
{
  //Получаем id текущей нити.
  int idx = threadIdx.x;
  
  //Расчитываем результат.
  result[idx] = left[idx] + right[idx];
}

* This source code was highlighted with Source Code Highlighter.


Таким образом, распараллеливание будет выполнено автоматически при запуске ядра. В этой функции так же используется встроенная переменная threadIdx и её поле x, которая позволяет задать соответствие между расчетом элемента вектора и нитью в блоке. Делаем расчет каждого элемента вектора в отдельной нити.

Пишем код, которые отвечает за 1 и 2 пункт в программе:

#define SIZE 512
__host__ int main()
{
  //Выделяем память под вектора
  float* vec1 = new float[SIZE];
  float* vec2 = new float[SIZE];
  float* vec3 = new float[SIZE];

  //Инициализируем значения векторов
  for (int i = 0; i < SIZE; i++)
  {
    vec1[i] = i;
    vec2[i] = i;
  }

  //Указатели на память видеокарте
  float* devVec1;
  float* devVec2;
  float* devVec3;

  //Выделяем память для векторов на видеокарте
  cudaMalloc((void**)&devVec1, sizeof(float) * SIZE);
  cudaMalloc((void**)&devVec2, sizeof(float) * SIZE);
  cudaMalloc((void**)&devVec3, sizeof(float) * SIZE);

  //Копируем данные в память видеокарты
  cudaMemcpy(devVec1, vec1, sizeof(float) * SIZE, cudaMemcpyHostToDevice);
  cudaMemcpy(devVec2, vec2, sizeof(float) * SIZE, cudaMemcpyHostToDevice);
  …
}

* This source code was highlighted with Source Code Highlighter.


Для выделения памяти на видеокарте используется функция cudaMalloc, которая имеет следующий прототип:
cudaError_t cudaMalloc( void** devPtr, size_t count ), где
  1. devPtr – указатель, в который записывается адрес выделенной памяти,
  2. count – размер выделяемой памяти в байтах.

Возвращает:
  1. cudaSuccess – при удачном выделении памяти
  2. cudaErrorMemoryAllocation – при ошибке выделения памяти

Для копирования данных в память видеокарты используется cudaMemcpy, которая имеет следующий прототип:
cudaError_t cudaMemcpy(void* dst, const void* src ,size_t count, enum cudaMemcpyKind kind), где
  1. dst – указатель, содержащий адрес места-назначения копирования,
  2. src – указатель, содержащий адрес источника копирования,
  3. count – размер копируемого ресурса в байтах,
  4. cudaMemcpyKind – перечисление, указывающее направление копирования (может быть cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyHostToHost, cudaMemcpyDeviceToDevice).

Возвращает:
  1. cudaSuccess – при удачном копировании
  2. cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен)
  3. cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте
  4. cudaErrorInvalidMemcpyDirection – неверное направление (например, перепутан источник и место-назначение копирования)

Теперь переходим к непосредственному вызову ядра для вычисления на GPU.

dim3 gridSize = dim3(1, 1, 1);    //Размер используемого грида
dim3 blockSize = dim3(SIZE, 1, 1); //Размер используемого блока

//Выполняем вызов функции ядра
addVector<<<gridSize, blockSize>>>(devVec1, devVec2, devVec3);


* This source code was highlighted with Source Code Highlighter.

В нашем случае определять размер грида и блока необязательно, так как используем всего один блок и одно измерение в блоке, поэтому код выше можно записать:
addVector<<<1, SIZE>>>(devVec1, devVec2, devVec3);

* This source code was highlighted with Source Code Highlighter.

Теперь нам остаеться скопировать результат расчета из видеопамяти в память хоста. Но у функций ядра при этом есть особенность – асинхронное исполнение, то есть, если после вызова ядра начал работать следующий участок кода, то это ещё не значит, что GPU выполнил расчеты. Для завершения работы заданной функции ядра необходимо использовать средства синхронизации, например event’ы. Поэтому, перед копированием результатов на хост выполняем синхронизацию нитей GPU через event.

Код после вызова ядра:
  //Выполняем вызов функции ядра
  addVector<<<blocks, threads>>>(devVec1, devVec2, devVec3);
  
  //Хендл event'а
  cudaEvent_t syncEvent;

  cudaEventCreate(&syncEvent);    //Создаем event
  cudaEventRecord(syncEvent, 0);  //Записываем event
  cudaEventSynchronize(syncEvent);  //Синхронизируем event

  //Только теперь получаем результат расчета
  cudaMemcpy(vec3, devVec3, sizeof(float) * SIZE, cudaMemcpyDeviceToHost);

* This source code was highlighted with Source Code Highlighter.

Рассмотрим более подробно функции из Event Managment API.

Event создается с помощью функции cudaEventCreate, прототип которой имеет вид:
cudaError_t cudaEventCreate( cudaEvent_t* event ), где
  1. *event – указатель для записи хендла event’а.

Возвращает:
  1. cudaSuccess – в случае успеха
  2. cudaErrorInitializationError – ошибка инициализации
  3. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
  4. cudaErrorInvalidValue – неверное значение
  5. cudaErrorMemoryAllocation – ошибка выделения памяти

Запись event’а осуществляется с помощью функции cudaEventRecord, прототип которой имеет вид:
cudaError_t cudaEventRecord( cudaEvent_t event, CUstream stream ), где
  1. event – хендл хаписываемого event’а,
  2. stream – номер потока, в котором записываем (в нашем случае это основной нулевой по-ток).

Возвращает:
  1. cudaSuccess – в случае успеха
  2. cudaErrorInvalidValue – неверное значение
  3. cudaErrorInitializationError – ошибка инициализации
  4. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
  5. cudaErrorInvalidResourceHandle – неверный хендл event’а

Синхронизация event’а выполняется функцией cudaEventSynchronize. Данная функция ожидает окончание работы всех нитей GPU и прохождение заданного event’а и только потом отдает управление вызывающей программе. Прототип функции имеет вид:
cudaError_t cudaEventSynchronize( cudaEvent_t event ), где
  1. event – хендл event’а, прохождение которого ожидается.

Возвращает:
  1. cudaSuccess – в случае успеха
  2. cudaErrorInitializationError – ошибка инициализации
  3. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
  4. cudaErrorInvalidValue – неверное значение
  5. cudaErrorInvalidResourceHandle – неверный хендл event’а

Понять, как работает cudaEventSynchronize, можно из следующей схемы:


Рис. 4. Синхронизация работы основоной и GPU прграмм.

На рисунке 4 блок «Ожидание прохождения Event’а» и есть вызов функции cudaEventSynchronize.

Ну и в заключении выводим результат на экран и чистим выделенные ресурсы.
  //Результаты расчета
  for (int i = 0; i < SIZE; i++)
  {
    printf("Element #%i: %.1f\n", i , vec3[i]);
  }

  //
  // Высвобождаем ресурсы
  //

  cudaEventDestroy(syncEvent);

  cudaFree(devVec1);
  cudaFree(devVec2);
  cudaFree(devVec3);

  delete[] vec1; vec1 = 0;
  delete[] vec2; vec2 = 0;
  delete[] vec3; vec3 = 0;

* This source code was highlighted with Source Code Highlighter.

Думаю, что описывать функции высвобождения ресурсов нет необходимости. Разве что, можно напомнить, что они так же возвращают значения cudaError_t, если есть необходимость проверки их работы.

Заключение


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

P.S.: Получилось не очень кратко. Надеюсь, что не утомил. Если нужен весь исходный код, то могу выслать на почту.
P.S.S: Задавайте вопросы.
Max K. @MaxFX
карма
40,0
рейтинг 0,0
Реклама помогает поддерживать и развивать наши сервисы

Подробнее
Реклама

Самое читаемое Разработка

Комментарии (21)

  • +2
    Большое спасибо за статью.
  • +1
    Пока не всё осилил, но статья интересная.

    P.S.
    Кстати в фотошопе можно не только nVidia задействовать, новые модели Radeon'a тоже неплохо ускоряют обработку.
  • +2
    Ам… дык было же уже?

    Сам только добавить могу, так как по работе уже применял сие =)
    Перенёс один метод под видео карту… получил реальное ускорение в 30(!!!) раз
    НО спешу добавить, что большой скачок получается только когда я воткнул в принципы группировки потоков в блок. и использование __constant__ и __share__ модификаторов переменных.

    Самая идея там если что-то можно объекдинить в блок — объединять, тогда каждый поток может скопировать например сперва в шаред массив свою ячейку, потом синхронизироваться со всеми, и потом работать только с локальной памятью.
    Чтение из обычной 400 клоков, чтение из локальной от 4х… так что вот и выйгрыш. Жаль что там только 16к на локальную память на блок, но пока и этого хватает за глаза.
    все параметры, что менять не надо писать в константную область памяти (её 64к на блок)!

    А так тема реальная!
    • 0
      А каким образом с хоста записать данные в константную память? И не подскажите, возможна ли запись двумерных массивов в эту память?
      • +1
        к примеру так:
        __constant__ int _threshold;

        cudaMemcpyToSymbol( _threshold, threshold, sizeof(int));

        Для работы с двумерными массивами там все есть, и специальные функции для выделения даже, чтобы выравнивание по адресам было!

        Но мне всегда было с линейным удобнее работать. (если уж оптимизировать, то уметь управление над процессом подсчета индекса)
        • 0
          Подробнее о памяти будет в следующей публикации.
    • +1
      Ещё раз отдельное спасибо минусующим! я вас так люблю! =)))))))
  • 0
    Спасибо за статью, как раз начал разбираться с CUDA :)
  • +2
    А продолжение будет? )
    Очень интересно как реализовывать хотя бы несколько более сложные задачи, чем Вы привели в примере. В частности интересно узнать про деление на блоки (как в этом случае нить узнает над какими данными она работает ?).
    Или какой-нибудь реальный пример из жизни перевода программы с CPU на GPU.
    За статью спасибо. Полезно.
    • 0
      Поставьте SDK там есть сэмплы, там есть пример умножения матриц, по разберите его и все станет ясно как Божий свет =)
  • 0
    Спасибо! Очень нужный маттериал.
  • 0
    MaxFX, так держать.
    Однако в Вашу статью вкралась идеологическая ошибка. Размер грида равный 1 никогда не позволит выжать максимум мощи GPU при такой 100% распараллеливаемой задаче как сложение векторов. Да, формально код написан верно, но скорость работы будет на порядок меньше, чем если написать правильно. А правильно будет загрузить все стрим процессоры задачей, задав меньше размер блока и больше размер грида.
    Вот например размер блока 32 треда (один варп), размер грида — 16 блоков. В таком случае разные части массива будут обрабатываться параллельно. Пока весь процесс обработки заключен в одном блоке — работает один из десятков процессоров, а остальные простаивают.
    • 0
      Знаю. Это сделано для облегчения понимания материала.
  • 0
    Картинки не грузятся
  • НЛО прилетело и опубликовало эту надпись здесь
    • 0
      Все в ваших руках. С удовольствием бы почитал.
      У самого под рукой нет видеокарт ATI, сам бы помучил Brook в свободное время.
  • 0
    По вашему мануалу переписал свою старую программку для рисования фигур лиссажу на OpenGL… все шикарно, но… оно работает слишком быстро, рисует с такой дикой скоростью, что по экрану ползают вверх-вниз горизонтальные полосы :) Как бы это дело замедлить? :)
    • 0
      по экрану ползают вверх-вниз горизонтальные полосы

      двойная буферизация?
      • 0
        Без двойной буфферизации у меня бы плыли полосы и на старой программе, сейчас fps за 2000 кадров и это выглядит примерно вот так
        Снимок-Безымянное_окно-5.png - upload images with Picamatic

        Тут конечно еще немного артефактов накинул сам скриншот, но суть ясна.
        • 0
          Ай-ай-ай, извиняюсь, был не прав, когда решил переписать старую прогу, то взял какую то очень давнишнюю версию, где двойной буфер не был включен, включил — все стало отлично.
          • 0
            Что и требовалось доказать :)

Только зарегистрированные пользователи могут оставлять комментарии. Войдите, пожалуйста.