如何在现行SYCL标准中编译并调用运行时生成的C代码到SYCL核函数？

针对你遇到的运行时生成C代码并在SYCL核函数中调用的问题，现行SYCL标准下可以通过以下几种方式解决（依赖主流SYCL实现的扩展接口，如DPC++、hipSYCL）：

方法一：用扩展接口编译运行时源码为内核Bundle #

主流SYCL实现提供了扩展接口替代已弃用的build_with_source，可以直接编译运行时生成的源码并创建可执行内核Bundle。以DPC++为例，使用sycl::ext::oneapi::experimental::compile接口：

1. 准备包含foo和SYCL核函数的源码字符串（如果需要复用foo，可以将其定义为设备端函数）：

// 读取运行时生成的foo.c内容，或者直接拼接成字符串
std::string source = R"(
extern "C"{
int foo(int a, int b) {
  return a + b;
}
}

// 定义可被SYCL调用的核函数，或者直接在SYCL lambda中引用foo
SYCL_EXTERN_C SYCL_EXTERNAL void compute_kernel(const int* in, int* out, size_t N) {
  size_t idx = sycl::this_item().get_id(0);
  if (idx < N) {
    out[idx] = foo(in[2*idx], in[2*idx+1]);
  }
}
)";

2. 编译源码为内核Bundle：

sycl::queue q;
// 编译源码，指定编译选项（如OpenCL版本）
auto kernel_bundle = sycl::ext::oneapi::experimental::compile(
  q.get_context(),
  {q.get_device()},
  source,
  "-cl-std=CL2.0"
);

3. 在submit中使用该Bundle并调用核函数：

q.submit([&](sycl::handler& h) {
  h.use_kernel_bundle(kernel_bundle);
  auto accessor_read = ...; // 你的输入Accessor
  auto accessor_write = ...; // 你的输出Accessor
  
  // 方式一：直接调用编译好的核函数
  h.parallel_for(sycl::range<1>(N_tasks), sycl::kernel_id("compute_kernel"),
    [=](sycl::id<1> id) {
      // 核函数逻辑已在编译的源码中，此处可直接映射参数
      compute_kernel(accessor_read.get_pointer(), accessor_write.get_pointer(), N_tasks);
    });
  
  // 方式二：在lambda中直接引用foo（需确保foo已被编译到Bundle）
  extern "C" int foo(int a, int b);
  h.parallel_for(sycl::range<1>(N_tasks), [=](sycl::id<1> id) {
    accessor_write[id] = foo(accessor_read[2*id], accessor_read[2*id+1]);
  });
});

方法二：编译运行时生成的代码为设备库，链接到SYCL程序 #

如果需要复用foo函数，可将其编译为设备目标文件，再链接到SYCL内核Bundle：

1. 将运行时生成的foo.c编译为设备目标文件（以DPC++为例，用icpx编译器）：

icpx -fsycl -c -o foo_device.o foo.c

（注：如果是运行时动态生成文件，可通过系统调用执行编译命令，或直接调用编译器API）

2. 读取目标文件二进制数据，创建库Bundle并链接为可执行Bundle：

// 读取foo_device.o的二进制内容
std::ifstream obj_file("foo_device.o", std::ios::binary);
std::vector<char> obj_data((std::istreambuf_iterator<char>(obj_file)), {});

sycl::queue q;
// 创建包含设备库的输入Bundle
auto lib_bundle = sycl::ext::oneapi::experimental::create_library_bundle(
  q.get_context(),
  {q.get_device()},
  {obj_data}
);

// 链接为可执行Bundle
auto exec_bundle = sycl::ext::oneapi::experimental::link(
  q.get_context(),
  {q.get_device()},
  {lib_bundle}
);

3. 在SYCL核函数中引用并调用foo：

q.submit([&](sycl::handler& h) {
  h.use_kernel_bundle(exec_bundle);
  extern "C" int foo(int a, int b); // 声明设备端函数
  
  auto accessor_read = ...;
  auto accessor_write = ...;
  
  h.parallel_for(sycl::range<1>(N_tasks), [=](sycl::id<1> id) {
    accessor_write[id] = foo(accessor_read[2*id], accessor_read[2*id+1]);
  });
});

方法三：内联简单函数到SYCL Lambda（适合逻辑简单的场景） #

如果foo的实现逻辑简单，可以直接将运行时生成的函数体嵌入到SYCL核函数的Lambda中，避免编译环节：

// 假设运行时生成的foo函数体是"return a + b;"
std::string foo_body = "return a + b;";

q.submit([&](sycl::handler& h) {
  auto accessor_read = ...;
  auto accessor_write = ...;
  
  h.parallel_for(sycl::range<1>(N_tasks), [=](sycl::id<1> id) {
    int a = accessor_read[2*id];
    int b = accessor_read[2*id+1];
    // 嵌入foo的实现
    accessor_write[id] = a + b;
  });
});

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