更好的排版请阅读: https://we5lw6jk7r.feishu.cn/wiki/GjQkwE6SciH2nhk3nERcwiC4nCq?from=from_copylink

SYCL+OneAPI+DPC++ 软件栈

教程

可视化的实例界面：https://tech.io/playgrounds/48226/introduction-to-sycl/sycl-kernel-programming—the-single_task-api

书籍：https://link.springer.com/book/10.1007/978-1-4842-5574-2

slide：
- https://0cch.com/uploads/2022/09/c++2022.pdf
- https://0cch.com/uploads/2022/11/sycl_coding_practice_and_compilation_design_analysis.pdf
Getting Started with oneAPI DPC++
- https://intel.github.io/llvm-docs/GetStartedGuide.html

背景

Intel i7 6700K，但是我们必须意识到其 GPU 的运算能力并不弱。从面积上来看，GPU 芯片的“占地面积”几乎和 CPU 平分秋色。也就是说，如果我们编写的程序只用到了 CPU 部分，那么是相当不划算的，因为大量“面积”的计算能力是没有用到的。
挑战：相同架构、指令集和语言的程序还比较容易写出来。那么不同架构，不同指令集，不同工具栏，不同语言的程序写起来就相当费劲了，这也是我们写异构程序时面临的重大挑战。

Intel oneAPI

oneAPI 就像操作系统的硬件抽象层，他让程序员不需要了解底层硬件有哪些，分别是什么。作为程序员，我们只需要知道我们的业务逻辑是什么，任务是什么，然后调用框架或者 oneAPI 接口即可。至于怎么和硬件交互，怎么把计算任务给到特定的加速器就不需要程序员关心了。当然了，既然 oneAPI 已经把硬件接管了，那么关于兼容性和移植性这些问题，也都是由 oneAPI 来处理的。

SYCL（更多的是一个规范）

介绍

SYCL 是一种用于编程异构系统的开放式行业标准。SYCL 的设计允许编写标准的 C++ 源代码，使其可以在异构设备或主机上运行。【一次编写，到处运行】
- SYCL 基于并行 API 或 OpenCL 等标准的基本概念、可移植性和效率，同时增加了单源 C++ 的易用性和灵活性。
- 使用 SYCL 的开发人员能够使用他们习惯的许多技术（如继承和模板）编写标准的现代 C++ 代码。同时，开发人员可以通过 SYCL 库的功能访问底层实现（如 OpenCL）的全部功能，必要时还可以通过 API 与直接使用底层实现编写的代码进行互操作。
SYCL 以几种方式扩展了 OpenCL 模型等标准中的概念，超越了 C++ 特性的一般用途【使用方式和 C++ 对齐】
- 支持使用 C++ 编写内核函数：在异构设备上执行并行内核既方便又灵活。常见的并行模式使用简单语法进行优先排序，通过一系列 C++ 类型，程序员可以在需要时表达额外的需求，如同步
- RAII 风格：当使用缓冲区和访问器时，SYCL 中的数据访问与数据存储分离。通过依赖 C++ 风格的资源获取即初始化（RAII）习惯用法来捕获设备代码块之间的数据依赖关系，运行时库可以跟踪数据移动并提供正确的行为，而无需手动管理内核实例之间的事件依赖关系，也无需程序员显式移动数据。这种方法使数据并行任务图（可能已经是执行模型的一部分）能够由 SYCL 程序员轻松安全地构建；
- 内存访问方案 1：统一共享内存（USM）：统一共享内存（USM）为显式数据分配和移动提供了一种机制。这种方法允许在异构设备上使用基于指针的算法和数据结构，并允许跨主机和设备增加代码的重用；

内存方案方案 2：缓冲器和访问器。将数据对象，这里是 vector，作为参数传给 buffer 的构造函数，用于构造 buffer 对象，然后在命令组范围内使用 buffer 构造 accessor，最后才能在内核函数中使用 accessor 去访问 vector 的数据。值得注意的是，buffer 必须使用一个作用域包括起来，因为 buffer 只有在析构的时候才会将加速设备的数据转移到主机内存中。

