SYCL lane 无法表达“SYCL + NVIDIA AOT”：架构集合只能挂在 cuda chunk 上 #887

写给：mcpp / mcpp-plugins 维护者
环境：Arch Linux x86_64，mcpp 2026.10.31.3，mcpp.plugins 0.19.1

TL;DR

--accel sycl 产出的设备代码是 SPIR-V，而 Unified Runtime 的 CUDA 后端不消费 SPIR-V，
所以在 NVIDIA 卡上必失败。规则自己的提示建议写 accel = "sycl, cuda12.9+{sm_89}"，
但这会打开 CUDA 后端轴：引擎随即要求编译所有 accel = "cuda" 的 .cu 设备源，
一个同时有 CUDA 和 SYCL 后端的工程就再也编不出纯 SYCL 构建。

期望的写法是让架构集合挂在 sycl chunk 上（sycl+{sm_89}），这样 accel 只点名
sycl 后端，却能驱动 -fsycl-targets=nvptx64-nvidia-cuda。我们验证了这个形态可以
工作，但需要两处改动：

1. rules/sycl.cppm 的 parse_target 从 sycl chunk 里读架构集合（下面是补丁）；
2. 插件的清单只在 all(accelerator = "sycl", accelerator = "cuda") 时声明
   xim:cuda-nvcc，所以项目还得自己声明这个 payload，否则规则报
   “the SYCL island needs payloads that are not installed”。

另外复现出两个独立问题：规则模块的源码改动不影响构建程序缓存；accel 的架构
集合不参与构建指纹（sycl 和 sycl+{sm_89} 共用同一个构建目录，会静默复用旧产物）。

同一台机器上，CMake + icpx 的 AOT 路径（-fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_89）是正常的，20 个设备测试全过——
所以问题不在适配器或硬件，而在 mcpp 规则/清单对这个形态的表达。

环境

工程里两个后端并存是常态：

[build]
sources = [
    { glob = "src/backend/cuda/device/**/*.cu",   accel = "cuda" },
    { glob = "src/backend/sycl/device/**/*.sycl", accel = "sycl" },
]

期望行为

mcpp test --features sycl --accel sycl --profile dev

在装了 CUDA 适配器的机器上直接选中 NVIDIA GPU；没装适配器时退回 CPU/其它设备。

实际行为

accel = "sycl" 让 tg.cuda_archs 为空，规则按 SPIR-V 编译。CUDA 后端拒绝 SPIR-V：

[ RUN      ] ...
No kernel named ... fill_kernel<float> was found

（在 ONEAPI_DEVICE_SELECTOR='cuda:*' 下最稳定复现；不带选择器时有时会变成 SYCL
调度器内部的 terminate called after throwing an instance of 'sycl::_V1::exception'，
程序自己接不住——规则里那段注释描述的就是这个。）

尝试 1：按提示写 --accel "sycl, cuda12.9+{sm_89}"

引擎把 cuda 当成激活的后端轴，于是要求 accel = "cuda" 的 glob 必须被某个
action 编译：

