к оглавлению

Технология CUDA

Введение

CUDA, Compute Unified Device Architecture — программно-аппаратная архитектура, позволяющая производить вычисления с использованием графических процессоров NVIDIA, поддерживающих технологию GPGPU (произвольных вычислений на видеокартах). Впервые появились на рынке с выходом чипа NVIDIA восьмого поколения — G80 и присутствует во всех последующих сериях графических чипов, которые используются в семействах ускорителей GeForce, Quadro и NVidia Tesla.

CUDA SDK позволяет программистам реализовывать на специальном упрощённом диалекте языка программирования Си алгоритмы, выполнимые на графических процессорах NVIDIA, и включать специальные функции в текст программы на Cи. CUDA даёт разработчику возможность по своему усмотрению организовывать доступ к набору инструкций графического ускорителя и управлять его памятью, организовывать на нём сложные параллельные вычисления.

ИСТОРИЯ

Направление вычислений эволюционирует от “централизованной обработки данных” на центральном процессоре до “совместной обработки” на CPU и GPU. Для реализации новой вычислительной парадигмы компания NVIDIA изобрела архитектуру параллельных вычислений CUDA, на данный момент представленную в графических процессорах GeForce, ION, Quadro и Tesla и обеспечивающую необходимую базу разработчикам ПО.

Говоря о потребительском рынке, стоит отметить, что почти все основные приложения для работы с видео уже оборудованы, либо будут оснащены поддержкой CUDA-ускорения, включая продукты от Elemental Technologies, MotionDSP и LoiLo.

Область научных исследований с большим энтузиазмом встретила технологию CUDA. К примеру, сейчас CUDA ускоряет AMBER, программу для моделирования молекулярной динамики, используемую более 60000 исследователями в академической среде и фармацевтическими компаниями по всему миру для сокращения сроков создания лекарственных препаратов.

На финансовом рынке компании Numerix и CompatibL анонсировали поддержку CUDA в новом приложении анализа риска контрагентов и достигли ускорения работы в 18 раз. Numerix используется почти 400 финансовыми институтами.

Показателем роста применения CUDA является также рост использования графических процессоров Tesla в GPU вычислениях. На данный момент более 700 GPU кластеров установлены по всему миру в компаниях из списка Fortune 500, таких как Schlumberger и Chevron в энергетическом секторе, а также BNP Paribas в секторе банковских услуг.

Благодаря ожидаемым в ближайшем будущем системам Microsoft Windows 7 и Apple Snow Leopard, вычисления на GPU займут свои позиции в секторе массовых решений. В этих новых операционных системах GPU предстанет не только графическим процессором, но также и универсальным процессором для параллельных вычислений, работающим с любым приложением.

Состав NVIDIA CUDA

CUDA включает два API: высокого уровня (CUDA Runtime API) и низкого (CUDA Driver API), хотя в одной программе одновременное использование обоих невозможно, нужно использовать или один или другой. Высокоуровневый работает “сверху” низкоуровневого, все вызовы runtime транслируются в простые инструкции, обрабатываемые низкоуровневым Driver API. Но даже “высокоуровневый” API предполагает знания об устройстве и работе видеочипов NVIDIA, слишком высокого уровня абстракции там нет. 

Есть и ещё один уровень, даже более высокий — две библиотеки:

CUBLAS — CUDA вариант BLAS (Basic Linear Algebra Subprograms), предназначенный для вычислений задач линейной алгебры и использующий прямой доступ к ресурсам GPU;

CUFFT — CUDA вариант библиотеки Fast Fourier Transform для расчёта быстрого преобразования Фурье, широко используемого при обработке сигналов. Поддерживаются следующие типы преобразований: complex-complex (C2C), real-complex (R2C) и complex-real (C2R).

Рассмотрим эти библиотеки подробнее. CUBLAS — это переведённые на язык CUDA стандартные алгоритмы линейной алгебры, на данный момент поддерживается только определённый набор основных функций CUBLAS. Библиотеку очень легко использовать: нужно создать матрицу и векторные объекты в памяти видеокарты, заполнить их данными, вызвать требуемые функции CUBLAS, и загрузить результаты из видеопамяти обратно в системную. CUBLAS содержит специальные функции для создания и уничтожения объектов в памяти GPU, а также для чтения и записи данных в эту память. Поддерживаемые функции BLAS: уровни 1, 2 и 3 для действительных чисел, уровень 1 CGEMM для комплексных. Уровень 1 — это векторно-векторные операции, уровень 2 — векторно-матричные операции, уровень 3 — матрично-матричные операции.

