OneAPI 编程练习

C 语言课程中有要用 OneAPI 编程的作业，OneAPI 是一个异构计算的框架；而异构计算现在总归是在各种各样的地方都有奇怪的用处，所以就借此机会熟悉一下这个东西到底是什么样的（

作业代码参见： ymd45921 / oneapi-homework

开始之前

在使用 OneAPI 之前，首先我们需要知道一些与之相关的概念；不然编程练习也无从谈起：

1. SYCL (Standard C++ for Parallelism):
   • 定义： SYCL 是一种基于 C++ 11 以及未来版本的编程模型，是一个用于支持异构计算的抽象层，具有多个实现。它允许开发人员使用标准 C++（没有拓展任何新的语法，和 CUDA/CX/Qt 不同）代码来编写并行应用程序，并使其运行在包括主机 CPU 内的各种加速器设备（如 GPU、FPGA 等）上。
   • 关系： SYCL 建立在 OpenCL 的概念之上，允许使用 C++ 编写的代码在不同类型的计算设备上运行；虽然设计之初是为了和 OpenCL 和 SPIR 一起使用，但是实际上已经相当可以独当一面了。
2. OpenCL (Open Computing Language):
   • 定义： OpenCL 是 Khronos Group 提出的（怎么又是你？）一种开放标准，旨在支持异构计算。它使用 C 语言风格的编程模型，和 OpenGL 一样，它自身并不包含任何实现，而只是一个 API 框架。
   • 关系： OpenCL 为异构计算提供了底层的编程框架，可与其他编程模型（如 SYCL）结合使用。
3. CUDA (Compute Unified Device Architecture):
   • 定义： CUDA 是由 NVIDIA 开发的并行计算平台和编程模型，可以让开发人员在 NVIDIA GPU 上使用专有语言编写并行应用程序。
   • 关系： CUDA 是一种特定于 NVIDIA GPU 的技术，而 CUDA 程序只能在支持 CUDA 的 NVIDIA GPU 上运行。但是由于 N 卡的广为流传使得这几乎已经成为高性能计算中的事实标准
4. OneAPI:
   • 定义： OneAPI 是一个由英特尔推动的开放、综合的软件编程环境，旨在简化和统一异构计算。它提供了跨不同类型的硬件加速器（如 CPU、GPU、FPGA）的通用编程模型、常用的库以及一系列实用工具，这些都被包含在 Intel oneAPI Base Toolkit 中了。
   • 关系： OneAPI 可以使用多种编程语言（比如 DPC++、Fortran）编写代码，并在不同类型的平台和硬件上运行（和 CUDA 不同！）。这次作业用的就是 DPC++。
5. DPC++ (Data Parallel C++):
   • 定义： DPC++ 是一种编程语言，属于 OneAPI 的一部分，也是 SYCL 抽象的一个实现。DPC++ 的代码需要使用 Intel 全新的基于 LLVM 的编译器进行编译。
   • 关系： DPC++ 是 OneAPI 中的编程语言之一，是 SYCL 的实现，也是使用 OneAPI 的最直接的方式。

综上所述，在异构计算领域中，万能的 Khronos Group 有 OpenCL 这一低级标准和 SYCL 这一高级抽象，然而都没有实现；而比较广泛的异构计算的设备是 GPU，又因为老黄猖獗（不），所以 CUDA 成为了实际上异构计算常用的技术 —— 它是用于在 N 卡上进行高性能计算的专用编程模型/语言；为了打破这一局势，Intel 搞出了一个叫做 OneAPI 的东西，它包含了一个算是 SYCL 实现的语言 DPC++，一些常用的异构开发库，以及一系列用于调优和方便开发的实用工具（比如 CUDA 到 DPC++ 的迁移工具）。

SYCL 的基本概念

和 Vulkan 中主机和 GPU 一样，在异构计算中也有着这样的主机-设备模型：主机先选择要派发作业的设备，然后将已经转换成设备端的代码的核代码发送到设备的作业队列中，然后的作业就由设备完成，主机则去干干别的，直到设备完成作业通知主机进行下一步的操作；这之中关键的概念有 ——

1. Host（主机）：
   • 定义： 主机是计算系统的中央处理单元（CPU）或主要计算资源。在异构计算环境中，主机负责整体控制和协调计算任务。通常，主机执行应用程序的主要部分，包括调度任务、管理内存、处理输入输出等。
