Ускорители вычислений.

Компании NVidia и ATI стали выпускать суперкомпьютеры на базе видеоускорителей NVidia GeForce и ATI Radeon. Суперкомпьютеры: NVidia Tesla и ATI Fire Stream. NVidia Tesla включает вычислительный процессор, приставной суперкомпьютер и вычислительный сервер. Суперкомпьютер и сервер построены на основе 2 и 4 процессоров. При этом сервер предназначен для построения кластерных решений.

CUDA 5.5 – Compute Unified Device Architecture (архитектура устройства для вычислений общего назначения). С точки зрения программиста видеокарта как вычислительное средство представляет собой набор независимых мультипроцессоров. Эта система позволяет создавать специальную функцию (ядра), выполняемые параллельно различными блоками и нитями. При запуске ядра блоки распределяются по доступным мультипроцессорам. Каждый блок состоит из множества нитей. CUDA предоставляет доступ к нескольким уровням памяти: глобальная, локальная, текстурная и т.д.

Ускоритель ClearSpeed представляет собой ускоритель операций над данными с плавающей точкой. Он является сопроцессором. Среда разработки программ основана на языке Си. Характеристики ускорителя: два процессора по 97 ядер, производительность – 66 ГигаФлопсов, память – 0,5 Гб на ядро, мощность – 43 Вт, поддержка нескольких ОС. Каждое ядро обладает своей локальной памятью. Главное приложение состоит из 2 частей: программа, выполняющаяся на процессоре; программа на ускорителе. Этот ускоритель использует для ускорения внутренних циклов программы. В зависимости от анализа команда (функция) может выполняться либо на главном процессоре, либо перехватывается ускорителем.

Причины возникновения видеокарт.

Развитие графических ОС – они выполняют операции с изображениями аппаратно.

Развитие трехмерных игр. Появляется видеокарта GeForce 256, использующая интерфейс OpenGL. Этот интерфейс впервые предлагает вычислять графические преобразования и освещение сцены в видеокарте. В 2001 году появляется GeForce 3, где впервые были использованы Direct X 8.0. Данный стандарт предлагал возможность программируемой обработки вершин и пикселей, так называемый шейдинг. В общем случае видеокарта получает координаты точки и должна выдать её цвет.

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

1. необходимость изучения шейдерных языков
2. необходимость формулировать любой решаемой задачи терминами и переменными традиционного рейбрендинга
3. ограниченный доступ к ресурсам видеокарты
4. отсутствие отладчиков.

В ноябре 2006 года NVidia объявила о выпуске видеокарты с поддержкой Direct X 10.0 (GeForce 8800). В ней были предусмотрены компоненты, предназначенные для вычислений общего назначения: включен унифицированный шейдерный конвейер, позволяющий задействовать любое АЛУ; АЛУ были сконструированы с учетом требований к арифметическим операциям над числами с плавающей точкой; исполняющим устройством видеокарты (GPU) разрешен произвольный доступ к памяти для чтения/записи, а также доступ к программно управляемому КЭШу; NVidia использовала стандартный язык Си и дополнила его новыми директивами, позволяющими задействовать ресурсы видеокарты.

Функция, исполняемая на устройстве, называется kernels.

#include <stdio.h>

#include <cuda.h>

#include …

_ _global_ _ void add (int a, int b, int *c)

{ *c = a+b;

}

int main (void)

{ int c, *dev_c;

cudaMalloc ((void **)&dev_c, sizeof(int))) ;

add<<<1,1>>>(2,7, dev_c);

cudaMemcpy (&c, dev_c, sizeof (int), cudaMemcpyDeviceToHost);

printf (“2+7 = %d\n”, c);

cudaFree (dev_c);

return 0;

}

Угловыми скобками обозначаются аргументы, которые мы собираемся передать исполняющей среде (видеокарте). Это не аргументы функции, а параметры, влияющие на то, как исполняющая среда будет запускать код для видеокарты!