device sources that no action compiles: src/backend/cuda/device/*.cu

而在纯 SYCL 构建里调用 rules::cuda::compile() 也不可行：CUDA 规则的产物按
libc++ 编，SYCL 单元链接 libstdc++，最终是一串 undefined reference to std::__1::...。

尝试 2：把 glob 收窄成 cuda13 再写 cuda12.9+{sm_89}

约束不是过滤器，而是覆盖要求：

... constrained to accel "cuda13+{}", which this build does not cover

尝试 3（成功）：--accel "sycl+{sm_89}" + 规则补丁

根因

1. parse_target 丢掉 sycl chunk 的架构集合。原实现只从 cuda/hip chunk 读
   {...}，sycl chunk 只置一个 bool。于是“SYCL 后端 + nvptx 目标”在 accel
   语法里无处表达；能表达它的唯一写法 sycl, cuda... 又绑定了 CUDA 后端轴。
2. split(accel, ',') 不是大括号感知的。cuda12.9+{sm_89,sm_90} 会被切成
   cuda12.9+{sm_89 和 sm_90}；parse_arch_set 要求同一 chunk 里既有 { 又有
   }，于是多架构集合静默变成空集——这是同一函数里的第二个 bug。
3. 插件的 payload 声明条件过窄。mcpp.toml 里
   [target.'cfg(all(accelerator = "sycl", accelerator = "cuda", linux))'.feature-xlings.rules-sycl]
   才声明 xim:cuda-nvcc。sycl+{sm_89} 的后端集合仍然只有 {sycl}，规则里
   payload("cuda-nvcc") 返回空，尽管包就装在 store 里。

修复

1. 规则补丁（rules/sycl.cppm）

让 parse_target 按大括号外的逗号切分，并从 sycl chunk 读架构集合：

 inline target parse_target(std::string_view accel) {
     target t;
-    for (auto raw : split(accel, ',')) {
-        auto chunk = trim(raw);
-        if (chunk.starts_with("sycl")) t.sycl = true;
-        else if (chunk.starts_with("cuda")) t.cuda_archs = parse_arch_set(chunk);
-        else if (chunk.starts_with("hip"))  t.amd_archs  = parse_arch_set(chunk);
+    // 逗号也可能出现在 `{sm_89,sm_90}` 里，按大括号外的逗号切。
+    std::size_t begin = 0;
+    while (begin <= accel.size()) {
+        std::size_t end = accel.size();
+        for (std::size_t j = begin; j < accel.size(); ++j) {
+            if (accel[j] == '{') {
+                auto close = accel.find('}', j);
+                if (close == std::string_view::npos) break;
+                j = close;
+            } else if (accel[j] == ',') {
+                end = j;
+                break;
+            }
+        }
+        auto chunk = trim(accel.substr(begin, end - begin));
+        if (chunk.starts_with("sycl")) {
+            t.sycl = true;
+            // 架构集合也可以写在 sycl chunk 自己身上（`sycl+{sm_89}`）：
+            // accel 里不出现 cuda，引擎就不会打开 CUDA 后端轴。
+            if (auto archs = parse_arch_set(chunk); !archs.empty()) t.cuda_archs = archs;
+        }
+        else if (chunk.starts_with("cuda")) t.cuda_archs = parse_arch_set(chunk);
+        else if (chunk.starts_with("hip"))  t.amd_archs  = parse_arch_set(chunk);
+        if (end == accel.size()) break;
+        begin = end + 1;
+    }
     return t;
 }

2. 项目侧声明 nvcc payload（临时绕过第 3 条）

[target.'cfg(accelerator = "sycl")'.xlings.workspace]
"xim:cuda-nvcc" = "12.9.86"

更合理的长期修法大概是：清单增加一个在 sycl 轴、但 accel 带架构集合时才生效的
声明；或由规则在 tg.cuda_archs 非空时主动声明 payload。

3. 验证结果

mcpp test device --features sycl --accel "sycl+{sm_89}" --profile dev

• 生成的 build.ninja 里，每个 .sycl 对象和 -fsycl-link 步骤都带：
  -fsycl-targets=nvptx64-nvidia-cuda --cuda-path=<xim:cuda-nvcc> -Xsycl-target-backend --cuda-gpu-arch=sm_89
• 设备测试二进制从 29 MB 涨到 60 MB（内嵌设备镜像）。
• ONEAPI_DEVICE_SELECTOR='cuda:*' 下：20 passed / 4 skipped（需要两块 SYCL 设备）/ 0 failed。
• 对照组：同样的选择器跑 SPIR-V 产物，报 No kernel named ... was found。

补丁验证完后已把 rules/sycl.cppm 还原（sha256 与原始文件一致），并删除了实验用的
构建目录。

顺带发现的两个缓存问题

A. 修改插件源码不会让构建程序重编

改完 ~/.mcpp/xpkgs/mcpp-x-plugins/0.19.1/.../rules/sycl.cppm 后，mcpp build
仍然跑旧规则（表现为：补丁里加的调试 warning 不打印，产物也没有新的 flag），
Finished dev ... in 0.5s。缓存键（~/.mcpp/build-cache/v1/pkg/mcpp/mcpp.plugins@0.19.1/<key>/entry.json）
只含 config/toolchain/profile，不含源文件摘要。清掉下面几处才生效：

rm -rf target/.build-mcpp/host \
       target/.build-mcpp/build.mcpp.bin \
       target/.build-mcpp/build.mcpp.cache \
       target/.build-mcpp/build.mcpp.ctx \
       ~/.mcpp/build-cache/v1/pkg/mcpp/mcpp.plugins@0.19.1

如果这是有意为之（插件包按版本缓存），至少应该按包内容摘要参与键值——否则发行版/
用户给插件打补丁的行为会被静默忽略，排查成本很高。

B. accel 的架构集合不进构建指纹

--accel sycl 与 --accel "sycl+{sm_89}" 落到同一个
target/x86_64-linux-gnu/<hash>/ 目录。先做过 sycl 构建后再换 accel，
mcpp 认为“已是最新”，直接运行旧的 SPIR-V 产物（我们第一次实验就踩了这个坑：
测试通过/失败结果来自旧产物，而非新规则）。切换 accel 应至少改变构建指纹，
或者 accel 字符串本身进入变体目录的哈希。

复现步骤

1. 装有 xim:dpcpp、xim:cuda-nvcc，且系统有可用的 UR CUDA 适配器；
2. 工程按上文配置 .cu / .sycl 两组设备源；
3. mcpp test --features sycl --accel sycl --profile dev → SPIR-V，
   ONEAPI_DEVICE_SELECTOR='cuda:*' 下崩溃/找不到内核；
4. 打上第 1 节补丁、清掉 A 中的缓存、在 mcpp.toml 声明 nvcc；
5. mcpp test --features sycl --accel "sycl+{sm_89}" --profile dev → 20 passed / 4 skipped。

建议的最终形态

• accel 里 sycl chunk 支持 +{arch,...}，语义与 cuda chunk 相同（驱动
  -fsycl-targets=nvptx64-nvidia-cuda，不影响后端轴）；
• parse_target 按大括号外逗号切分，修掉多架构集合被截断的问题；
• 架构集合进入构建指纹与变体目录哈希；
• 插件模块的构建缓存键包含源文件摘要；
• 规则现有的 advisory 里补一句：accel = "sycl+{sm_89}"（打补丁后）是纯 SYCL
  工程在 NVIDIA 上的推荐写法。