SYCL 2020 — 用一份标准 C++ 让 GPU/CPU/加速器一起跑

待复核

是什么

SYCL（读作 “sickle”，镰刀）是 Khronos 给异构计算（GPU + CPU + FPGA + 加速器一起干活）定的一套纯标准 C++ 编程标准。

日常类比：以前你想让 CPU 干一段、GPU 干一段，得写两份代码——一份普通 C++ 给 CPU、一份特殊语言（CUDA / OpenCL）给 GPU，再手动搬数据。SYCL 说：“不行，太难维护。我让你一份 .cpp 文件、一个编译器，CPU 和 GPU 的代码写在一起，谁跑哪段我帮你切。”

最关键的一句：没有新关键字，写法就是 C++17/20 的库 + lambda。实现上编译器仍要多走一步，把设备代码提取出来特化（像把同一份菜谱拆成灶台版和烤箱版）。GPU kernel 是一个 lambda（匿名小函数），不是 __global__ void 这种 CUDA 扩展。

queue q;
q.parallel_for(1024, [=](id<1> i) {
    data[i] = data[i] * 2;
});

普通 C++ 编译器看着就是个 lambda；SYCL 编译器再把它编成 GPU/加速器代码。

为什么重要

不理解 SYCL，下面这些事都没法解释：

为什么 Intel 的 oneAPI / DPC++、AMD 侧的可移植栈、欧洲 exascale 超算项目都讨论它，而不是只押 CUDA
为什么”CUDA 想被替代”是个真命题——锁死 NVIDIA 硬件的代价越来越大
为什么 2020 年代的高性能计算（HPC）、AI 训练框架下层都在讨论”可移植 GPU 抽象”
为什么”单源 C++“是一个比”再发明一种语言”更聪明的工程选择

核心要点

SYCL 干的事可以拆成 三个支柱：

单源（single-source）：host 和 device 代码在同一个 .cpp 文件里。类比：一份菜谱同时写明灶台步骤和烤箱步骤。编译器扫两遍——先当普通 C++ 给 CPU，再把 kernel lambda 提取成设备代码。对比 OpenCL：host C++ + 一段字符串里的 OpenCL C，你得手动管两份。
三大抽象 = queue + buffer/accessor + parallel_for：
- queue：提交任务的入口，绑一个设备（像快递柜）
- buffer：托管数据容器，运行时决定何时在 host/device 之间拷（像自动行李传送）
- parallel_for：把一段 lambda 分发到 N 个工作项并行执行
后端可换：SYCL 2020 砍了对 OpenCL 后端的硬绑定。同一份代码可编到 NVIDIA（CUDA 后端）、AMD（HIP）、Intel（Level Zero，Intel 的底层驱动接口）、CPU（OpenMP）。这是它和 CUDA 最本质的差异。

加上 SYCL 2020 的 USM（Unified Shared Memory）——直接 malloc_device 拿指针、像 CUDA 那样用——入门四块就够。

实践案例

案例 1：单源 + lambda kernel 的最小例子

#include <sycl/sycl.hpp>
#include <vector>
using namespace sycl;

int main() {
    constexpr int N = 1024;
    std::vector<float> v(N, 1.0f);
    queue q;
    {
        buffer buf(v.data(), range{N});
        q.submit([&](handler& h) {
            accessor a(buf, h, read_write);
            h.parallel_for(N, [=](id<1> i) {
                a[i] = a[i] * 2.0f;
            });
        });
    }  // buffer 析构：等设备做完并把数据写回 host（RAII：对象销毁时自动收尾）
}

逐部分解释：

queue q：选默认设备，之后任务都往这里交
buffer + accessor：把 v 托管给运行时；kernel 里用 a[i]，不要直接摸 v
handler / parallel_for / id<1>：提交一次并行循环，每个工作项拿到自己的下标 i
大括号结束才保证写回——submit 本身不代表立刻同步

案例 2：USM 让 SYCL 写起来像 CUDA

constexpr int N = 1024;
std::vector<float> host(N, 1.0f);
queue q;
float* d = malloc_device<float>(N, q);
q.memcpy(d, host.data(), N * sizeof(float)).wait();
q.parallel_for(N, [=](id<1> i) { d[i] *= 2.0f; }).wait();
q.memcpy(host.data(), d, N * sizeof(float)).wait();
free(d, q);

