В этой статье делается попытка проиллюстрировать тот факт, что NVRTC + динамическое создание экземпляров может быть мощной парадигмой программирования CUDA для создания библиотек GPU, которые можно повторно использовать из произвольного языка, поддерживающего взаимодействие C/C++.
Это было мотивацией моего проекта с открытым исходным кодом ThrustRTC и ряда подобных проектов.
Важность и ограничения шаблонов в библиотеках GPU
Библиотеки с интенсивными вычислениями в основном написаны на C/C++ или другом скомпилированном языке. Интерпретируемые языки обычно также могут извлечь выгоду из этих библиотек за счет повторного использования библиотек.
Этот шаблон также применим к некоторым библиотекам GPU. Пока все функциональные возможности доступны через конкретные хост-API, эти библиотеки графических процессоров ничем не отличаются от других библиотек C/C++.
Однако есть исключения, известные как «библиотеки шаблонов». Эти библиотеки в основном представляют собой исходный код, который не компилируется при использовании. В программировании на GPU шаблоны особенно интересны и незаменимы из-за неэффективности динамического полиморфизма в моделях программирования на GPU. Одним из известных примеров такого рода библиотек является Thrust. Поскольку эти библиотеки предоставляются в нескомпилированном виде, их невозможно повторно использовать с языка, отличного от того, на котором эти библиотеки запрограммированы.
template<typename ForwardIterator , typename T >
void thrust::replace(ForwardIterator first,
ForwardIterator last,
const T & old_value,
const T & new_value)Это всего лишь простой пример функции, предоставляемой Thrust. Все параметры шаблонны. «T» может быть чем угодно, что имеет определение, а «ForwardIterator» должен быть «T*» или чем-то совместимым. Подобные функции могут быть очень мощными и полезными, но они доступны только для C++.
Компиляция во время выполнения и динамическое создание экземпляров
Компиляция во время выполнения не является чем-то новым для программистов GPU. В графическом программировании мы используем его для компиляции шейдеров для адаптации к различным средам выполнения. В OpenCL мы также используем его как способ компиляции кода устройства по умолчанию. Однако в программировании CUDA у нас не было компиляции во время выполнения до CUDA 7.x, новый модуль CUDA SDK называется NVRTC.
Теперь давайте посмотрим, как можно использовать компиляцию во время выполнения для создания библиотек столь же мощных, как и библиотеки шаблонов, но при этом их можно повторно использовать на другом языке.
Мощный, как шаблоны
Во-первых, нетрудно представить, что большинство проблем, которые могут быть решены с помощью шаблонного программирования, также могут быть решены путем автоматической модификации строк исходного кода. «Экземпляр» шаблона — это не более чем замена параметров шаблона конкретными типами и значениями.
Во-вторых, в большинстве случаев нам не нужно использовать шаблоны для кода хоста. Потому что динамический полиморфизм хорошо поддерживается в основном коде. Это просто проблема эффективности, и здесь мы пытаемся использовать GPU для интенсивной вычислительной части.
Следовательно, все, что нам нужно сделать, это запустить процесс «создания экземпляра» кода устройства во время выполнения, чтобы мы могли заставить его работать вместе с динамическим кодом хоста. Для этого требуется 2 части: одна — автоматическая модификация строк, другая — компиляция кода устройства во время выполнения.
Вот функция ThrustRTC, соответствующая функции traffic::replace:
bool TRTC_Replace(DVVectorLike& vec,
const DeviceViewable& old_value,
const DeviceViewable& new_value);DVVectorLike и DeviceViewable являются ведущими классами. Это абстрактные классы, и их подклассы могут содержать данные разных типов, а информация о типах записывается членами объектов времени выполнения.
Есть 2 функции интерфейса, которые реализует каждый объект, видимый на устройстве:
class DeviceViewable
{
public:
...
virtual std::string name_view_cls() const = 0;
virtual ViewBuf view() const = 0;
};Функция name_view_cls() возвращает, как этот объект представлен в коде устройства GPU, что должно быть распознано компилятором GPU. Строка будет задействована в процессах модификации строки при использовании объекта.
Функция view() возвращает байтовый буфер, содержащий данные, которые можно скопировать на устройство в переменную типа name_view_cls().
В библиотеке у нас может быть встроенный код (в виде строки), например:
template<T_Vec, T_Value>
extern "C" __global__
void replace(T_Vec view_vec,
T_Value old_value,
T_Value new_value)
{
uint32_t tid = threadIdx.x + blockIdx.x*blockDim.x;
if (tid>=view_vec.size()) return;
if (view_vec[idx] == (decltype(view_vec)::value_t)old_value)
view_vec[idx] = (decltype(view_vec)::value_t)new_value;
}Во время выполнения, когда мы пытаемся запустить ядро с некоторыми конкретными параметрами, его можно легко «создать», например:
template<class _T> struct VectorView { typedef _T value_t; typedef _T& ref_t;value_t* _data; size_t _size;__device__ size_t size() const { return _size; }__device__ ref_t operator [](size_t idx) { return _data[idx]; } };extern "C" __global__ void replace(VectorView<float> view_vec, float old_value, float new_value) { uint32_t tid = threadIdx.x + blockIdx.x*blockDim.x; if (tid>=view_vec.size()) return; if (view_vec[idx] == (decltype(view_vec)::value_t)old_value) view_vec[idx] = (decltype(view_vec)::value_t)new_value; }
Вы можете видеть, что здесь нам не нужно менять тело функции, а нужно изменить только заголовок функции.
Почему это портативно
Во-первых, сам компилятор не представляет собой ничего особенного как библиотека C, за исключением того, что он доступен только как общая библиотека, не знаю, почему NVIDIA не предоставляет статическую версию. Он используется вместе с API-интерфейсом драйвера CUDA, в котором тоже нет ничего особенного.
Во-вторых, в библиотечных API тоже нет ничего особенного. Все, что он предоставляет, — это скомпилированные API-интерфейсы C/C++.
Эти скомпилированные API-интерфейсы можно вызывать из любого языка программирования, поддерживающего C-interop. Такие классы, как «DeviceViewable», могут быть обернуты как классы на целевых языках. Для этого требуется, чтобы целевой язык поддерживал объектно-ориентированное программирование, что в настоящее время делает большинство из них.
Таким образом, вы можете найти очень похожие функции в версиях ThrustRTC для C++/Python/C#/JAVA:
// C++ bool TRTC_Replace(DVVectorLike& vec, const DeviceViewable& old_value, const DeviceViewable& new_value);# Python def Replace(vec, old_value, new_value): ...// C# using ThrustRTCLR;namespace ThrustRTCSharp { public partial class TRTC { ... public static bool Replace(DVVectorLike vec, DeviceViewable old_value, DeviceViewable new_value) ... } }// JAVA package JThrustRTC;public class TRTC { ... public static boolean Replace(DVVectorLike vec, DeviceViewable old_value, DeviceViewable new_value) ... }
Альтернативный способ создания библиотек CUDA?
Неудивительно, что большинство существующих библиотек CUDA основано на среде выполнения CUDA. Важной причиной является то, что это действительно очень удобно. Смешивание кода хоста и кода устройства в одном исходном файле на «одном» языке привлекательно и заставляет людей поверить, что знания программирования на C++ достаточно для того, чтобы программировать GPU. Среда выполнения CUDA + статическая компиляция + шаблоны также являются официально рекомендуемой парадигмой программирования CUDA. До появления NVRTC выбор, похоже, невелик. Даже после появления NVRTC люди все еще редко осознают, как можно изменить игру.
Здесь мы видим, что NVRTC + динамическое создание экземпляров может быть серьезной альтернативной парадигмой для программирования CUDA в целом. Он такой же мощный, как и шаблоны, и портативный, что важно для библиотек. Кроме того, это сокращает время компиляции библиотеки и генерирует очень тонкий двоичный файл. За это приходится платить тем, что приложение может работать довольно медленно при первом запуске (время компиляции перенесено сюда!).
Тот факт, что CUDA так широко используется в наши дни, делает ее обширной областью для изучения:
Как альтернативная парадигма соотносится со средой выполнения CUDA в каждой из областей применения CUDA?
На момент написания этой статьи. Я только закончил ThrustRTC, и мы впервые можем сравнить его с Thrust.
Следующие области, которые я собираюсь исследовать, включают:
- Матричные операции
- Нейронные сети
- Визуализация
Не удивляйтесь, если увидите библиотеки с именами BlasRTC, DNNRTC или VisRTC.
