Автоматизация ввода-вывода CUDA

Причина реализации: память под данные, передаваемые между хост и устройством, должна быть выделена на обоих сторонах. Однако это вызывает сложности в реализации взаимодействия с передачей комплексных объектов, которые имеют поля с динамической памятью, поскольку в CUDA не реализован механизм глубокого копирования. Для автоматизации выделения и инициализации памяти были созданы потоки cuda_input_buffer_stream для входных параметров (объекты, передающиеся с хоста на устройство) и cuda_output_buffer_stream для выходных параметров (объекты, передающиеся с устройства на хост).

cuda_input_buffer_stream – автоматизация передачи входных параметров

Диаграмма классов input_buffer_stream

cuda_input_buffer_stream – основной класс-поток по типу std::ostream, например, std::cout. Класс реализуется только для хост. В данный поток записываются некоторые данные obj с помощью operator<<().

get_cuda_buf() возвращает cuda_input_buffer, в котором данные представлены в бинарном виде.

cuda_input_buffer – буфер, содержащий данные в бинарном виде. Используется для трансфера данных. Метод release() используется для передачи владения внутренним буфером pBuf.

Метод data() возвращает его содержимое.

cuda_input_buffer_stream_d – реализация для девайса потока. Используется для распаковки данных, полученных с хоста.

Метод read_as позволяет из cuda_input_buffer_stream_d воссоздать класс SomeClass.

cuda_input_buffer_stream_d создается на основе буфера cuda_input_buffer.

Механизм передачи входного параметра с использованием потоков CUDA

Для хост объекта определяется кастомизируемый метод operator<<(), осуществляющий запись каждого поля данного объекта в структуру типа cuda_input_buffer_stream. Для выбора корректной реализации и кастомизации оператора может использоваться как механизм разрешения перегрузок функций и/или шаблонов функций, так и механизм SFINAE (substitution failure is not an error – неудачная подстановка не является ошибкой). После записи всех полей объекта необходимо выделить память на устройстве и инициализировать данными хост объекта. Для этого был создан метод get_cuda_buf(), использующий методы cudaMalloc() и cudaMemcpy() для выделения памяти на устройстве и копирования данных. Затем объект класса cuda_input_buffer, удовлетворяющий требованиям StandardLayoutType, передается в качестве входного параметра в ядро. Чтение объекта на устройстве осуществляется с помощью метода read_as(), который считывает структуру из потока и который кастомизируется для произвольных типов с помощью SFINAE.

Схема получение класса с хоста: cuda_input_buffer_stream -> cuda_input_buffer_stream::get_cuda_buf() -> cuda_input_buffer -> cuda_input_buffer_stream_d -> cuda_input_buffer_stream_d::read_as.

cuda_output_buffer_stream – автоматизация передачи выходных параметров

Диаграмма классов cuda_output_buffer_stream

cuda_output_buffer_stream – реализация для хоста выходного потока, используется для получения данных с устройства на хост. read_as() используется для создания из потока cuda_output_buffer_stream класса SomeClass.

cuda_output_buffer_d – реализация буфера для хост-девайс взаимодействия. Все методы (за исключением get_host_buf()) могут вызываться на стороне хост и девайс.

cuda_output_buffer – реализация выходного буфера для хоста (который может создаваться и на хосте, и на девайсе). Хост буфер создается из cuda_output_buffer_d.

cuda_output_buffer_stream_d – основной класс потока по типу std::cout для девайса. В данный поток записываются некоторые данные obj с помощью operator<<(). get_cuda_buf() возвращает cuda_input_buffer_d, в котором данные представлены в бинарном виде.

Механизм обработки выходного параметра с использованием потоков CUDA