并行语法：分层并行语法提供了一种以易于理解的现代 C++ 形式表达数据并行性的方法，类似于 OpenCL 设备或 OpenMP 目标设备执行模型。它更清晰地分层并行循环和同步点，以避免代码碎片，并更有效地映射到 CPU 风格的架构。

实例

一个执行向量加法的内核，在主程序种使用 SYCL 队列和缓冲区将数据传输到 GPU 设备上运行
- 我们首先为要操作的数据设置主机存储。我们的目标是计算 c = a + b，其中变量是向量。
Host Code
Device Code
Queue：队列 queue 用于提交命令组(command group)到 SYCL 运行时执行，它是一种将工作(work)提交到设备的机制。一个队列可以映射(map)到一个设备，多个队列(multiple queue)可以映射到同一设备。queue 的成员函数 parallel_for 可以将命令提交到队列，不过这只是一种简化写法，更加完整的写法是调用 submit 函数，submit 函数将提交一个命令组 lambda 表达式，在表达式中命令组对象的成员函数再调用 parallel_for 函数，将内核函数提交给队列。
Selector：队列可以将命令提交给设备，但是具体提交给什么设备应该是可以指定的。所以 SYCL 提供了几种方法指定设备，一种是直接指定设备类别，比如 default_selector、cpu_selector、gpu_selector 等，这种选择器都是 SYCL 标准中预设好的。不过这种选择设备的方式也有局限性，因为如果同一类型存在多种不同的设备，上述方式就没有选择。

自定义：这个函数中，可以通过参数 device 来获取设备信息，比如类型和名称，然后通过这些信息选择目标设备。比如这里的代码示例，就是选择一个 Intel 的 GPU 设备并返回 100，数值越大优先级越高

#include <iostream>
#include <CL/sycl.hpp>

class vector_addition;

int main(int, char**) {
   // Host Storage
   // 我们首先为要操作的数据设置主机存储。我们的目标是计算 c = a + b，其中变量是向量。
   cl::sycl::float4 a = { 1.0, 2.0, 3.0, 4.0 };
   cl::sycl::float4 b = { 4.0, 3.0, 2.0, 1.0 };
   cl::sycl::float4 c = { 0.0, 0.0, 0.0, 0.0 };

   // 在 SYCL 中，有不同的方法来配置和选择我们想要使用的设备。 
   // SYCL 提供了一个默认选择器，尝试选择系统中最合适的设备。
   // 可以使用自定义选择器，但由于我们只有一台设备，所以我们使用默认选择器。
   cl::sycl::default_selector device_selector;

   // 为了发送要在目标设备上计划和执行的任务，我们需要使用 SYCL 队列。
   // 我们对其进行设置并向其传递我们的选择器，以便它知道在运行任务时选择什么设备。
   cl::sycl::queue queue(device_selector);
   std::cout << "Running on "
             << queue.get_device().get_info<cl::sycl::info::device::name>()
             << "\n";
   {
      // 发送到设备内存上
      cl::sycl::buffer<cl::sycl::float4, 1> a_sycl(&a, cl::sycl::range<1>(1));
      cl::sycl::buffer<cl::sycl::float4, 1> b_sycl(&b, cl::sycl::range<1>(1));
      cl::sycl::buffer<cl::sycl::float4, 1> c_sycl(&c, cl::sycl::range<1>(1));
  
      // 提交一个lambda函数，异步执行
      queue.submit([&] (cl::sycl::handler& cgh) {
         // 为数据设置权限 read, discard_write
         auto a_acc = a_sycl.get_access<cl::sycl::access::mode::read>(cgh);
         auto b_acc = b_sycl.get_access<cl::sycl::access::mode::read>(cgh);
         auto c_acc = c_sycl.get_access<cl::sycl::access::mode::discard_write>(cgh);

         // 定义了需要在设备上执行的内核函数
         cgh.single_task<class vector_addition>([=] () {
         c_acc[0] = a_acc[0] + b_acc[0];
         });
      });
   }
   std::cout << "  A { " << a.x() << ", " << a.y() << ", " << a.z() << ", " << a.w() << " }\n"
        << "+ B { " << b.x() << ", " << b.y() << ", " << b.z() << ", " << b.w() << " }\n"
        << "------------------\n"
        << "= C { " << c.x() << ", " << c.y() << ", " << c.z() << ", " << c.w() << " }"
        << std::endl;
   
   // SYCL 的特点之一是它使用 C++ RAII（资源获取即初始化）。
   // 这意味着没有显式清理，一切都是通过 SYCL 对象析构函数完成的。

   return 0;
}

