CUDA.NET для .NET-разработчика (ч.3)

С точки зрения платформы .NET вплоть до недавно выпущенной 4-й версии фактически не было средств, облегчающих параллельную обработку данных. Однако ещё 4 июля 2008 года увидела свет версия 1.1 alpha библиотеки CUDA.NET. Она работала только под Windows XP, и поддерживала CUDA версии 1.1. С тех пор проект довольно динамично развивался, и на данный момент актуальной является релиз 3.0.0 - кроссплатформенный (в Linux и Mac OS использующий Mono), совместимый с CUDA версии 3.0, вышел он 17 июня 2010 года. Этого пакета вполне достаточно для изучения CUDA и создания каких-либо приложений на её основе: примером может служить программа для конвертации музыки в формат FLAC - FlaCuda. Фактически, CUDA.NET представляет собой объектно-ориентированную обёртку для использования API, предоставляемых драйвером CUDA.

Давайте попробуем написать простую программу, взаимодействующую с CUDA API. Качаем архив с библиотекой, внутри есть документация в виде CHM-файла и собственно сама сборка CUDA.NET.dll. Напишем небольшую консольную утилиту, отображающую сведения о CUDA-совместимых устройствах.

Чтобы начать, создадим в Visual Studio новый С#-проект - консольное приложение и сошлёмся на сборку CUDA.NET:

Основная точка взаимодействия с API драйвера CUDA - это класс GASS.CUDA.CUDA. Создадим объект этого класса:

CUDA cuda = new CUDA( true );

Параметр true конструктора позволяет инициализировать драйвер, "включая" механизмы CUDA. В общем-то дальше всё довольно просто:

Console.WriteLine();

Console.WriteLine( "Список CUDA-оборудования:" );

Console.WriteLine();

var devices = cuda.Devices;

Console.WriteLine( "Количество устройств: " + devices.Count() );

Console.WriteLine();

foreach ( var dev in devices )

{

  Console.WriteLine( "Название чипа видеоадаптера: " + dev.Name );

  Console.WriteLine( "Объём доступной видеопамяти: {0}Кб", dev.TotalMemory / 1024 );

  Console.WriteLine( "CUDA Compute Capability: " + dev.ComputeCapability );

  Console.WriteLine();

  Console.WriteLine( "Расширенные свойства:" );

  var props = dev.Properties;

  Console.WriteLine( "\tЧастота работы шейдерного домена: {0}МГц", props.ClockRate / 1000 );

  Console.WriteLine( "\tМаксимальное число нитей в блоке: " + props.MaxThreadsPerBlock );

  Console.WriteLine( "\tРегистров в мультипроцессоре: {0} шт.", props.RegistersPerBlock );

  Console.WriteLine( "\tОбъём общей памяти в мультипроцессоре: {0} байт(а)", props.SharedMemoryPerBlock );

  Console.WriteLine( "\tРазмер варпа: " + props.SIMDWidth );



  Console.WriteLine();

  Console.WriteLine( "Нажмите ENTER..." );

  Console.ReadLine();

}

}

catch ( CUDAException ex )

{

Console.ForegroundColor = ConsoleColor.Red;

Console.WriteLine( "CUDAException" );

Console.WriteLine( ex.Message );

}

ne();
}
}
catch ( CUDAException ex )
{
Console.ForegroundColor = ConsoleColor.Red;
Console.WriteLine( "CUDAException" );
Console.WriteLine( ex.Message );
}[/code]
Код утилиты очень прост: мы обращаемся к коллекции описаний устройств, и выводим информацию о каждом из них. Отмечу, что в целях упрощения здесь отображается не вся доступная из API информация. Вот какой результат я получил, запустив утилиту на своём компьютере:

У меня, очевидно, трудится GeForce 8800GT с 512 Мб памяти на борту (чип G92b), потоковые процессоры работают на частоте 1620МГц. Наверняка многим знакомы утилиты CPU-Z и GPU-Z, отображающие информацию о CPU и GPU соответственно. Имеется и open-source проект CUDA-Z - это программа, выводящая описание CUDA-"железа", доступного в данный момент. Она содержит любопытный встроенный бенчмарк, позволяющий оценить производительность видеокарты. Вот что я получил, запустив её на своём компьютере:

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