2. Device（设备）：
   • 定义： 设备是加速器，例如图形处理单元（GPU）、协处理器或其他专用硬件。设备通常用于执行并行计算任务，它们具有高度并行的架构，能够在处理大规模数据时提供显著的性能优势。
   • 关系： 设备是主机的扩展，主机通过将计算任务分派给设备来加速执行；异构计算中的设备可以是不同架构的硬件，如 GPU、FPGA 或其他加速器。
3. Kernel（核）：
   • 定义： 核是在设备上执行的并行计算任务的代码单元。通常，这些计算任务会被实现为一个或多个核函数，它们描述了如何在设备上执行并行计算。在 SYCL 的 API 中，还区分 Basic Parallel Kernel 和 ND-range Kernel。
   • 关系： 主机将计算任务划分为多个并行的核函数，然后将这些核函数分配给设备执行。核函数是针对设备架构优化的，并且通过并行执行来充分利用设备上的并行计算能力。
4. Device Selector（设备选择器）:
   • 定义： 设备选择器是 SYCL 中的一种机制，用于选择在运行时要使用的计算设备。它实际上是 Functor；SYCL 2020 之前，设备选择器需要继承自 sycl::device_selector 并重载 operator() 方法，该方法为作为参数传入的设备打分；在那之后，设备选择器可以是一个任意可调用物，作用依然是打分。
   • 作用： 可以通过 OneAPI 提供的设备选择器或者是用户自定义的设备选择器，在创建队列的时候选择一个用户所期望的设备。
5. Queue（队列）:
   • 定义： 队列是异构计算中常见的概念，表示要执行操作的执行上下文；队列被创建在一个设备上，或者说队列总是关联了一个设备；主机可以通过队列将需要设备执行的任务发送到设备，并控制设备和主机之间的数据传输。
   • 作用： 在 SYCL 中，所有的并行操作都在队列上执行。队列会将任务提交到设备，使得设备上的核函数能够执行，并在必要时与主机之间进行数据传输。

至于到了 SYCL 代码中，一个这样的模型就如下所示：

这段代码中，主机首先先通过 sycl::buffer 来管理内存，然后选择设备（这里省略了）创建命令队列，然后使用 submit 函数将任务提交到队列中；在提交任务的 Lambda 表达式中，首先通过 accessor 说明作业和数据的关联（这是作业相当重要的部分，DPC++ 编译器会根据这个来决定作业在设备上运行的顺序，来保证得到的结果是可预测的），然后再通过 parallel_for 函数发起并行 —— 这里的 Lambda 表达式就是核函数了，它会像 Shader 那样被加速设备的每个单元并行执行以加速计算。

正因为有明确的 Host-executed Scope 和 Device-executed Scope，才使得这两段代码出现在一个文件里成为了可能；但是在编译阶段，编译器仍然会将这两部分代码分开处理，并且根据 Command 的信息对在设备上的运行进行优化以尽可能提高并行度 —— 这也自然的导致了，在设备上执行的核代码的 C++ 必然是受到了限制的；毕竟加速硬件大规模的代价就是泛用度的降低嘛。

因为涉及到了作业的递交，所以就也会出现主机端的数据和设备端的数据的问题；对于 SYCL 而言，主机端的数据有两种模型：一是上面用到的更接近数组声明的 Buffer & Accessor 模型 —— Buffer 将持有内存的管理权，而 Accessor 声明作业对内存的访问权（OneAPI 大概会在设备上开辟一片这样的空间将数据传输过去？不过这里的细节我就不太了解了）—— 但是一旦一块主机内存交给 Buffer 管理（比如你声明的数组，分配的内存），即使是主机访问也需要通过声明 sycl::host_accessor，否则程序的行为是 UB；二是看起来比较简单且更接近 C 语言指针分配内存的 USM (Unified Shared Memory) 模型，这种情况下首先通过 sycl::malloc_shared API 分配主机和设备都能访问的内存，因此用户需要保证主机和设备对数据的分开访问 —— 其实也很简单，就是等待作业完成后主机再访问这片内存就好了。

更多的像是 barrier 啊、fence 啊这些同步机制 SYCL 当然也是有提供的，但是这也要说那也要说博客就写不完了所以就算了。

异构计算 vs OpenMP

OpenMP 也是一个挺常用的并行计算的东西，和 DPC++ 这种需要 icpx 支持的东西不同，几乎所有的常见/系统自带编译器都支持这个玩意；因为 Intel DevCloud 给我的服务器只有 CPU 和 FPGA 模拟器，而根本没有什么 GPU，我还思考了这个问题一段时间，这里给出解释——

