Как создать общую или статическую библиотеку с SYCL с помощью DPC++

Я пытаюсь создать общую библиотеку Linux, которую можно будет распространять и связывать, как любую обычную общую библиотеку. Недавно мы перенесли наши процедуры HPC GPU с CUDA на SYCL, чтобы обеспечить кросс-вендорный подход и создать единый код для GPU и CPU. До сих пор мы использовали DPC++.

Однако нам не удалось заставить DPC++ выполнить фактическую компиляцию перед этапом компоновки. В созданном файле .so отсутствуют все ядра, и его нельзя связать другим компилятором.

Ответ, найденный в этой теме от 2021 года, предполагает, что для создания окончательного двоичного файла следует использовать DPC++:

Создайте статическую или общую библиотеку из программы sycl с помощью dpc++

Однако это нежизнеспособное предложение по нескольким причинам:

1. Мы не можем требовать от наших конечных пользователей установки проприетарных пакетов компиляторов для связи с нашим кодом.
2. Это делает невозможным распространение и обновление общей библиотеки (нашего продукта) отдельно от программ, которые ее используют (продуктов наших пользователей).
3. Поскольку ядра не генерируются до момента компоновки, время компиляции в сотни раз замедляется даже для нашего собственного программного обеспечения (сотни небольших модульных тестов и прикладных программ, которые раньше компоновались за миллисекунды, теперь каждый из которых вызывает перекомпиляцию нашей огромной библиотеки, как .so — это всего лишь оболочка.

В общем, «общие библиотеки», которые мы можем создать с помощью DPC++, вообще не являются разделяемыми библиотеками.

Мы хотим просто сгенерировать фактический окончательный код (ядра хоста+устройства для выбора целей графического процессора и параллельного процессора) в общую библиотеку с API-интерфейсом только для хоста, который никоим образом не раскрывает SYCL (т. е. нет необходимости в компоновщик, чтобы узнать о SYCL). Можно ли создать общую библиотеку, содержащую код SYCL, но являющуюся реальной распространяемой общей библиотекой с полностью скомпилированными ядрами?

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

Есть ли способ создать настоящие общие библиотеки с помощью SYCL? Мы уже проделали большую работу по переносу нашего кода на SYCL, прежде чем обнаружили эту кирпичную стену, и нам не хотелось бы тратить все эти усилия на отказ от SYCL и перенос всего на третью модель программирования. Но сейчас это действительно похоже на остановку шоу.

Если невозможно создать распространяемую (разделяемую или статическую) библиотеку, содержащую ядра SYCL, с помощью DPC++, но это можно сделать с помощью другой реализации, такой как AdaptiveCpp, мне также очень интересно узнать, как это сделать.

Огромное спасибо за любую помощь, которую вы можете предложить!

Минимальный пример:

тест.хх:

void hello();

нормальная-lib.cc:

#include "test.hh"

void hello(){ }

sycl-lib.cc:

#include <sycl/sycl.hpp>
#include "test.hh"

void hello(){
    sycl::queue q;
   
    q.submit([](sycl::handler &cgh){
        cgh.single_task([](){
            // We don't need to do anything to trigger the failure,
            // it happens as soon as any SYCL kernel is defined. 
        });
    }).wait();
}

тест-main.cc:

#include "test.hh"

int main() {
    hello();
    return 0;
}

Как воспроизвести:

$ icpx -fPIC -fsycl -fsycl-unnamed-lambda -c sycl-lib.cc 
$ c++ -shared normal-lib.o -o libnormal-lib.so
$ c++ -o normal-test test-main.cc -L. -lnormal-lib
$ ./normal-test 
$ # Perfect, no problem here. Normal libraries work.
$ icpx -fPIC -fsycl -c sycl-lib.cc 
$ icpx -shared sycl-lib.o -o libsycl-lib.so
$ c++ -o sycl-test test-main.cc -L. -lsycl-lib -lsycl
$ ./sycl-test 2>&1 | c++filt
terminate called after throwing an instance of 'sycl::_V1::runtime_error'
  what():  No kernel named typeinfo name for hello()::{lambda(sycl::_V1::handler&)#1}::operator()(sycl::_V1::handler&) const::{lambda()#1} was found -46 (PI_ERROR_INVALID_KERNEL_NAME)
# Problem is here: Why was the kernel code not generated when building the shared library with "icpx -shared"?