Пусть наша программа увеличивает на 42 все целые числа из входного массива элементов. Собственно инкремент будет производиться силами GPU. Этот пример довольно надуманный, зато он позволит "на пальцах" пояснить, как можно создать ядро, передать ему набор параметров и получить результат. Снова создадим консольное приложение, сделаем ссылку на сборку CUDA.NET, и добавим файл mykernel.cu в состав нашего проекта:

В этом файле объявим функцию-ядро следующим образом:

[code]extern "C" __global__ void test_func(int *g_data, int inc_value)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
g_data[idx] = g_data[idx] + inc_value;
}[/code]
Функция написана на CUDA С - обратите внимание на ключевое слово __global__ - оно обозначает, что функций является ядром (kernel). Наше ядро принимает набор параметров - массив в виде указателя на целое и величину инкремента. test_func будет увеличивать каждый элемент массива на заданную величину.

В первой строке тела функции мы вычисляем индекс текущего обрабатываемого элемента. Делается это так: исходная задача (grid) делится на определённое число блоков заданного размера. Блоки располагаются последовательно по цепочке. Значит, чтобы найти индекс очередного обрабатываемого элемента, достаточно индекс текущего блока (blockIdx) умножить на его размер (blockDim) и прибавить индекс текущей нити (threadIdx). Поскольку нити в блоке и грид блоков одномерные, используется только X-индексы. После этого во второй строке увеличиваем элемент массива на значение переданного аргумента.

Полученный *.cu-файл надо скомпилировать при помощи nvcc. Для этого в свойствах проекта в Post-build actions указываем следующую командную строку:

В командной строке nvcc указана опция "--cubin" - она означает, что на выходе мы получим бинарный файл mykernel.cubin, готовый к исполнению. Соответственно, компиляция ядра будет проходить только после успешной сборки хост-приложения. Уже сейчас можно попробовать собрать проект - если он собирается, значит всё ок, nvcc настроен должным образом.

Небольшой совет для пользователей 64-разрядных версий Windows: после установки CUDA Toolkit и SDK компиляция *.cu файлов может не проходить успешно, с сообщением "Cannot open include file: 'crtdefs.h': No such file or directory". Это легко исправить модификацией файла nvcc.profile, расположенный умолчанию по такому пути: C:\CUDA\bin64\nvcc.profile. Нужно поменять значение INCLUDE на следующее:

[code]INCLUDES += "-I$(TOP)/include" "-I$(TOP)/include/cudart" "-IC:/Program Files (x86)/Microsoft Visual Studio 9.0/VC/include" $(_SPACE_)[/code]
Все изменения сводятся к добавлению ещё одного пути, где nvcc будет искать заголовочные файлы. Собственно решение найдено на форуме NVidia.

Теперь пришло время заняться хост-приложением. Сначала инициализируем драйвер видеокарты, и создадим исходный массив:

[code]CUDA cuda = new CUDA( 0, true );
Int32[] array = Enumerable.Range( 0, 4096 ).ToArray();[/code]
Я сразу начал работу с устройством под номером 0 - это первая и единственная видеокарта GeForce 8800GT в моей системе. Входной массив содержит 4096 элементов - от 0 до 4095. Скопируем его в память видеоадаптера:

[code]CUdeviceptr d_input = cuda.CopyHostToDevice<Int32>( array );[/code]
API, по-моему, достаточно наглядный и в комментариях не нуждается. После этого загрузим подготовленную на стадии компиляции сборку mykernel.cubin:

[code]cuda.LoadModule( "mykernel.cubin" );[/code]
и получим дескриптор нашей функции:

[code]CUfunction func = cuda.GetModuleFunction( "test_func" );[/code]
после чего укажем параметры запуска:

