SyCL ComputeCpp：如何在运行时同时支持SPIR与PTX比特码

一、现成实现方案 #

ComputeCpp官方支持通过多目标编译+运行时设备检测实现跨NVIDIA/AMD/Intel的适配，无需创建两个独立项目：

• 利用CMake的SYCL_TARGET_BACKEND编译选项，为同一内核源文件分别编译生成SPIR-V（适配AMD/Intel）和PTX（适配NVIDIA）格式的.sycl二进制文件
• 将多个.sycl文件作为资源嵌入可执行文件，或部署到程序可访问的目录
• 运行时通过cl::sycl::platform的信息接口（如get_info<cl::sycl::info::platform::name>()）识别当前GPU厂商，手动加载对应格式的.sycl二进制创建program

二、修改CMakeLists.txt的可行步骤 #

你提出的生成不同命名.sycl文件的思路完全可行，具体操作如下：

• 对含SyCL内核的源文件，添加两次add_custom_command：
  • 第一次指定-sycl-target spirv64，生成kernel_spirv.sycl
  • 第二次指定-sycl-target ptx64，生成kernel_ptx.sycl
• 通过configure_file或平台专属工具（如Windows的rc文件、Linux的objcopy）将两个.sycl文件打包到可执行文件中
• 在代码中通过平台检测逻辑选择加载对应的.sycl文件：

  cl::sycl::platform plat = cl::sycl::platform::get_platforms()[0];
  std::string plat_name = plat.get_info<cl::sycl::info::platform::name>();
  std::vector<char> binary_data;
  
  if (plat_name.find("NVIDIA") != std::string::npos) {
      // 加载kernel_ptx.sycl的二进制数据
  } else if (plat_name.find("AMD") != std::string::npos || plat_name.find("Intel") != std::string::npos) {
      // 加载kernel_spirv.sycl的二进制数据
  }
  
  cl::sycl::program prog(queue.get_context(), queue.get_device(), binary_data.data(), binary_data.size());
  prog.build();
  

三、更简便的替代方案 #

如果不想手动管理多个.sycl文件，可尝试ComputeCpp的Fat Program特性：将多个后端的二进制合并到同一个program对象中，运行时ComputeCpp会自动选择适配当前设备的二进制。注意：

• 需确保编译时同时指定多个-sycl-target参数（部分版本需额外配置）
• 合并后的program体积会增大，但无需手动编写设备检测逻辑

四、运行时热修补.sycl镜像的可行性 #

理论上可实现，但不推荐：

• .sycl是ComputeCpp封装的专用二进制格式，热修补需修改可执行文件的资源段或内存镜像，涉及平台相关底层操作
• 易引发稳定性问题，比如破坏二进制结构、触发系统安全机制（如ASLR、签名校验）
• 相比运行时动态选择加载的方案，实现复杂度高且维护成本大

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