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

Материал из CAMaaS preliminary wiki
Перейти к навигации Перейти к поиску

Контекст

Включение 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__ и __host__ - если не определен __CUDACC__. Таким же образом производим включение заголовочных файлов CUDA и Thrust.

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

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

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

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

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

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

Первый способ

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

  1. Генерируем матрицу D интенсивностей, которые надо сложить с существующей матрицей М. Код CUDA. На вход библиотеки CUDA передается источник PatternBasedSource, плоскость вывода результатов CPlain, коллекция отражаюших объектов CPoly.
  2. Генерируем вторичные источники (CPU);
  3. plain.mutex.lock() (CPU);
  4. Сложить: M += D. (CUDA/CPU);
  5. cudaStreamSynchronize (CUDA);
  6. plain.mutex.unlock() (CPU);
Классы, которые нужно реализовать для CUDA.

Второй способ

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

Модификация node-ов

Нужно реализовать Config::CudaCPUThreads().

По реализации

По автоматизации ввода-вывода см. статью Автоматизация ввода-вывода CUDA. Представленные ниже черновые наброски устарели, но пока оставлены на всякий случай.

Пусть X - некоторый класс, который реализуется для CUDA устройства и для CPU. Поскольку не всегда есть гарантия совместимости представлений класса для CPU и GPU ([1]), такие классы должны быть приводимы к формату, который можно передать ядру CUDA, с помощью метода X::kernel_data get_device_object(bool async = false, cudaStream_t = 0) const;, который должен возвращать объект вложенного типа X::kernel_data, удовлетворяющий требованиям StandardLayoutType. Этот метод, возможно перегруженный, должен 

struct X
{
    //...
    struct kernel_data
    {
        typedef X class_type;
        __device__ class_type get_object() const &;
        __device__ class_type get_object() &&;
        //...
    };
    kernel_data get_device_object(bool async = false, cudaStream_t = 0) const &;
    kernel_data get_device_object(bool async = false, cudaStream_t = 0) &&;
};

Вложенный тип X::kernel_data должен определять тип X::kernel_data::class_type, эквивалентный X, а также выполняемый устройством метод X::kernel_data::get_object без параметров, который возвращает восстановленное для устройства представление исходного экземпляра X.

Замена 

struct X
{
    //...
    //Тип для записи X в устройство
    struct kernel_data
    {
        typedef X class_type;
        __host__ kernel_data(const class_type&, bool async = false, cudaStream_t = 0);
        __host__ kernel_data(class_type&&, bool async = false, cudaStream_t = 0); //Опционально
        //...
    };
    __device__ X(kernel_data&&);

    //Тип для считывания X из устройства
    struct output_kernel_data
    {
        typedef X class_type;
        __device__ output_kernel_data(const class_type&);
        __device__ output_kernel_data(class_type&&); //Опционально
        //...
    };
    __host__ X(output_kernel_data&&, bool async = false, cudaStream_t = 0);
};

Если класс X_dev, заданный для устройства, отличен от соответствующего ему класса X_host, заданного для хоста, то...

TODO: Описать из CudaUtils.h


Пример
__global__ void my_kernel(bool fail, std::size_t host_thread_id, _Out_ unsigned* pVector)
{
    auto iThread = threadIdx.x + blockIdx.x * blockDim.x;
    if (fail && (iThread & 1))
        cuda_abort_with_error(-1, "CUDA error");
    pVector[iThread] = iThread;
}
void test(cudaStream_t stream)
{
    try
    {
        unsigned pVector[100];
        auto pDeviceVector = make_cuda_unique_ptr<unsigned>(sizeof(pVector));
        cuda_kernel_call(my_kernel_2, 10, 10, 0, stream, true, thread_id, pDeviceVector.get());
        cuda_runtime_call(cudaMemcpyAsync, pVector, pDeviceVector.get(), sizeof(pVector), cudaMemcpyDeviceToHost, stream);
        cuda_stream_synchronize(stream);
    }
    catch(cuda_user_exception& ex)
    {
        std::cerr << "Encountered an expected user-defined exception #" << std::hex << ex.code() << ": " << ex.what() << "\n";
    }
    catch(cuda_exception& ex)
    {
        std::cerr << "CUDA exception (code 0x" << std::hex << ex.code << ", description: " << ex.what() << ").\n";
    }
    catch (std::exception& ex)
    {
        std::cerr << "std::exception: " << ex.what() << "\n";
    }
}
Источник — http://51.250.0.107/w/index.php?title=Архитектурная_акустика/Реализация/CUDA&oldid=8058