[code]cuda.SetFunctionBlockShape( func, 512, 1, 1 );
cuda.SetParameter( func, 0, ( uint )d_input.Pointer );
cuda.SetParameter( func, IntPtr.Size, ( uint )42 );
cuda.SetParameterSize( func, ( uint )( IntPtr.Size + sizeof( uint ) ) );[/code]
Здесь стоит остановиться подробнее. Для начала определим размер блока - пусть это будет совокупность из 512 нитей - т.е. обработка входной последовательности будет производиться блоками по 512 элементов. Делается это при помощи вызова SetFunctionBlockShape, блоки будут одномерными, поэтому y- и z-размерности устанавливаем в 1. Далее передадим параметры нашей функции. Для этого передаём указатель на скопированный в DRAM GPU массив при помощи SetParameter, и тем же способом задаём приращение (пусть это будет 42). В конце следует указать общий размер параметров функции при помощи вызова SetParameterSize. После чего запускаем функцию на исполнение и копируем результат обратно в память хоста:

[code]cuda.Launch( func, array.Length / 512, 1 );
cuda.CopyDeviceToHost<int>( d_input, array );[/code]
При вызове метода Launch определим размеры грида. В нашем случае удобно сделать одномерный грид - фактически, просто цепочку блоков заданного размера. Параметры блока мы задали ранее, размер его - 512 нитей. Значит, размер грида будет равна числу частному от деления числа элементов во входном массиве (array.Length) на размер блока (512). Грид одномерный, поэтому y-размер его равен 1. Стоит отметить, что данный вызов метода является синхронным. Т.е. возврат из метода не будет произведён, пока наше ядро не обработает все элемента массива. Впрочем, библиотека CUDA.NET позволяет производить вызовы в асинхронном режиме. После того как функция-ядро успешно отработает, можно забрать результаты с устройства - сделаем это посредством вызова CopyDeviceToHost - скопируем результаты в тот же массив, который содержал исходные данные

Затем освободим ранее выделенную память (метод Free) и выведем на экран результат:

[code]cuda.Free( d_input );
foreach ( var item in array )
{
Console.WriteLine( item );
}[/code]
На этом всё! В результате ожидается увидеть на консоли ряд целых чисел от 0 до 4095, увеличенные на 42. Что и получаем:

Отлично! Значит, наше ядро работает правильно!

Как видим, использование CUDA в .NET приложениях является делом довольно простым. Библиотека CUDA.NET здесь приходится весьма кстати - впрочем, ничто не мешает написать собственную обёртку, если что-то в CUDA.NET не устраивает. В комплекте с CUDA.NET версии 3 примеров почему-то нет, однако в архиве с версией 2.3.7 библиотеки есть несколько интересных проектов - рекомендую ознакомиться. Я не стану приводить здесь более сложные примеры - для тех, кто заинтересовался темой и видит применение CUDA в своём реальном проекте, найти их не составит труда. Остальным же, надеюсь, этот небольшой экскурс в мир GPGPU был интересен.

Вместо углубления в тему неграфических вычислений на видеокартах, предлагаю рассмотреть поближе весьма занимательный продукт компании NVidia, имеющий непосредственное отношение к CUDA - имя ему NVIDIA Parallel Nsight. Это весьма интересная система, с помощью которой можно отлаживать, профилировать, различными способами анализировать CUDA C, OpenCL, DirectCompute, Direct3D и OpenGL приложения. Очень важным, на мой взгляд, моментом является интеграция всех этих средств прямо в Visual Studio 2008. На данный момент продукт находится в стадии RC, и распространяется среди закрытой группы пользователей. Впрочем, подать заявку на вступление в неё может любой желающий - всё здесь.