• OpenMP 定义： OpenMP 是一种用于共享内存多处理系统的并行编程 API。用户可以在循环、函数等作用域上标记 #pragma，以指示编译器在运行时并行执行这些代码。

简而言之，OpenMP 只支持在 Host 上并行运行代码，只应对了并行编程中共享内存并行编程这一领域；而异构计算则是不只要在主机上跑，还要在主机之外的计算平台上跑 —— 就算是用 OpenGL 这些，在 GPU 上跑的代码都得用 Shader Language 写完了编译完了再给它送过去，而 SYCL 代码两边都是 C++！只是在设备上跑的部分会受到限制而已，确实是可以更为方便地编写高性能计算！

SYCL vs Vulkan

Vulkan 和 SYCL、OpenCL 一样，也是属于 Khronos Group 的标准；不过主要是用与图形渲染的，不过现在的图形 API 中都会有 Compute Shader 这个概念，所以就会让我这种图形学学徒产生一种万能感 —— 那我为什么不去用美丽的 Vulkan 呢？想了想如果有原因的话大概是这样：

• 实现 Vulkan API 的主要是 GPU，而计算平台不只是 GPU，还包括 FPGA 和核弹级 CPU；这就导致了 Vulkan 的 Compute Shader 只能在一部分设备上运行，不如专用的异构计算框架来的广泛，移植性更好。
• SYCL 写起来的手感比 Compute Shader 更爽：首先，SYCL 的核代码也是标准 C++，甚至还和主机代码写在一起，分派到设备的转化都是编译器的事情；而且 SYCL 提供的抽象更高级，大量的优化都是交给编译器编译期进行，即方便开发也不会拖累运行时速度；而 Vulkan 本身就是个低级 API，什么都要手动指定就不说了，你觉得用 Shader Language 写一个 Computer Shader 再编译成 SPIR-V 很炫酷嘛？

综上所述，Vulkan 也可以用于异构计算，但是它的异构仅限于使用支持的 GPU；尽管 GPU 是异构计算中最为常见的设备，但是它代表不了异构计算的全部！所以我的评价是，用专门的库干专门的事，不要搞七搞八（）虽然 SYCL 还是有不少过程和 Vulkan 蛮像的，但是人家毕竟还是高级抽象层！不能比！

动手之前

然而，已经被 Vulkan 配环境气晕的我一点也不像再在电脑上配置一个新的环境；那这种情况呢就可以使用 Intel DevCloud 服务：这样吧，首先在浏览器的地址栏中输入 https://devcloud.intel.com/oneapi/get_started/，然后点击登陆，注册一个新的账号就可以获得 Intel DevCloud 的几个月的访问权了！

使用 Intel Devcloud

获得访问权限之后，就可以根据指示配置电脑上的 SSH 了；Windows 需要手动而 Linux/Mac 有脚本可以用；总之最后 Intel 会在 SSH 配置中增加它的登录节点、计算节点以及 VSCode 服务器；登陆节点用于提交任务到计算节点，并查看任务运行的状态，并不能运行很多东西，不然就会：

而要开启一个计算节点，则需要根据官网中所说先在登陆节点中使用 qsub sleep.sh 提交一个睡眠任务：

1
2
3
4
5
6
7
8

#!/bin/bash

#PBS -q batch@v-qsvr-1
#PBS -l select=1:batch:ppn=2
#PBS -l walltime=04:00:00
#PBS -d .

sleep 14400

然后通过 qstat -n -1 命令就可以看到为我们分配的计算节点的名字：

图中选中的部分就是计算节点的 ID，可以看到我运行了两个睡眠任务（两个脚本只是名字不一样），一个的状态是 Q 说明还在排队，而另一个是 R 说明已经分配了计算节点并且正在运行。然后因为 Intel 已经配置好了我们的 SSH 客户端，所以我们可以直接通过 ssh <node-id>.aidevcloud 就可以连接到这个计算节点，保持这个连接打开的情况下，就可以通过 VSCode 的 Remote-SSH 插件连接到 devcloud-vscode 了：

用服务器开发那确实方便啊…… 可是这种先让服务器睡眠，然后再用终端骚扰计算节点的方法，怎么那么像当年那会 Github Actions 刚出来的时候一堆“聪明老哥”嫖 Github 服务器的样子呢？（不）