CUFFT — CUDA вариант функции быстрого преобразования Фурье — широко используемой и очень важной при анализе сигналов, фильтрации и т.п. CUFFT предоставляет простой интерфейс для эффективного вычисления FFT на видеочипах производства NVIDIA без необходимости в разработке собственного варианта FFT для GPU. CUDA вариант FFT поддерживает 1D, 2D, и 3D преобразования комплексных и действительных данных, пакетное исполнение для нескольких 1D трансформаций в параллели, размеры 2D и 3D трансформаций могут быть в пределах [2, 16384], для 1D поддерживается размер до 8 миллионов элементов.

Среда программирования

В состав CUDA входят runtime библиотеки:

Основной процесс приложения CUDA работает на универсальном процессоре (host), он запускает несколько копий процессов kernel на видеокарте. Код для CPU делает следующее: инициализирует GPU, распределяет память на видеокарте и системе, копирует константы в память видеокарты, запускает несколько копий процессов kernel на видеокарте, копирует полученный результат из видеопамяти, освобождает память и завершает работу.

В качестве примера для понимания приведем CPU код для сложения векторов, представленный в CUDA: 

Функции, исполняемые видеочипом, имеют следующие ограничения: отсутствует рекурсия, нет статических переменных внутри функций и переменного числа аргументов. Поддерживается два вида управления памятью: линейная память с доступом по 32-битным указателям, и CUDA-массивы с доступом только через функции текстурной выборки.

Программы на CUDA могут взаимодействовать с графическими API: для рендеринга данных, сгенерированных в программе, для считывания результатов рендеринга и их обработки средствами CUDA (например, при реализации фильтров постобработки). Для этого ресурсы графических API могут быть отображены (с получением адреса ресурса) в пространство глобальной памяти CUDA. Поддерживаются следующие типы ресурсов графических API: Buffer Objects (PBO / VBO) в OpenGL, вершинные буферы и текстуры (2D, 3D и кубические карты) Direct3D9.

Стадии компиляции CUDA-приложения: 

Файлы исходного кода на CUDA C компилируются при помощи программы NVCC, которая является оболочкой над другими инструментами, и вызывает их: cudacc, g++, cl и др. NVCC генерирует: код для центрального процессора, который компилируется вместе с остальными частями приложения, написанными на чистом Си, и объектный код PTX для видеочипа. Исполнимые файлы с кодом на CUDA в обязательном порядке требуют наличия библиотек CUDA runtime library (cudart) и CUDA core library (cuda).

Практические выводы

CUDA — это доступная каждому разработчику ПО технология, её может использовать любой программист, знающий язык Си. Придётся только привыкнуть к иной парадигме программирования, присущей параллельным вычислениям. Но если алгоритм в принципе хорошо распараллеливается, то изучение и затраты времени на программирование на CUDA вернутся в многократном размере.

Вполне вероятно, что в силу широкого распространения видеокарт в мире, развитие параллельных вычислений на GPU сильно повлияет на индустрию высокопроизводительных вычислений. Эти возможности уже вызвали большой интерес в научных кругах, да и не только в них. Ведь потенциальные возможности ускорения хорошо поддающихся распараллеливанию алгоритмов (на доступном аппаратном обеспечении, что не менее важно) сразу в десятки раз бывают не так часто.

Определение функций

Перед функциями в .cu файле могут стоять следующие “модификаторы”:

__device__ — это означает, что функция выполняется только на видеокарте. Из программы выполняющейся на обычном процессоре(host) её вызвать нельзя.

__global__ — Эта функция — начало вашего вычислительного ядра. Выполняется на видеокарте, но запускается только с хоста.

__host__ — Выполняется и запускается только с хоста (т.е. обычная функция C++). Если вы перед функцией укажите например __host__ и __device__ — то будут скомпилированы 2 версии функции (не все комбинации допустимы).

Определение данных

__device__ — означает что переменная находится в глобальной памяти видеокарты (т.е. которой там 512-1024Мб и выше). Очень медленная память по меркам вычислений на видеокарте(хоть и быстрее памяти центрального процессора в несколько раз), рекомендуется использовать как можно реже. В этой памяти данные сохраняются между вызовами разных вычислительных ядер. Данные сюда можно записывать и читать из host-части с помощью

cudaMemcpy(device_variable, host_variable, size, cudaMemcpyHostToDevice); //cudaMemcpyDeviceToHost - в обратную сторону

__constant__ — задает переменную в константной памяти. Следует обратить внимание, что значения для констант нужно загружать функцией cudaMemcpyToSymbol. Константы доступны из всех тредов, скорость работы сравнима с регистрами(когда в кеш попадает).

__shared__ — задает переменную в общей памяти блока тредов (т.е. и значение будет общее на всех). Тут нужно подходить с осторожностью — компилятор агрессивно оптимизирует доступ сюда(можно придушить модификатором volatile), можно получать race condition, нужно использовать __syncthreads(); чтобы данные гарантированно записались. Shared memory разделена на банки, и когда 2 потока одновременно пытаются обратиться к одному банку, возникает bank conflict и падает скорость.