DPC++（Intel 的一个开源?项目, 扩展了 SYCL）

代码放置到设备上执行

单源代码: 在 host 和 device 上执行的代码可以放在同一个源文件中

主机代码：管理数据和管理依赖关系
设备代码
- 异步执行：本质上就是不阻塞主机代码的执行方式
- 内存限制：不支持动态内存分配
- API 限制：部分函数和查询功能只能在设备代码中使用
动作 Action：把设备代码提交给设备队列的过程叫做 Action，这种动作不只是代码的执行提交，还包括内存数据的移动。【这个也太像 Legion 的核心目标了！】

C++ Lambda Functions in SYCL

[ capture-list ] ( params ) -> ret { body }

capture-list： SYCL uses [=] as kernel does not support capturing variable by reference
params： In SYCL, it can be unique 1D id, or 2D/3D id
ret：SYCL kernels must always have a void return type thus it is not specified in SYCL kernel
body： SYCL kernel body does not have any return statement

在任何设备上都可以运行

当我们不关心我们的设备代码将在哪里运行时，我们可以直接让运行时系统进行选择。
如果没有指定队列应该绑定的设备，就会在运行时选择可用的设备，SYCL 保证至少有一个设备总是可用的，即主机设备本身。

设备选择接口 selector

内置的设备选择接口
- default_selector
- host_selector
- cpu_selector
- gpu_selector
- accelerator_selector
- DPC++ 提供： ext::intel::fpga_selector
自定义选择器接口
- my_selector

#include <CL/sycl.hpp>
#include <sycl/ext/intel/fpga_extensions.hpp> // For fpga_selector
#include <iostream>
using namespace sycl;


int main() {
    queue my_gpu_queue( gpu_selector{} );
    queue my_fpga_queue( ext::intel::fpga_selector{} );
    
    std::cout << "Selected device 1: " <<my_gpu_queue.get_device().get_info<info::device::name>() << "\n";
    std::cout << "Selected device 2: " <<my_fpga_queue.get_device().get_info<info::device::name>() << "\n";
    
    return 0;
}

class my_selector : public device_selector {
public:
    int operator()(const device &dev) const override {
        if (
            dev.get_info<info::device::name>().find("Arria")!= std::string::npos &&
            dev.get_info<info::device::vendor>().find("Intel")!= std::string::npos
        ) {
        return 1;
    }
    
    return -1;
    }
};

任务流图调度

核心要素: Action+Dependence

命令组 Command Group = Action+Dependence+…

Command Group 包含的内容
- Action
- Dependence：当某个 kernel 的依赖关系被满足，kernel 就会被执行
- 其他主机代码
Command Group 会在 Host 上被执行，在提交到队列的时候执行，且只会执行一次。

Action

Device Code Execution（就是 Kernel）：通过 parallel_for 或 single_task 将工作提交到 device 的工作队列中
Explicit Memory Operation（显式内存操作）：USM 的 memcpy, memset, fill 操作、缓冲区的 copy, fill, update_host 操作

Dependence: 顺序队列和 wait+depends_on 和 accessor 隐式指定

顺序队列：左图
wait+depends_on，但是用 wait()和 depends_on()表示执行依赖：右图
accessor 隐式指定数据依赖
- Read After Write：B 任务需要依赖于 A 任务计算得到的结果，此时会出现隐式的数据依赖
- Write After Read：当 A 任务读取了某个数据，可能执行还没结束的时候，B 任务需要写这个数据，此时会有隐式的数据依赖
- Write After Write：当 A 任务在写某个数据的时候，B 任务也需要写某个数据

和 host 同步方案

方案 1：阻塞等待队列的完成 Q.wait() 和 Q.wait_and_throw()

方案 2：对事件进行同步，允许程序在特定的命令组上同步
方案 3：使用 host_accessor，确保被复制回 host 的数据是 device 计算完成后的数值
方案 4（少见）：使用缓冲区的属性 use_mutex 进行信号量同步

数据管理