配置项目

因为只能用 VSCode 远程开发，所以显然是没得 CLion 用了。习惯来说我会使用 CMake 管理项目，可是 VSCode 的 C++ 插件一直都挺奇怪的，似乎不是很和 CMake 联动…… 解决方案是让 CMake 构建的时候输出编译命令，这样就会被 VSCode 的 C++ 插件读取到正确的配置了。

1

set(CMAKE_EXPORT_COMPILE_COMMANDS ON)

此外，还需要强调使用 icpx 作为编译器，并且需要显式标记使用 sycl 库；因为编译器已经锁定，所以所有的编译选项只要直接加入编译命令中就好了。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16

if(UNIX)
    set(CMAKE_CXX_COMPILER icpx)
else()
    include (CMakeForceCompiler)
    CMAKE_FORCE_CXX_COMPILER (icx-cl IntelDPCPP)
    include (Platform/Windows-Clang)
endif()
set(target_compile_flags "-fsycl -Wall")
set(target_link_flags "-fsycl")
if (WIN32)
    set(target_compile_flags "${target_compile_flags} /EHsc")
endif ()

// for executable targets...
set_target_properties(${target_fullname} PROPERTIES COMPILE_FLAGS "${target_compile_flags}")
set_target_properties(${target_fullname} PROPERTIES LINK_FLAGS "${target_link_flags}")

接下来编译并运行示例项目 vector-add，看起来成功运行了：

接下来就可以开始写作业了。

作业项目

作业一共有三个选题。都是密集计算中的经典问题，很适合拿来入门异构加速计算。

完整实现参见 Github： ymd45921 / oneapi-homework

矩阵乘法

矩阵乘法是并行计算中的经典计算密集型问题，很适合并行加速；之前用 OpenMP 的时候就知道给三重循环中的一层或者多层套上并行，在这里只是换成了用 SYCL 来描述，道理都是一样的：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16

q.submit([&](sycl::handler &h) {
    sycl::accessor a(a_buf, h, sycl::read_only);
    sycl::accessor b(b_buf, h, sycl::read_only);
    sycl::accessor c(c_buf, h, sycl::write_only, sycl::no_init);

    int a_width = a.get_range()[1];

    h.parallel_for(sycl::range(M, P), [=](sycl::id<2> index) {
        float sum = 0;
        auto row = index[0], col = index[1];
        for (int i = 0; i < a_width; i++) {
            sum += a[row][i] * b[i][col];
        }
        c[index] = sum;
    });
});

这里使用核函数并行地执行前两层循环，同时计算矩阵多个位置的值。但是访问的局部性并不好，因为计算矩阵的每个单元都需要访问整个矩阵的一整行和一整列，这里的数据跨度显然是很大的。

此外，虽然使用二维数组的方式直接声明矩阵有各种方便，但是因为栈空间是有限的，想要处理大数据还是应该动态地分配内存，所以实现了 RAII 矩阵类 my::mat，内部使用动态分配的二维数组存储数据。而且 DevCloud 的服务器内存*这——么——*多，不用白不用不是？

优化

对于并行矩阵乘法这个问题来说，分块优化是一个很经典的优化；虽然分块矩阵乘法和普通的暴力矩阵乘法的复杂度是一样的，但是分块矩阵的每次会将同行和同列的一整块的子矩阵相乘，较小的矩阵更容易全部放进缓存中，从而避免了缓存缺失，并具有了更好的局部性，从而更快的完成计算。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28

q.submit([&](sycl::handler &h) {
    sycl::accessor a(a_buf, h, sycl::read_only);
    sycl::accessor b(b_buf, h, sycl::read_only);
    sycl::accessor c(c_buf, h, sycl::read_write);

    sycl::local_accessor<float, 2> a_t(sycl::range<2>(matrix_unit_size, matrix_unit_size), h);
    sycl::local_accessor<float, 2> b_t(sycl::range<2>(matrix_unit_size, matrix_unit_size), h);

    h.parallel_for(sycl::nd_range<2>(global_size, local_size), [=](sycl::nd_item<2> item) {
        auto local_row = item.get_local_id(0), local_col = item.get_local_id(1);
        auto global_row = item.get_group(0) * matrix_unit_size + local_row;
        auto global_col = item.get_group(1) * matrix_unit_size + local_col;
        auto tiles = matrix_size / matrix_unit_size;
        float acc = 0;
        for (int i = 0; i < tiles; i++) {
            auto tile_col = i * matrix_unit_size + local_col;
            auto tile_row = i * matrix_unit_size + local_row;
            a_t[local_row][local_col] = a[global_row][tile_col];
            b_t[local_row][local_col] = b[tile_row][global_col];
            item.barrier(sycl::access::fence_space::local_space);
            for (int j = 0; j < matrix_unit_size; j++) {
                acc += a_t[local_row][j] * b_t[j][local_col];
            }
            item.barrier(sycl::access::fence_space::local_space);
        }
        c[global_row][global_col] = acc;
    });   
});