Все локальные переменные которые вы определеили в ядре (__device__) — в регистрах, самая высокая скорость доступа.

Как поток узнает над чем ему работать

Допустим, надо сделать какую-то операцию над картинкой 200x200. Картинка разбивается на куски 10x10, и на каждый пиксел такого кусочка запускаем по потоку. Выглядить это будет так:

dim3 threads(10, 10);//размер квардатика, 10*10

dim3 grid(20, 20);//сколько квадратиков нужно чтобы покрыть все изображение

your_kernel<<<grid, threads>>>(image, 200,200);//Эта строка запустит 40'000 потоков (не одновременно, одновременно работать будет 200-2000 потоков примерно).

В отличии от Brook+ от AMD, где мы сразу определяем какому потоку над какими данными работать, в CUDA все не так: передаваеиые kernel-у параметры одинаковые для всех потоков, и поток должен сам получить данные для себя, чтобы сделать это, потоку нужно вычислить, в каком месте изображения он находится. В этом помогают магические переменные blockDim, blockIdx.

const int ix = blockDim.x * blockIdx.x + threadIdx.x;

const int iy = blockDim.y * blockIdx.y + threadIdx.y;

В ix и iy — координаты, с помощью которых можно получить исходные данные из массива image, и записать результат работы.

Оптимизация

