Архитектурная акустика/Реализация/CUDA

Контекст

Включение CUDA-параллелизма условно. Зависит от:

• доступности CUDA-устройств и CUDA-драйвера. Проверяется наличием библиотеки cudart32_xx.dll/cudart64_xx.dll (xx - версия CUDA) (Windows) и libcudart.so/libcuda.so (Linux).
• версия CUDA cudaRuntimeGetVersion. Не менее 9.1 (возвращаемое значение 9010).
• доступности устройств с Compute Capability как минимум 3.0. См. ф-ции cudaDeviceGetAttribute и cudaGetDeviceCount.
• входного флага, включающего CUDA-параллелизм.

Поэтому реализацию CUDA нужно будет проводить в отдельной разделяемой библиотеке, которая загружается в адресное пространство arch_ac динамически (LoadLibrary/GetProcAddress в Windows и dlopen/dlsym в Linux). В свою очередь, библиотека с CUDA компонуется с cudart.lib (Windows) либо cuda и cudart (Linux) статически. Тогда отсутствие в клиентской системе CUDA приведет к детектируемым ошибкам функций LoadLibrary и dlopen. Все вызовы CUDA должны быть инкапсулированы в этой библиотеке, включая проверку версии CUDA.

Точка возможной загрузки библиотеки с CUDA - статический метод INodeImpl::Construct, на этапе выполнения выбирающий нужную реализацию вычислений в предметной области. К настоящему времени выбор реализации сводится к выбору специализации шаблона CNodeModel на основе целочисленного идентификатора расчетной модели, которая определяет алгоритм моделирования, но не способ реализации алгоритма. Поэтому шаблон CNodeModel нужно дополнить шаблонным параметром, задающим политику реализации выбранного алгоритма - с помощью CUDA или без. Более простым будет включение флага, вроде use_cuda в набор параметров, но более гибким - включение целочисленного идентификатора политики, возможно типа, наследующего std::integral_constant.

template <CalculationModel calc_id, class execution_policy>
class CNodeModel;

Код библиотеки во многом будет дублировать код arch_ac для CPU, но с учетом ограничений для CUDA C++, а также с учетом того, что интерфейс библиотеки не должен включать мета-кода с шаблонами. Использование нешаблонных классов возможно. Также возможно включение общих заголовочных файлов. Если заголовочные файлы содержат (мета) код, который имеет смысл использовать на устройстве CUDA, в начале такого файла условно определяем макросы __device__ и __code__ - если не определен __CUDACC__. Таким же образом производим включение заголовочных файлов CUDA и Thrust.

Регистрация источника на плоскости вывода результатов моделирования

Метод, который реализуется для CUDA: CNodeModel::ProcessSourcePlainInteraction (для всех расчетных моделей кроме CRayTracingSecondarySource), а также метод ProcessSource (для CKuliginRelaySecondarySource). Для узлов с CUDA данные методы являются точками входа в библиотеку CUDA.

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

Балансировка нагрузки между CUDA и CPU

Балансировка нагрузки сейчас осуществляется классом CLocalLoadBalancer, который уже включает в себя пул рабочих CPU-потоков, которым передаются поступающие задачи.

Существуют следующие варианты подачи нагрузки.

Первый способ - во всех рабочих CPU потоках, в реализации ProcessSource, асинхронно вызывать функции CUDA для регистрации источника на плоскостях вывода результатов, после чего создавать вторичные источники и, затем, синхронизировать выполнение с CUDA. Такой узел будет содержать в себе пустое определение CNodeModel::ProcessSourcePlainInteraction, т.к. регистрация будет выполняться в библиотеке CUDA, обращение к которой осуществляется из метода ProcessSource. Кроме того, каждый рабочий поток будет ассоциирован с выделенным для него потоком (stream) CUDA для передачи данных и команд на устройство. Этот stream логично создавать в локальной памяти потока (см. thread_local).

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