这里使用到了数据同步机制之 barrier：它要求当前工作组内所有的工作项都到达 barrier 所在的位置后才能继续执行下去；这里使用了两次 barrier —— 第一次使用它的原因是再利用 a_t 和 b_t 计算之前，需要保证每个工作项都正确加载所需要的局部数据；也就是说必须都完成了对 a_t 和 b_t 的写操作之后，才能读取他们的值用来计算 acc；同样，第二次使用它的原因是需要保证所有的工作项都完成了这一轮对 acc 的贡献才能进行下一轮贡献的计算，保证写入全局矩阵 c 之前已经完成了所有必要的计算；总之，就是避免出现数据竞争。

可以说是最经典的学以致用环节了。一直在教科书中都很抽象的局部性在这里得到了鲜明的体现。在这之上还有 Cannon 算法，不过懒得学了，之后再说（）

执行结果

可以看出并行的矩阵乘法比暴力不知道高到哪里去了，而分块优化后时间还能再进一步的优化 50% 左右的时间，只能说快中快。

关于 Device 时间和 Kernel 时间：前者指的是主机开始向队列提交任务到任务完成的时间，是从主机的角度统计的时间，除了计算消耗的时间，还包括了调用 SYCL 抽象层和传输数据等的开销；后者则指的是在创建队列的时候开启了 sycl::property::queue::enable_profiling 之后，命令的执行时间，可以看作真实的计算时间。

归并排序

归并排序是一个很经典的递归问题，但是在并行执行下显然递归有点抽象；第一想到的就是把递归改成从下到上的合并；然后等待每一层完成合并了再去合并下一层：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29

for (auto stride = 2; stride <= size; stride *= 2) {
    auto event = queue.submit([&](sycl::handler& cgh) {
        auto data = buffer.template get_access<sycl::access::mode::read_write>(cgh);
        auto groups = (size + stride - 1) / stride;
        cgh.parallel_for(sycl::range<1>{groups}, [=](sycl::item<1> item) {
            size_t idx = item.get_id(0);
            size_t left = idx * stride;
            size_t mid = std::min(left + stride / 2, size);
            size_t right = std::min(left + stride, size);

            std::vector<T> tmp(right - left);
            size_t i = left, j = mid, k = 0;

            while (i < mid && j < right) {
                tmp[k++] = (data[i] < data[j]) ? data[i++] : data[j++];
            }
            while (i < mid) {
                tmp[k++] = data[i++];
            }
            while (j < right) {
                tmp[k++] = data[j++];
            }
            for (k = 0, i = left; i < right; ++i, ++k) {
                data[i] = tmp[k];
            }
        });
    });
    event.wait();
});

所以新的任务必须得等待上一轮并行运行结束后才能发起…… 那这段时间 CPU 就只能空等了啊！略蠢。

优化

思路不是特别清晰，但是有一个大致的方向是尽量减少在等待的线程 —— 因为归并的问题数量逐渐减少，问题规模逐渐增大，所以到最后就会只有几个线程在工作，而其他的线程都在旁观。在网上看到了很暴力的两头归并，只能说感觉不是很有意义== 有没有什么更加行之有效的，一看就很爽的方法？

还有就是这里的代码是每层完成了合并后才会开启下一层的合并，这显然可以优化成前面部分合并足够上一层了就触发上层合并。但是一方面感觉带来的优化空间不多，另一方面也完全不知道该怎么写，干脆摆烂了==

一般而言，在并行计算中的归并排序更多的则是 Odd-Even-Merge-Sort；在 NVIDIA 的 CUDA 实例代码中似乎是有的，还被 Intel 拿来示范使用转 DPC++ 工具。可惜我是个土地瓜，之前也没听说过这个排序，也没有学过别的异构计算，一时半会不是很能看得懂（瘫）

执行结果