Пару слов о том, как не сделать вашу программу очень медленной (написать программу работающую медленее чем CPU намного проще, чем работающую в 10 раз быстрее.

•Как можно меньше используйте __global__ память.

•При работе с __shared__ памятью избегайте конфлктов банков (впрочем многие задачи могут быть решены без shared памяти).

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

•Используйте как можно меньше памяти. Чем меньше памяти вы используете, тем агрессивнее компилятор и железо смогут запускать ваш kernel (например он может взять 100 тредов, и используя в 100 больше регистров запустить одновременно на одном MP, радикально уменьшая задержки)

Примеры Программ на CUDA C++

Исходные данные:

A[n][n] – матрица размерности n x n;

b[n] – вектор, состоящий из n элементов.

Результат:

c[n] – вектор из n элементов.

//Последовательный алгоритм умножения матрицы на вектор

for (i=0; i<n; i++)

{

c[i]=0;

for (j=0; j<m; j++)

{

c[i]=A[i][j]*b[j];

}

}

Теперь рассмотрим решение этой задачи на видеокарте. Следующий код иллюстрирует пример вызова функции CUDA:

// инициализация CUDA

if(!InitCUDA()) { return 0; }

int Size = 1000;

// обычные массивы в оперативной памяти

float *h_a,*h_b,*h_c;

h_a = new float[Size*Size];

h_b = new float[Size];

h_c = new float[Size];

for (int i=0;i<Size;i++) // инициализация массивов a и b

{

for (int k=0;k<Size;k++)

{

h_a[i*Size+k]=1;

}

h_b[i]=2;

}

// указатели на массивы в видеопамяти

float *d_a,*d_b,*d_c;

// выделение видеопамяти

cudaMalloc((void **)&d_a, sizeof(float)*Size*Size);

cudaMalloc((void **)&d_b, sizeof(float)*Size);

cudaMalloc((void **)&d_c, sizeof(float)*Size);

// копирование из оперативной памяти в видеопамять

CUDA_SAFE_CALL(cudaMemcpy(d_a, h_a, sizeof(float)*Size*Size,

cudaMemcpyHostToDevice) );

CUDA_SAFE_CALL(cudaMemcpy(d_b, h_b, sizeof(float)*Size,

cudaMemcpyHostToDevice) );

// установка количества блоков

dim3 grid((Size+255)/256, 1, 1);

// установка количества потоков в блоке

dim3 threads(256, 1, 1);

// вызов функции

MatrVectMul<<< grid, threads >>> (d_c, d_a, d_b,Size);

// копирование из видеопамяти в оперативную память

CUDA_SAFE_CALL(cudaMemcpy(h_c, d_c, sizeof(float)*Size,

cudaMemcpyDeviceToHost) );

// освобождение памяти

CUDA_SAFE_CALL(cudaFree(d_a));

CUDA_SAFE_CALL(cudaFree(d_b));

CUDA_SAFE_CALL(cudaFree(d_c));

В принципе структуру любой программы с использованием CUDA можно представить аналогично рассмотренному выше примеру. Таким образом, можно предложить следующую последовательность действий:

1. инициализация CUDA

2. выделение видеопамяти для хранения данных программы

3. копирование необходимых для работы функции данных из оперативной памяти в видеопамять

4. вызов функции CUDA

5. копирование возвращаемых данных из видеопамяти в оперативную

6. освобождение видеопамяти

Пример функции, исполнимой на видеокарте

extern "C" __global__ void MatrVectMul(float *d_c, float *d_a, float *d_b, int Size)

{

int i = blockIdx.x*blockDim.x+threadIdx.x;

int k;

d_c[i]=0;

for (k=0;k<Size;k++)

{

d_c[i]+=d_a[i*Size+k]*d_b[k];

}

}

Здесь:

threadIdx.x – идентификатор потока в блоке по координате x,

blockIdx.x – идентификатор блока в гриде по координате x,

blockDim.x – количество потоков в одном блоке.

Пока же следует запомнить, что таким образом получается уникальный идентификатор потока (в данном случае i), который можно использовать в программе, работая со всеми потоками как с одномерным массивом.

Важно помнить, что функция, предназначенная для исполнения на видеокарте, не должна обращать к оперативной памяти. Подобное обращение приведет к ошибке. Если необходимо работать с каким-либо объектом в оперативной памяти, предварительно его надо скопировать в видеопамять, и обращаться из функции CUDA к этой копии.

Среди основных особенностей CUDA следует отметить отсутствие поддержки двойной точности (типа double). Также для функций CUDA установлено максимальное время исполнения, отсутствует рекурсия, нельзя объявить функцию с переменным числом аргументов.

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

Для синхронизации потоков в блоке существует функция __syncthreads(), которая ждет, пока все запущенные потоки отработают до этой точки. Функция __syncthreads() необходима, когда данные, обрабатываемые одним потоком, затем используются другими потоками.

Замер времени в CUDA

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

unsigned int timer;

// создание таймера

cutCreateTimer(&timer);

// запуск таймера

cutStartTimer(timer);

… // код, время исполнения которого необходимо замерить

// остановка таймера

cutStopTimer( timer);

// получение времени

printf("time: %f (ms)\n", cutGetTimerValue(timer));

// удаление таймера

cutDeleteTimer( timer);

Заключение

Технология CUDA позволяет ощутимо увеличить производительность персональных компьютеров лишь в определенных задачах. Но сфера применения подобной технологии будет расширяться, а процесс наращивания количества ядер в обычных процессорах свидетельствует о росте востребованности параллельных многопоточных вычислений в современных программных приложениях. Не зря в последнее время все лидеры индустрии загорелись идеей объединения CPU и GPU в рамках одной унифицированной архитектуры (вспомнить хотябы разрекламированный AMD Fusion). Возможно CUDA это один из этапов в процессе данного объединения.

Источники:

  1. http://www.nvidia.ru/object/what_is_cuda_new_ru.html - Что такое CUDA
  2. http://www.overclockers.ua/software/nvidia-cuda/ - Знакомство с технологией CUDA
  3. http://fevt.ru/load/cuda/47-1-0-162 - Технология Nvidia CUDA (Л/Р) 2009
  4. http://www.ixbt.com/video3/cuda-1.shtml - NVIDIA CUDA — неграфические вычисления на графических процессорах
к оглавлению

Знаете ли Вы, что, как и всякая идолопоклонническая религия, релятивизм ложен в своей основе. Он противоречит фактам. Среди них такие:

1. Электромагнитная волна (в религиозной терминологии релятивизма - "свет") имеет строго постоянную скорость 300 тыс.км/с, абсурдно не отсчитываемую ни от чего. Реально ЭМ-волны имеют разную скорость в веществе (например, ~200 тыс км/с в стекле и ~3 млн. км/с в поверхностных слоях металлов, разную скорость в эфире (см. статью "Температура эфира и красные смещения"), разную скорость для разных частот (см. статью "О скорости ЭМ-волн")

2. В релятивизме "свет" есть мифическое явление само по себе, а не физическая волна, являющаяся волнением определенной физической среды. Релятивистский "свет" - это волнение ничего в ничем. У него нет среды-носителя колебаний.

3. В релятивизме возможны манипуляции со временем (замедление), поэтому там нарушаются основополагающие для любой науки принцип причинности и принцип строгой логичности. В релятивизме при скорости света время останавливается (поэтому в нем абсурдно говорить о частоте фотона). В релятивизме возможны такие насилия над разумом, как утверждение о взаимном превышении возраста близнецов, движущихся с субсветовой скоростью, и прочие издевательства над логикой, присущие любой религии.

4. В гравитационном релятивизме (ОТО) вопреки наблюдаемым фактам утверждается об угловом отклонении ЭМ-волн в пустом пространстве под действием гравитации. Однако астрономам известно, что свет от затменных двойных звезд не подвержен такому отклонению, а те "подтверждающие теорию Эйнштейна факты", которые якобы наблюдались А. Эддингтоном в 1919 году в отношении Солнца, являются фальсификацией. Подробнее читайте в FAQ по эфирной физике.

Bourabai Research Institution home page

Bourabai Research - Технологии XXI века Bourabai Research Institution