逐部分解释：malloc_device 在设备上分配；两次 memcpy 搬数据；.wait() 自己卡同步。代价：失去 buffer/accessor 的自动依赖图（运行时根据谁读写同一块数据自动排队），得手写 .wait()。

案例 3：为什么”一份代码跑多硬件”是真的

icpx -fsycl main.cpp                                    # Intel GPU/CPU
icpx -fsycl -fsycl-targets=nvptx64-nvidia-cuda main.cpp # NVIDIA
acpp main.cpp                                           # AdaptiveCpp：AMD/NVIDIA/CPU

换工具链 + 目标参数，同一份源码可落到不同硬件；CUDA 二进制做不到这一点。

踩过的坑

kernel lambda 里不能用任意 C++：new / std::vector 内部分配、iostream、虚函数、异常都不行——GPU 没那些运行时。
buffer/accessor 的依赖图是隐式的：两次 submit 共用同一 buffer，SYCL 会自动串行；你以为并发就会调不动性能。
USM 的 device 指针不能在 host 解引用：malloc_device 在 host 上只是个地址数字，乱写会段错误；host 也要用就选 malloc_shared。
不同实现行为有差：DPC++、AdaptiveCpp 等都自称 SYCL 2020，扩展与性能差很大，“标准代码”不保证每家都最优。

适用 vs 不适用场景

适用：

HPC / 科学计算要一份代码跑多套硬件（多厂商集群、exascale 常见）
团队已是现代 C++ 栈，不想引入 CUDA 扩展语法
长期项目，担心被单一供应商锁死

不适用：

只用 NVIDIA 单卡、要挤尽每一滴性能 → CUDA + cuBLAS / cuDNN 仍最快
Python / AI 训练为主 → PyTorch / JAX 已包好异构，不必手写 SYCL
嵌入式、极简环境 → SYCL 运行时体积大，OpenCL 1.2 更轻

历史小故事（可跳过）

2014 年：Khronos 发布 SYCL 1.2，绑死 OpenCL 后端，工业界反响一般。
2017-2019 年：Intel 启动 oneAPI，把 SYCL 当核心；Codeplay 做出 ComputeCpp。
2020 年：SYCL 2020 临时草案在 IWOCL 公布——砍掉 OpenCL 硬绑定 + 加 USM。
2021 年 2 月：SYCL 2020 正式 ratify。
2022 年起：Aurora（Intel GPU）等推 oneAPI/SYCL；LUMI（AMD MI250X）应用层更常走 HIP，SYCL 是可移植选项之一。

学到什么

“单源”是关键工程选择——复用 C++/LLVM 成本低，比再发明 DSL 更划算
可移植 != 必然很慢——AdaptiveCpp 在部分常见内核上可接近原生 CUDA 的八成到九成五，仍视 workload 而定
委员会标准 vs 单家扩展——CUDA 快但锁定；SYCL 开放但生态更碎
USM 的退让说明：理论纯洁 < 用户体验——buffer 更安全，USM 让 CUDA 用户更快上手

关联

opencl-2010 —— OpenCL 是 host/device 双源，SYCL 把它升级成单源 C++
ampere-architecture-2020 —— SYCL 想抽象掉的硬件细节，Ampere 是其中一种
llvm —— DPC++ 和 AdaptiveCpp 都基于 LLVM/Clang 做 device 代码提取
cuda —— 单厂商对照物：快、生态深，但二进制锁 NVIDIA
thrust-2010 —— 同属”用 C++ 抽象并行”路线，Thrust 更偏 NVIDIA 算法库

反向链接

aurora-exascale-2024 —— Aurora 2024 — 不用 NVIDIA 也能造 2 EFLOPS 超算
cuda-streams-concurrency-2018 —— CUDA Streams 并发量化研究 — 为什么 SM 利用率拉不满
fpga-hls-2011 —— FPGA HLS 2011 — 把 C 代码自动翻译成芯片电路的范式
kokkos-2014 —— Kokkos — 一份 C++ 代码同时跑 CPU、GPU、Xeon Phi
nickolls-dally-2010-cuda-era —— Nickolls-Dally 2010 — GPU 怎么从画三角形变成跑 AI
owens-2007-gpgpu-survey —— Owens 2007 GPGPU 综述 — CUDA 之前 GPU 通用计算的黑魔法时代
thrust-2010 —— Thrust — 让 GPU 编程像写 STL 一样一行调用
filament —— Filament — Google 跨平台 PBR 引擎