虽然没有优化，但是这样的速度也已经薄纱手写的归并排序和 STL 的 std::stable_sort 了，只能说摆了；等着有懂哥来教我怎么办了==

图像卷积

图像卷积，或者说卷积运算，实际上就是用一个一般是正方形奇数边长的小矩阵作为”卷积核“，对于图像中的每个像素，滑动卷积核并将其每个元素与图像的局部区域对应位置的值进行乘法运算，然后将结果相加，从而生成新的图像。这样可以提取出图像中的特征，并根据需求增强/减弱它。作为数字图像处理的经典问题，当然是可以通过并行来加速的。

最暴力的实现和暴力的矩阵乘法差不多，也就是加速外两层循环，同时让多个计算单元计算多个像素位置的卷积结果。核函数如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30

queue.submit([&](sycl::handler& cgh) {
        auto accessor_input = buffer_input.get_access<sycl::access::mode::read>(cgh);
        auto accessor_output = buffer_output.get_access<sycl::access::mode::write>(cgh);
        auto accessor_kernel = buffer_kernel.get_access<sycl::access::mode::read>(cgh);
    
        cgh.parallel_for(sycl::range<2>(height, width), [=](sycl::item<2> item) {
            int y = item.get_id(0);
            int x = item.get_id(1);
            float sum_r = 0.0f, sum_g = 0.0f, sum_b = 0.0f, sum_a = 0.0f;
            float sum_weight = 0.0f;
            const auto kernel_offset = kernel_size / 2;
            for (int i = 0; i < kernel_size; ++i) {
                for (int j = 0; j < kernel_size; ++j) {
                    int inputX = x + i - kernel_offset;
                    int inputY = y + j - kernel_offset;
                    auto inputCoord = sycl::id<2>{(unsigned)inputY, (unsigned)inputX};
                    if (inputX >= 0 && inputX < width && inputY >= 0 && inputY < height) {
                        auto weight = accessor_kernel[{(unsigned)i, (unsigned)j}];
                        sum_r += accessor_input[inputCoord].r * weight;
                        sum_g += accessor_input[inputCoord].g * weight;
                        sum_b += accessor_input[inputCoord].b * weight;
                        sum_a += accessor_input[inputCoord].a * weight;
                        sum_weight += weight;
                    }
                }
            }
            sum_r /= sum_weight, sum_g /= sum_weight, sum_b /= sum_weight, sum_a /= sum_weight;
            accessor_output[item] = my::make_pixel_rgba(sum_r, sum_g, sum_b, sum_a);
        });
    });

中途曾经出现过很抽象的结果，一度让我怀疑我对图像卷积的理解是不是出现了问题，然后才发现我傻逼了：我们一般描述图像是 宽×高，然而当在 C++ 中使用二维数组表示图像时，其实代表图像有 高 行和 宽 列，也就是坐标什么的得反过来…… 只能说确实图像处理代码写少了（）

优化

一种很显然的思路是像矩阵乘法那样，尽量让一个组内的工作项处理一块连续的图像 —— 对图像分块就好了。将图像划分成每个小方块，然后每个工作组处理其中的一小块，具有更好的局部性；但是我懒了…… 嗯。至于别的思路只能说毫无想法，有没有懂哥教教==

执行结果

锐化似乎就是这样……？从性能上看更多的成本在数据传输&抽象层上，实际计算很迅速。

后记

那么这次的作业就暂时告一段落了。虽然很难说我会用 OneAPI 了，但是至少透过这样的一个标准浅浅窥探了一眼异构计算的领域到底是什么样的，想必对之后写 Compute Shader 或者是学 CUDA 一定大有裨益吧（笑）

逼逼赖赖

为什么后进队列的任务会被先排到计算节点？（

虽然说 Vulkan 是低级 API，而 SYCL 是高级的抽象层，但是感觉现在写的 DPC++ 基本上也是一直在描述模型啊，优化什么的都是编译器给干的；现在尚且还能写一些时空局部性的优化，除了这之外也算是毫无想法，至于别的什么高级的用指令优化呢，那是和用户没有一点关系…… 而且又是 LLVM 啊== 让我想到了之前和友人吹牛

唉，于是我是不是也该考虑面对 LLVM 呢？要学的东西太多了！

参考

• oneAPI-samples/Tools/Advisor/matrix_multiply_advisor at master · oneapi-src/oneAPI-samples · GitHub
• GitHub - kevinknow/OneAPI_Matrix_acceleration