Я пытаюсь оптимизировать мой симулятор, используя компиляцию во время выполнения. Мой код довольно длинный и сложный, но я определил определенную функцию __device__, чьи характеристики можно значительно улучшить, удалив все глобальные обращения к памяти.Функции NVRTC и __device__
Поддерживает ли CUDA динамическую компиляцию и связывание одной функции __device__ (а не __global__), чтобы «переопределить» существующую функцию?
Я уверен, что действительно короткий ответ - нет.
Несмотря на то, что у CUDA есть поддержка динамических/JIT-устройств, важно помнить, что сам процесс связи остается статическим.
Таким образом, вы не можете отложить загрузку определенной функции в существующей компилируемой полезной нагрузке GPU во время выполнения, как вы можете, в обычной среде загрузки динамической ссылки. И компоновщик по-прежнему требует, чтобы во время соединения присутствовал один экземпляр всех кодовых объектов и символов: a priori или во время выполнения. Таким образом, вы могли бы свободно связывать вместе предварительно скомпилированные объекты с разными версиями одного и того же кода, если только один экземпляр всего присутствует, когда сеанс завершен, и код загружается в контекст. Но это насколько можно.
Спасибо за ваш добрый ответ. Чтобы быть уверенным, вы бы так любезны проанализировать этот простой пример, который я нашел в Интернете? https://groups.google.com/forum/#!topic/thrust-users/GGWw_uHPfGs Он должен быть JIT компиляции, связывания и переопределение устройства функтор передается упорном –
я предполагал, что, когда вы сказали, «Разрешает ли CUDA ...», что вы говорили о части определенной/опубликованной функциональности. В приведенном ниже примере используются неопубликованные функции таким образом, которые могут поддерживаться или не поддерживаться в будущем. Если вы хотите создать свое программное обеспечение на этом, оно может работать, но оно также может сломаться. И если вы ищите документацию о том, как это сделать, вы ее не найдете. Процесс замены связанной функции другой связанной функцией во время выполнения таким образом также по существу ортогонален nvRTC. –
Хорошо, понял. Я буду придерживаться другого подхода (возможно, перекомпилировав весь бинарный файл при каждом изменении модели). Я все еще думаю, что эта функциональность будет классной, btw :) Спасибо за разъяснение! –
Похоже, у вас есть «основное» ядро с частью, которая «переключается» во время выполнения.
Вы можете определенно сделать это, используя nvrtc. Вам нужно будет делать что-то вроде этого:
Это примерно. Ключ заключается в том, чтобы отложить компиляцию основного ядра до тех пор, пока оно не понадобится во время выполнения. Вы также можете каким-то образом кэшировать свои ядра, чтобы вы в конечном итоге компилировались только один раз.
Есть одна проблема, которую я предвижу. nvrtc не может найти вызовы устройства curand, которые могут вызвать некоторые проблемы. Одна работа должна состоять в том, чтобы посмотреть на заголовок, в котором находится вызов функции устройства, и использовать nvcc для компиляции соответствующего ядра устройства в ptx. Вы можете сохранить полученный ptx как текст и использовать cuLinkAddData для связи с вашим модулем. Дополнительную информацию вы можете получить в this section.
Это ближе к тому, что я имел в виду. Итак, ключевым моментом является то, что я должен скомпилировать все ядро - вместе с новой сгенерированной функцией - во время выполнения. И да, вы были правы, CURAND в настоящее время является проблемой, и ваше решение, вероятно, является наиболее подходящим подходом. Спасибо! –
Согласен с ответом @talonmies. Замена функции, которая уже скомпилирована и связана с объектом, невозможна, я не думаю. Однако я думаю, что должно быть возможно загружать скомпилированное ядром устройство во время выполнения, которое имеет вызов функции '__device__', а затем выбирает и компилирует и связывает во время выполнения определенную функцию устройства, которую вы будете использовать для этого вызова. –
К сожалению, я не могу выбрать определенную функцию, потому что код функции будет сгенерирован в соответствии с моделируемой модели. То есть, загрузка разных моделей подразумевает различные уравнения, которые должны быть оценены в функции '__device__'. Вот почему я мог использовать компиляцию JIT, предоставляемую NVRTC. –
Вы можете выбрать функцию после загрузки модели, и вы знаете, какую функцию вы хотите использовать (или какую функцию вы только что создали). –