Замечание: мы можем передавать параметры ядру программы как любой другой функции. Мы должны выделить память, если желаем, чтобы видеокарта выполнила некие действия. &dev_c – указатель на область памяти, где будет храниться адрес выделенной области памяти, sizeof – размер выделенной области памяти.

Правило использования указателей.

Разрешается: передавать указатели на память, выделенную cudaMalloc, функциям, исполняемым устройством; использовать указатели на память, выделенную cudaMalloc, для чтения и записи в коде, исполняемым устройством; передавать указатели на память функциям, исполняемым центральным процессором.

Не разрешается: использовать указатели на память для чтения и записи в коде, исполняемым центральным процессором.

Чтобы освободить память, выделенную cudaMalloc, обращаются к функции cudaFree. Попытка обратиться к памяти центрального процессора, адресуемой указателем, из кода, исполняемого устройством, закончится ошибкой. Имеется функция cudaGetDeviceCount (&count), где count имеет тип int.

#include …

int main (void)

{ cudaGetDeviceProp prop;

int dev;

cudaGetDevice (&dev);

printf (“Id текущего устройства: %d\n”, dev);

memset (&prop, 0, sizeof(cudaDeviceProp));

prop.major = 1;

prop.minor = 3;

cudaChooseDevice (&dev, &prop);

printf (“Id ближайшего к уровню 1.3: %d\n”, dev);

cudaSetDevice (dev); // установка активного устройства

}

Cuda построена на концепции, считающей видеокарту массивно-параллельным сопроцессором центральному процессору. Программа на Cuda задействует как центральный, так и видеокарту. Обычный код (последовательный) исполняется на центральном процессоре, а для параллельных расчетов используется код, исполняемый на видеокарте. Он представляет собой набор одновременно выполняющихся нитей или потоков (thread). Таким образом видеокарта рассматривается как устройство, являющимся сопроцессором с собственном памятью и обладающим возможностью вызова параллельных нитей. Нити, исполняемые на видеокарте, обладают следующими особенностями: нити имеют очень малую стоимость создания, управления и уничтожения; для эффективной загрузки видеокарты необходимо использовать много тысяч отдельных нитей.

Схема использования Cuda.

1. выделяется память на видеокарте
2. копируются данные из памяти центрального процессора в выделенную память видеокарты
3. осуществляется запуск ядра или последовательно запускаются несколько ядер
4. результаты расчетов копируются обратно в память центрального процессора
5. освобождается память видеокарты

Все нити выполняются параллельно. Каждая нить может получить информацию через встроенные переменные. SIMT (Thread) – одна команда, множество нитей. При этом нити разбиваются на группы по 32 нити в каждой. Эти группы называются warp. Только нити в пределах одного warpа выполняются физически одновременно. Все выполняемые нити организованы в иерархию.

Сетка соответствует всем нитям, выполняющим данное ядро. Она представляет собой одномерный и двумерный массив блоков. При этом каждый блок – это одномерный, двумерный или трехмерный массив нитей. Все блоки в сетке имеют одинаковую размерность и размер. Каждый блок в сетке имеет свой адрес из одного или двух неотрицательных целых чисел. Каждая нить внутри блока имеет свой собственный адрес (одно, два или три целых неотрицательных числа). Поскольку одно ядро выполняется на множестве нитей, для определения номера нити используются встроенные переменные threadIdx.x, blockIdx.x. Каждая из этих переменных является трехмерным целочисленным вектором. Ядро может получить размеры сетки через встроенные переменные gridDim, blockDim. Общим приемом использования Cuda заключается в разбиении исходной задачи на набор отдельных подзадач, выполняемых независимо.

