JIT-компиляция функций CUDA __device__

Question

У меня фиксированное ядро, и я хочу иметь возможность включать определенные пользователем функции устройства для изменения вывода. Пользовательские функции всегда будут иметь одинаковые входные аргументы и всегда будут выводить скалярное значение. Если бы я знал пользовательские функции во время компиляции, я мог бы просто передавать их в виде указателей ядру (и иметь функцию устройства по умолчанию, которая работает на входе, если не задана функция). У меня есть доступ к коду PTX пользовательской функции во время выполнения, и мне интересно, могу ли я использовать что-то вроде jitify от NVIDIA для компиляции PTX во время выполнения, получить указатель на функцию устройства, а затем передать эту функцию устройства предварительно скомпилированной функции ядра,

Я видел несколько сообщений, близких к ответу на этот вопрос ( Как генерировать, компилировать и запускать ядра CUDA во время выполнения ), но большинство предлагают компилировать все ядро ​​вместе с функцией устройства вво время выполнения. Учитывая, что функция устройства имеет фиксированные входы и выходы, я не вижу причин, по которым функция ядра не может быть скомпилирована заранее. Часть, которую мне не хватает, это как скомпилировать только функцию устройства во время выполнения и получить указатель на нее, чтобы затем передать ее функции ядра.

Answer 1

Вы можете сделать это, выполнив следующие действия:

1. Сгенерируйте свой проект cuda с помощью - сохраните и найдите сгенерированный ptx или cubin для вашего проекта cuda.
2. Во время выполнения сгенерируйте ваш ptx (в нашем эксперименте нам нужно было сохранить указатель функции в области памяти устройства, объявив глобальную переменную).
3. Создать новый модуль во время выполнения, начиная с cuLinkCreate , добавив сначала ptx или кубин из вывода --keep, а затем ваш ptx, сгенерированный во время выполнения, с помощью cuLinkAddData .
4. Наконец, вызовите ваше ядро. Но вам нужно вызывать ядро, используя только что сгенерированный модуль и не используя нотацию <<< >>>. В последнем случае это будет в модуле, где указатель функции неизвестен. Этот последний этап должен быть выполнен с использованием API драйвера (вы можете попробовать API времени выполнения cudaLaunchKernel).

Основной элемент - убедиться, что ядро ​​вызывается из сгенерированного модуля, а не измодуль, магически связанный с вашей программой.

Answer 2

У меня есть доступ к коду PTX пользовательской функции во время выполнения, и мне интересно, могу ли я использовать что-то вроде jitify от NVIDIA для компиляции PTX во время выполнения, получить указатель на функцию устройства, а затем передать это устройствофункция к скомпилированной функции ядра.

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

Вы можете выполнить связывание во время выполнения предварительно скомпилированного ядра (PTX или cubin) с функциями устройства, которые вы во время исполнения компилируете, используя NVRTC . Однако это можно сделать только через драйвер API модуля . Эта функциональность не предоставляется API времени выполнения (и, исходя из моего понимания того, как работает API времени выполнения, она, вероятно, не может быть представлена ​​без каких-либо серьезных архитектурных изменений в способе внедрения встроенного статически скомпилированного кода во время выполнения).