Для работы с выходным параметром необходимо на хосте в памяти устройства выделить память под структуру, которая в результате работы устройства CUDA должна описывать данные выходного параметра. Созданный на устройстве выходной объект записывается в выходной поток cuda_output_buffer_stream_d с помощью кастомизируемого оператора «<<». Выходной поток агрегирует динамический буфер с данными выходного параметра. Описание этого буфера присваивается структуре, выделенной ранее на хосте. Для получения непосредственно данных на хосте создается хост версия выходного буфера и поток чтения, из которого будет формироваться хост-версия выходного параметра host_object кастомизируемым методом read_as().

Схема получение класса с девайса: cuda_output_buffer_d (предварительно созданный на хост) -> cuda_ouput_buffer_stream_d -> cuda_ouput_buffer_stream_d.get_cuda_buf() -> cuda_output_buffer_d -> cuda_output_buffer -> cuda_output_buffer_stream -> cuda_output_buffer_stream.read_as().

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

Использование средств автоматизации

Для использования автоматизации в классе SomeClass необходимо добавить флаг typedef std::true_type custom_cuda_device_read в device версию класса SomeClassCuda или специализировать cuda_specially_read для SomeClass.

Для записи из SomeClass необходимо определить перегрузку operator<<() для данного класса:

__host__ cuda_input_buffer_stream& operator<<(cuda_input_buffer_stream& is, const SomeClass& obj);

Далее необходимо записать каждое поле класса. Если тип члена класса динамическая структура (например, vector), то в поток необходимо предварительно записать количество элементов в контейнере. Для каждого составного типа необходимо определить свой оператор operator<<.

is << obj.container.size();

Для чтения данных из буфера и создания SomeClassCuda необходимо определить дружественный метод read_as:

template <class T>
__device__ auto read_as(cuda_input_buffer_stream_d& is) -> std::enable_if_t<std::is_same<SomeClassCuda, T>::value, T>

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

auto containerSize = read_as<std::size_t>(is);
for (std::size_t i = 0; i < containerSize; ++i)
		result.container.emplace_back(read_as<ComplexObject>(is));

Пример внедрения автоматизации

Use cases

Представленные варианты использования были использованы для создания средств автоматизации

• Передача динамического массива на устройство

Пример передачи динамического массива на устройство

• Передача динамического массива на хост

Пример передачи динамического массива на хост

• Передача вектора на устройство

Пример передачи вектора на устройство

• Передача вектора на хост

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

• Передача списка векторов на устройство

Пример передачи списка векторов на устройство

• Передача списка векторов на хост

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

Дополнительные классы

struct cuda_exception
{
	cuda_exception() = default;
	cuda_exception(cudaError_t err) :m_err(err) {}
private:
	cudaError_t m_err;
};

cuda_exception – класс-обертка для хоста, содержащий код ошибки cudaError_t.

struct cuda_input_buffer_datum
{
	virtual ~cuda_input_buffer_datum() {}
	virtual const std::uint8_t* data() const noexcept = 0;
	virtual std::uint8_t* data() noexcept = 0;
	virtual std::size_t size() const noexcept = 0;
};

cuda_input_buffer_datum – класс, реализующий буфер для хоста.

template <class DatumType>
struct cuda_input_buffer_datum_holder :cuda_input_buffer_datum
{
	alignas(DatumType) std::uint8_t datum[sizeof(DatumType)];
	template <class ArgType, class = std::enable_if_t<std::is_constructible<DatumType, ArgType&&>::value>>
	cuda_input_buffer_datum_holder(ArgType&& arg);
	virtual inline const std::uint8_t* data() const noexcept override final;
	virtual inline std::uint8_t* data() noexcept override final;
	virtual std::size_t size() const noexcept override final;
};

cuda_input_buffer_datum_holder – класс, реализующий буфер для хоста с выравниванием.

template <class T, class = std::true_type> struct cuda_specially_read :std::false_type {};

template <class T> struct cuda_specially_read<T, typename T::custom_cuda_device_read> :std::true_type {};

Метафункции для определения необходимости читать на cuda особым образом (на host и device могут существовать разные версии одного и того же класса).