Каждой подзадаче соответствует набор нитей. Разбиение нитей на warpы происходит отдельно для каждого блока. Все нити одного warpа принадлежат одному блоку. Нити могут взаимодействовать между собой только в пределах одного блока. Данный подход является компромиссом. Имеется два механизма взаимодействия нитей между собой: разделяемая память и синхронизация. Cuda предоставляет возможность барьерной синхронизации. Она осуществляется с помощью функции _ _ syncthreads ( ). Эта функция блокирует нити блока до тех пор, пока все нити блока не войдут в эту функцию. Расширения языка Си, применяемые в интерфейсе Cuda включают: спецификаторы функций показывают, где она будет исполняться и откуда она может быть вызвана; спецификаторы переменных, задающие тип памяти для этих переменных; директивы для ядра; встроенные переменные, содержащие информацию о текущей нити; дополнительные типы данных.

Спецификаторы функций и переменных.

Спецификатор global обозначает ядро. На функцию, исполняемой видеокартой, накладываются ограничения: не поддерживается рекурсия; не поддерживается переменное число входных аргументов. Для задания переменных в память видеокарты используются спецификаторы: _ _device_ _, _ _constant_ _, _ _shared_ _. На их использование накладываются ограничения: переменные могут использоваться только в пределах одного файла; запись в переменные типа constant могут осуществляться только центральным процессором; переменные shared не могут инициализироваться при объявлении.

Добавленные типы.

В Cuda есть одно-, двух-, трех-, четырехмерные наборы базовых типов (int1, float4). Обращения к компонентам вектора идет по именам: x, y, z, w. Для создания значений векторов служит конструкция: make _<typename>. Пример: int2 a = make_int2(1,7), float3 u = make_float3(1,2,3.4f). Для этих типов не поддерживаются векторные покомпонентные операции. Добавлен тип dim3. Этот тип используется для задания размерности. Пример: dim3 blocks (16,16), dim3 grid (256).

Добавленные элементы.

gridDim размерности dim3, blockDim размерности dim3, blockIdx размерности uint3, threadIdx размерности uint3, warpSize размерности int.

KernelName <<< Dg, Db, Ns, S>>> (args); Dg – переменная типа dim3, задающая размерность и размер сетки в блоках, Db – размер и размерность блока в нитях, Ns – необязательная переменная задает дополнительный объем разделяемой памяти в байтах, которая может быть выделена динамически каждому блоку, S – переменная типа cudaStream_t задает поток cuda, в котором должен произойти вызов функции.

myKernel <<<dim3(n/256), dim3 (16,16), 512, myStream>>> (a,n);

Добавленные функции.

Cuda поддерживает все функции из стандартной библиотеки Си. Для ряда функций предусматривается возможность округления. При этом используемый способ округления задается при помощи одного из следующих суффиксов. rn – Округление к ближайшего, rz – округление к нулю, ru – вверх, rd – вниз.

Основы Cuda Host API.

Система Cuda предоставляет программисту функции, которые могут быть использованы только ЦП. Эти функции отвечают за управление видеокартой, за работу с памятью, за управление выполнением кода, за взаимодействие с OpenGL, Direct3D и работу с текстурами. Эта система выступает в двух формах: низкоуровневая форма (cuda driver API) и высокоуровневая (cuda runtime API).

Эти два уровня являются взаимоисключающими. Низкоуровневый дает больше возможностей. К числу его недостатков относятся больший объем кода, требование явной инициализации и отсутствие режима эмуляции. Высокоуровневый не требует явной инициализации, поддерживает эмуляцию и возможность использования дополнительных библиотек (CUFFT, CUBLAS, CUDPP). При необходимости код, основанный на cuda runtime, может быть переписан на driver. Многие функции cuda асинхронны. Управление возвращается до реального завершения операций. К числу асинхронных операций относятся: запуск ядра, некоторые из функций копирования памяти, функции инициализации памяти. Для синхронизации текущей нити используется функция cudaThreadSyncronize. Эта функция дожидается завершения всех операций, вызванных с данной нити.

cudaError_t cudaThreadSyncronize (void);

Каждая функция cuda возвращает значение типа cudaError_t. При успешном выполнении функции возвращается cudaSuccess. При неудаче возвращается код ошибки. Можно получить описание ошибки по её коду при помощи char * cudaGetErrorString (cudaError_t code);