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

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

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

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

Cuda_input_buffer_stream – основной класс-поток по типу 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 [3] и cuda_input_buffer_stream показан на рисунке 1.

Рисунок 1 – Использование cuda_input_buffer_stream Для хост объекта определяется кастомизируемый метод operator<<(), осуществляющий запись каждого поля данного объекта в структуру типа cuda_input_buffer_stream.Для выбора корректной реализации и кастомизации оператора может использоваться как механизм разрешения перегрузок функций и/или шаблонов функций, так и механизм SFINAE (substitution failure is not an error – неудачная подстановка не является ошибкой) [1]. После записи всех полей объекта необходимо выделить память на устройстве и инициализировать данными хост объекта. Для этого был создан метод get_cuda_buf(), использующий методы cudaMalloc() и cudaMemcpy() для выделения памяти на устройстве и копирования данных. Затем объект класса cuda_input_buffer, удовлетворяющий требованиям StandardLayoutType [2], передается в качестве входного параметра в ядро. Чтение объекта на устройстве осуществляется с помощью метода 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_output_buffer_stream показан на рисунке 2.

Рисунок 2 – Использование cuda_output_buffer_stream

Для работы с выходным параметром необходимо на хосте в памяти устройства выделить память под структуру, которая в результате работы устройства 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.