Parallel Nsight логически имеет разделение на 2 части - исполняемую на хосте и на целевой машине. Хост - это тот компьютер, с которого будет производиться отладка, а целевая машина - та, на которой находится CUDA-совместимая видеокарта. Деление чисто логическое, поскольку роль хоста и целевой машины вполне может играть и один компьютер, однако при этом на нём нужно иметь как минимум 2 GPU на чипе G92 или GT200. В случае разделения ролей по разным компьютерам на целевой машине также потребуется адаптер G92 или выше, а какая видеокарта будет на хосте - не важно. Обмен информацией будет идти по сети. Именно эту схему я и применял - на десктопной машине у меня располагается 8800GT (на чипе G92b), а отладку я веду с ноутбука с интегрированным в чипсет видео. Как видим, здесь присутствуют более жесткие ограничения - видеокарты на G80 (8500GT, 8600GT и др.) не поддерживаются. Хотя я не пробовал - возможно, заработает и на G80 :) Впрочем, наличие даже 2 карточек на G92 сейчас может позволить себе, наверное, любой заинтересованный разработчик.

Итак, у меня работает Nsight монитор на целевом компьютере, и я попробую запустить отладку примера CUDA-приложения. В этом примере производится умножение матриц. В настройках NSight User Properties я задам IP-адрес целевой машины (у меня это 192.168.1.248):

Затем соберу проект и в меню Visual Studio выберу NSight -> Start CUDA debugging:

После этого произойдёт соединение с целевой машиной, и в случае успеха получаем соответствующее уведомление:

Я заблаговременно поставил breakpoint в одном из *.cu-файлов, и в окне Locals вижу доступные сейчас переменные и могу проконтролировать их значения:

Видим привычный инструмент отладки, как будто мы исполняем код на обычном CPU! Можно посмотреть расширенные свойства CUDA-устройства:

А вот так выглядит окно Call Stack:

Это лишь часть, касающаяся CUDA, и только отладка. Как запустить в режиме отладки написанное нами ранее .NET приложение, я не нашёл - пока Nsight User Properties доступны только для Native C++ приложений. Однако, ничто не мешает написать и отладить функцию-ядро в native приложении, а потом лишь интегрировать её в managed-код - благо, разработку хоста и ядер можно вести в известной степени независимо друг от друга.

Да, и что касается цены Parallel Nsight - существует 2 версии: Standard (бесплатная) и Professional ($349 за рабочее место в первый год и $299 в последующие года). При этом Standard интегрируется в VS2008, умеет отлаживать CUDA-ядра и HLSL-шейдеры, поддерживает удалённую отладку. Профессиональная версия + к этому может профилировать ядра, включает тех. поддержку и имеет ещё ряд дополнений. Подробное сравнение версий можно найти тут. Небольшая подсказка: текущая RC-версия является Professional и истекает 1 сентября - т.е. у вас есть возможность при желании оценить все прелести полной версии Parallel Nsight!


Итог
Итак, мы рассмотрели подход компании NVidia к выполнению вычислений общего назначения на графических адаптерах - технологию CUDA. По моим личным ощущениям потенциал у технологии весьма хороший - дешевые потоковые процессоры таят в себе такие вычислительные мощности, которые несравненно больше предоставляемых любым современным CPU. С другой стороны, чтобы их утилизировать, нужно изначально адаптировать алгоритм к выполнению на GPGPU - большое количество задач в эти рамки не укладывается. Просто мало где требуется столь громадные объёмы вычислений, чтобы загрузить "под завязку" все SP. Накладные расходы, связанные с копированием данных из хоста в устройство и обратно, и др. при малых объёмах данных могут нивелировать преимущества от внедрения CUDA. Однако, по-моему стоит "держать руку на пульсе" и следить за развитием событий!

Существование CUDA.NET существенно упрощает жизнь .NET-разработчику, который решил применить CUDA в своём проекте - не нужно вручную оборачивать native библиотеки, всё уже готово. Ну а вместе с Parallel Nsight отладка и тестирование CUDA-функций становятся ещё проще.
И ещё, не думаю, что правильно противопоставлять CPU и GPGPU - это разные классы устройств. Стоит скорее рассматривать GPGPU как очень мощный "сопроцессор", буквально машину для выполнения тяжелых вычислений - и использовать оба этих устройства по назначению.

Спасибо за внимание, успехов в разработке! :)