Функции NVRTC и __device__

Question

Я пытаюсь оптимизировать мой симулятор, используя компиляцию во время выполнения. Мой код довольно длинный и сложный, но я определил конкретный __device__ функция, производительность которой может быть значительно улучшена путем удаления всех глобальных обращений к памяти.

Позволяет ли CUDA динамическую компиляцию и связывание одного __device__ функция (не __global__), чтобы "переопределить" существующую функцию?

Answer 1

Я уверен, что действительно короткий ответ - нет.

Хотя в CUDA имеется поддержка компоновщика динамических устройств /JIT, важно помнить, что сам процесс связывания все еще остается статичным.

Таким образом, вы не можете откладывать загрузку определенной функции в существующую полезную нагрузку скомпилированного графического процессора во время выполнения, как в обычной среде загрузки динамических соединений. И компоновщик по-прежнему требует, чтобы один экземпляр всех объектов и символов кода присутствовал во время соединения, будь то априорно или во время выполнения. Таким образом, вы можете свободно связывать воедино предварительно скомпилированные объекты с разными версиями одного и того же кода, если только один экземпляр всего присутствует, когда сеанс завершается и код загружается в контекст. Но это так далеко, как вы можете пойти.

Answer 2

Похоже, у вас есть "основное" ядро ​​с частью, которая "переключается" во время выполнения.

Вы определенно можете сделать это с помощью nvrtc. Вам нужно будет сделать что-то вроде этого:

• Вместо того, чтобы компилировать основное ядро ​​заранее, сохраните его как строку для компиляции и связывания во время выполнения.
• Допустим, основное ядро ​​вызывает "myFunc", которое является ядром устройства, которое выбирается во время выполнения.
• Вы можете сгенерировать соответствующее ядро ​​"myFunc" на основе уравнений во время выполнения.
• Теперь вы можете создать программу nvrtc, используя несколько источников, используя nvrtcCreateProgram.

Вот и все. Ключ заключается в том, чтобы отложить компиляцию основного ядра, пока оно не понадобится во время выполнения. Вы также можете как-то кэшировать свои ядра, чтобы компилировать их только один раз.

Я предвижу одну проблему. nvrtc может не найти вызовы устройства curand, что может вызвать некоторые проблемы. Один из способов обойти это - посмотреть заголовок, в котором находится вызов функции устройства, и использовать nvcc для компиляции соответствующего ядра устройства в ptx. Вы можете сохранить полученный ptx как текст и использовать cuLinkAddData для связи с вашим модулем. Вы можете найти больше информации в этом разделе.