跳转到内容
Jason's Study

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。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 在推的 ROCm 替代、欧洲 exascale 超算项目都选了它，而不是继续用 CUDA
  • 为什么”CUDA 想被替代”是个真命题——锁死 NVIDIA 硬件的代价越来越大
  • 为什么 2020 年代的高性能计算（HPC）、AI 训练框架下层都在讨论”可移植 GPU 抽象”
  • 为什么”单源 C++“是一个比”再发明一种语言”更聪明的工程选择

核心要点

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

  1. 单源（single-source）：host 和 device 代码在同一个 .cpp 文件里。SYCL 编译器扫两遍——第一遍当普通 C++ 编译给 CPU，第二遍把标了 kernel 的 lambda 提取出来、再编成 GPU/加速器代码。对比：OpenCL 是 host C++ + device 一段单独的字符串 OpenCL C，你得手动管两份。

  2. 三大抽象 = queue + buffer/accessor + parallel_for：

    • queue：你提交任务的入口，绑一个设备
    • buffer：托管的数据容器，SYCL 自己决定什么时候在 host/device 之间拷
    • parallel_for：把一段 lambda 分发到 N 个工作项并行执行

  3. 后端可换：SYCL 2020 砍了对 OpenCL 后端的硬绑定。同一份 SYCL 代码可以编到 NVIDIA GPU（用 CUDA 后端）、AMD GPU（用 HIP）、Intel GPU（用 Level Zero）、CPU（用 OpenMP）。这是它和 CUDA 最本质的差异。

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

实践案例

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

#include <sycl/sycl.hpp>
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;        // 这段会在 GPU 上跑
            });
        });
    }                                      // buffer 析构时把数据写回 host
    // v 现在是全 2.0
}

注意：整段是标准 C++。lambda、模板、RAII、构造析构都是 C++ 自带的。SYCL 编译器只是多读一次这个文件，把 lambda 拎出来特化。

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

queue q;
float* d = malloc_device<float>(N, q);   // 设备上分配
q.memcpy(d, host_ptr, N * sizeof(float)).wait();
q.parallel_for(N, [=](id<1> i) {
    d[i] = d[i] * 2.0f;
}).wait();
q.memcpy(host_ptr, d, N * sizeof(float)).wait();
free(d, q);

这就是 SYCL 2020 的让步——CUDA 用户已经习惯了 cudaMalloc / cudaMemcpy / kernel<<<>>>，SYCL 提供一套类似的指针式 API（USM）让迁移更轻松。代价：失去了 buffer/accessor 那套自动依赖图，得自己 .wait()。

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

Terminal window
# 同一份 main.cpp
icpx -fsycl main.cpp                                    # Intel GPU/CPU
icpx -fsycl -fsycl-targets=nvptx64-nvidia-cuda main.cpp # NVIDIA GPU
acpp main.cpp                                           # AdaptiveCpp，AMD/NVIDIA/CPU

切换工具链 + 一两个目标参数，二进制就跑在不同硬件上。这是 CUDA 做不到的（CUDA 二进制只能跑 NVIDIA）。

踩过的坑

  1. kernel lambda 里不能用任意 C++：动态内存（new / std::vector 内部分配）、iostream、虚函数、异常都不行。GPU 没那些东西。编译器一般会报错，但偶尔放过去运行时才崩。

  2. buffer/accessor 的依赖图是隐式的：两次 submit 都用同一个 buffer，SYCL 自动串行；如果你以为它们并发就大错。新人调性能时常被这个迷惑。

  3. USM 的 device-allocation 不能在 host 解引用：malloc_device 拿到的指针在 host 上只是个数字，写代码时编译器不查，跑起来段错误。要 host 也能用得用 malloc_shared。

  4. 不同实现行为有差：DPC++（Intel）、AdaptiveCpp（社区）、以前的 ComputeCpp 都自称 SYCL 2020，但扩展、调度器、性能差别很大。“标准 SYCL”代码不一定每家都最优。

适用 vs 不适用场景

适用：

  • HPC / 科学计算需要写一份代码跑多套硬件（exascale 超算就是这种场景）
  • 团队已经是现代 C++ 栈，不想引入 CUDA 这种”半 C 半扩展”的语法
  • 长期项目，担心被单一供应商锁死

不适用：

  • 只用 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。这两步让 CUDA 用户能迁移过来。
  • 2021 年 2 月：SYCL 2020 正式 ratify。
  • 2022 年起：美国 Aurora（Intel GPU）、欧洲 LUMI（AMD MI250X）等 exascale 超算的应用层主推 SYCL/oneAPI/HIP 三家可移植抽象。

学到什么

  1. “单源”是关键工程选择——比起再发明 DSL，复用 C++ 编译器框架（LLVM/Clang）成本低、生态可继承
  2. 可移植 != 性能最差——AdaptiveCpp 在 NVIDIA 上的性能能做到接近原生 CUDA 的 80-95%
  3. 委员会标准 vs 单家扩展——CUDA 的速度是优势，但锁定是劣势；SYCL 慢但开放
  4. USM 的退让说明：理论纯洁 < 用户体验。buffer/accessor 更安全但学习曲线陡，USM 难看但能让 CUDA 用户立刻上手

延伸阅读

关联

  • opencl-2010 —— OpenCL 是 host/device 双源，SYCL 把它升级成单源 C++
  • ampere-architecture-2020 —— SYCL 想抽象掉的硬件细节，Ampere 就是其中一种
  • llvm —— DPC++ 和 AdaptiveCpp 都基于 LLVM/Clang 实现 device 代码提取
  • hindley-milner —— 同样是”用编译器自动化”的思路，HM 自动推类型，SYCL 自动切代码