三种管理方式
- 统一共享内存 Unified Shared Memory, USM：基于指针管理，支持 USM 的设备会有统一的虚拟地址空间，任何 host 上的 USM 动态内存分配函数返回的指针可以被 device 访问

缓冲区【Legion 的数据访问方式，但是没有 Legion 那么优秀】：代表一个或者多个 C++ 类型的对象，可以是标量数据（int、float、double），向量数据，用户定义的 struct。缓冲区变量不是具体的内存地址，需要用 accessor 对象进行访问

Images：专门用于图像处理的缓冲区管理
存储器的管理方式
- 显式管理
- 运行时隐式管理：由运行时系统自动完成，可以保证正确性。缺点：性能差

统一共享内存 Unified Shared Memory, USM

三种方式
- 设备分配：直接在 device 上分配空间，host 无法访问，如果 host 需要这些数据需要使用 USM 的 memcpy 复制过来
- 主机分配：在主机的内存中分配空间，host 和 device 都可以直接访问，但是不能直接迁移到 device 的存储器中，device 的访问是远程内进行的，通常需要经过 PCI-E 总线
- 共享分配：可以被 host 和 device 访问，自行进行数据迁移，但是自动数据迁移会有延迟
C 风格内存分配：分配完后得到 void*类型的数据，然后再进行强制类型转换
- malloc_device => aligned_alloc_device(用于返回对齐的内存指针)
- malloc_host => aligned_alloc_host
- malloc_shared => aligned_alloc_shared

C++ 风格内存分配：可以使用模板函数，返回的是对应类型的指针
- 支持 C++ allocator 进行内存释放

数据初始化
- Memset
- Fill
数据移动
- 显式数据移动：memcpy
- 隐式数据移动：

缓冲区

核心思想：本身只能代表数据，如何管理数据（存储和移动）都是运行时的工作
创建缓冲区
- 直接创建

从标量数组创建

从 vector 创建

使用缓冲区数据 Accessor：无法直接使用缓冲区的数据，必须定义对应数据的 accessor 才可以继续使用
- 为 Accessor 设置访问标记：read_write, read_only, write_only, read_constant

Kernel 的三种表达方式

基本数据并行 SPMD(单程序多数据流)

基本的数据并行 kernel 的功能是通过三个 C++ 类表示的：range、id 和 item
- range 表示一个一维、二维或三维的范围，维度需要在编译期确定，但每个维度的大小可以是动态的
- id 表示一个一维、二维或三维范围内的索引。尽管我们可以构造一个 id 来代表一个任意的索引，但为了获得与特定 kernel 实例相关的索引，我们必须将其作为 kernel 函数的一个参数。
- item 代表了一个 kernel 函数的单个实例，同时封装了其执行范围和实例在该范围内的索引。和 id 的主要区别在于 item 给出了额外的函数来查询范围与线性化的索引。

ND-range 数据并行

三种划分方式
- Work-group：表示一个 kernel 函数的实例，可以按任何顺序执行，除了对全局内存的原子内存操作外，不能相互通信或同步。每个 work-group 中的 work-item 可以访问 work-group 的本地内存，可以映射到一些设备上的本地存储器，可以用 work-group barriers 进行同步。[同一 work-group 内的 work-item 可以保证同时执行，不同 work-group 内的 work-item 不能保证同时执行]
- Sub-group：用于进一步的局部调度，可以使用编译器向量化的功能使得 sub-group 内的 work-item 并行执行，sub-group 没有自己的本地内存，可以用 shuffle 操作直接交换数据
- work-items：

分层数据并行(不太详细)

通信与同步: 针对 ND-range 数据并行

基于 Barrier 确保 work-group 内的 work-item 的同步

同一个 work-group 内的四个 work-items 通过 barrier 进行同步，尽管每个 work-item 的执行时间可能不同，但没有任何 work-item 可以跨过 Barrier 执行，直到所有 work-item 都遇到了 Barrier。执行 Barrier 功能之后，所有工作项就有了一致的内存。

基于 Work-group 的本地内存进行同步

通信可以通过 USM 或缓冲区进行，但可能效率不高，因而可以专门划分一部分内存用于通信，作为 work-group 的本地内存。
每个 work-group 都可以访问全局内存或自己的本地内存