如何用DPC++构建含SYCL内核的可分发共享/静态库

问题核心 #

我们正将HPC GPU程序从CUDA迁移至SYCL（基于DPC编译器），目标是生成可独立分发、能被普通编译器链接的Linux共享库。但当前用DPC编译的.so文件缺失内核代码，导致运行时抛出内核未找到的错误。此前要求用DPC++链接最终二进制的方案不可行，原因如下：

• 无法要求终端用户安装专有编译器套件来链接代码；
• 无法独立于用户程序分发和更新共享库；
• 链接阶段生成内核会大幅增加自身软件的编译时间。

我们需要共享库包含主机代码及预编译的GPU/并行CPU目标内核，对外仅暴露主机API，无需用户链接器知晓SYCL细节。

DPC++解决方案 #

DPC++默认在链接阶段才实例化并编译SYCL内核，直接编译共享库时内核不会被嵌入。要解决此问题，需强制内核在编译共享库阶段就完成实例化与嵌入：

关键编译选项调整 #

1. 显式指定目标设备：用-fsycl-targets指定要预编译的设备，例如针对NVIDIA GPU和x86 CPU：

   icpx -fPIC -fsycl -fsycl-targets=nvptx64-nvidia-cuda,x86_64-unknown-linux-gnu -c sycl-lib.cc
   

2. 支持匿名Lambda内核：添加-fsycl-unnamed-lambda选项，兼容匿名Lambda形式的内核定义。
3. 链接共享库时保留内核：必须用DPC++（icpx）链接共享库，并带上SYCL相关选项，确保内核嵌入到.so中：

   icpx -shared -fsycl -fsycl-targets=nvptx64-nvidia-cuda,x86_64-unknown-linux-gnu sycl-lib.o -o libsycl-lib.so
   

代码调整确保内核实例化 #

若匿名Lambda仍存在问题，可将内核改为具名函数对象并显式实例化，强制编译器提前生成内核代码：

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

struct HelloKernel {
    void operator()() const {}
};

void hello(){
    sycl::queue q;
   
    q.submit([](sycl::handler &cgh){
        cgh.single_task<HelloKernel>(HelloKernel{});
    }).wait();
}

// 显式实例化内核，确保编译阶段生成
template void sycl::handler::single_task<HelloKernel>(const HelloKernel&);

用户用普通GCC/Clang链接时，只需链接生成的共享库和SYCL运行时库：

c++ -o sycl-test test-main.cc -L. -lsycl-lib -lsycl

AdaptiveCpp（原HIPSYCL）替代方案 #

AdaptiveCpp对共享库的支持更灵活，默认情况下只要编译时指定目标设备，内核就会被嵌入到共享库中：

编译命令 #

# 编译针对NVIDIA GPU和CPU的共享库
acpp -fPIC --hipsycl-targets=nvidia,gcc -c sycl-lib.cc
acpp -shared --hipsycl-targets=nvidia,gcc sycl-lib.o -o libsycl-lib.so

用户链接 #

用户可直接用普通编译器链接：

c++ -o sycl-test test-main.cc -L. -lsycl-lib -lhipsycl_rt

AdaptiveCpp的优势在于运行时更轻量化，对不同设备的支持更统一，无需用户安装Intel专有套件。

原方法失败原因 #

原命令中用icpx -shared sycl-lib.o -o libsycl-lib.so时，未添加-fsycl和目标设备选项，导致编译器未将内核编译并嵌入到共享库中。运行时SYCL runtime找不到对应的内核二进制，因此抛出PI_ERROR_INVALID_KERNEL_NAME错误。

复现示例 #

头文件 test.hh #

void hello();

普通库代码 normal-lib.cc #

#include "test.hh"

void hello(){ }

SYCL库代码 sycl-lib.cc #

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

struct HelloKernel {
    void operator()() const {}
};

void hello(){
    sycl::queue q;
   
    q.submit([](sycl::handler &cgh){
        cgh.single_task<HelloKernel>(HelloKernel{});
    }).wait();
}

template void sycl::handler::single_task<HelloKernel>(const HelloKernel&);

主程序 test-main.cc #

#include "test.hh"

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

正确编译步骤 #

# 编译SYCL库
icpx -fPIC -fsycl -fsycl-targets=nvptx64-nvidia-cuda,x86_64-unknown-linux-gnu -fsycl-unnamed-lambda -c sycl-lib.cc
# 生成共享库
icpx -shared -fsycl -fsycl-targets=nvptx64-nvidia-cuda,x86_64-unknown-linux-gnu sycl-lib.o -o libsycl-lib.so
# 用户编译主程序
c++ -o sycl-test test-main.cc -L. -lsycl-lib -lsycl
# 运行
./sycl-test

内容的提问来源于stack exchange，提问作者James