write: heterogeneous programming model

YaeBlog/source/posts/cncc-2024.md

@@ -0,0 +1,67 @@
+---
+title: 2024中国计算机大会
+date: 2024-11-03T14:06:36.4212070+08:00
+tags:
+- 杂谈
+---
+
+2024年的中国计算机大会于10月24日到10月26日在浙江省金华市东阳市横店镇举办，而鄙人在下不才我，有幸受到实验室资助前去参观学习。
+
+<!--more-->
+
+首先开幕式镇楼。
+
+![image-20241102212738598](./cncc-2024/image-20241102212738598.png)
+
+## 学术上
+
+大会每天的日程是上午的大会特邀报告和大会论坛，下午的各个分论坛讨论。老实说，大会上午的报告和论坛我都没有特别感兴趣，因此这里将重点放在我参加的三个分论坛上。
+
+### AI时代的异构融合操作系统：聚散终有时，融合亦有期
+
+第一个报告是华为庞加莱实验室秦彬娟老师的《异构智算时代的操作系统演进》。报告高屋建瓴，从比较宏观的角度上介绍了当前异构融合操作系统诞生的背景、发展的方向。在报告中重点介绍了一种异构融合操作系统的设计思路：通过三层架构，基于互联池化技术，构建AI时代的融合算力系统。系统中的三层包括：（1）池化基础底层，包括多设备的融合和池化设备虚拟化；（2）异构融合核心子系统，例如异构融合调度系统、异构融合内存和异构融合存储系统；（3）异构核心服务。总的来说，这个报告在一定程度上勾勒出了未来一个异构融合操作系统应有的各项功能，但是显然这一操作系统的实现还存在着明显的困难。
+
+![image-20241102211959206](./cncc-2024/image-20241102211959206.png)
+
+下面一个报告是较为有干货的报告，北京航空航天大学刘瀚骋老师的《异构融合OS及多样性内存管理框架》。报告中介绍了一个称作`FMMU`的系统，是对于异构融合操作系统中内存管理系统的探索。报告中首先介绍了内存池化技术对于异构融合操作系统的重要性，指出分布式共享内存（Distributed Shared Memory）可能是实现内存池化技术的未来。然后介绍了将部分内存管理中的计算卸载到可编程网络硬件中来加速分布式内存访问的新思路。最后在报告中提到了内存管理技术如何解决错误预测和错误回复的问题。虽然在听的时候没太注意，但是现在总结的时候才发现这个报告的思路似乎有点混乱，尤其是最后一点和内存管理系统并没有什么直接的关系，而且这个内存管理系统似乎不是**异构系统**的内存管理，反而是分布式系统的内存管理。不过总的来说，这个报告还是非常实际的，介绍了不少当前异构融合操作系统中的内存管理面临的问题和解决问题的探索。
+
+![image-20241102212355390](./cncc-2024/image-20241102212355390.png)
+
+第三个报告是国防科技大学李东升老师的《异构计算环境下的分布式深度学习训练》。报告首先从李老师的主业——并行计算起手，介绍了深度学习训练过程中主要的各种并行方法，例如数据并行、模型并行和混合并行等，指出目前大模型的并行训练存在着计算/存储/通信难的问题。因此，提出了一个智能模型训练并行任务划分方法：（1）基于符号算子的计算图定义方法；（2）面向Transformer模型的流水线并行任务划分方法；（3）异构资源感知的流水线并行任务划分方法。然后针对分布式模型训练中通信调度存在的通信墙、数据依赖关系复杂等的问题，提出综合词嵌入表的稀疏通信调度技术、流水线并行的P2P通信调度技术、模型计算的统一操作执行引擎和网络链路感知的通信执行引擎的通信调度技术。最后提到了智能模型训练 的内存优化技术，针对现有重计算技术（re-computing）和存储交换（swapping）技术存在的问题，提出了一种面向大型智能模型训练的细粒度内存优化方法`DELTA`。
+
+最后一个报告是上海交通大学杜冬冬老师的《软硬芯异构融合操作系统的多个维度》。报告伊始，杜老师就抛出一个问题：操作系统的演进应该是提供新的抽象还是兼容现有的抽象？在回答这个问题之前，杜老师首先介绍他们一个异构融合操作系统的设计思路：层OS架构的思路，通过设置两个层次——全局OS和本地OS，全局OS在本地OS的基础上提供一层跨`XPU`的能力。杜老师设计的这个系统称作`XPU-Shim`，在设计这个系统时就面对着前面的问题，是提供新的抽象还是兼容现有的抽象。`XPU-Shim`的回答是兼容现有的抽象，在底层的CXL、UB等内存语义总线的基础上实现了传统的Socket抽象，提供了低时延、高吞吐的协同能力。在操作系统的抽象问题之外，杜老师还就云上GPU应用的启动时延问题进行了讨论，深入解释了通过状态复用完全跳过初始化阶段从而加速应用冷启动过程的思路。
+
+Plane讨论没有参加。
+
+### 编译系统前沿技术与应用
+
+第一个报告是清华大学陈文光老师的《神经网络全同态编译器》。这个报告可以说证明了“编译技术的人才活跃在各行各业”，报告中的主要内容就是编译技术如何助力机密计算中的全同态加密应用在神经网络的推理中。全同态加密算法实现了“数据可用不可见”的概念，允许程序直接在密文上进行乘法和加法运算，但是限制也是只能进行加法和乘法运算，而且过多的乘法操作会造成计算之后解密失败。该编译器成为`ANT-ACE`，首先通过设计新的五层中间表示（IR）实现了自动化全同态加密程序生成和面向性能的优化设计，在实现基本的编译工作之外，`ANT-ACE`提供了一定的调试支持，通过部分支持对于模型的部分加密支持和运行时校验为解决加密之后程序推理准确率下降的问题。
+
+接下来三个报告都是关于如何将人工智能技术同编译技术解决起来。计算所冯晓兵老师的报告《人工智能编译领域的应用探索》，介绍了大模型同编译后端的两个结合方向：（1）使用大模型生成编译器的后端代码；（2）使用大模型替换编译器的后端，直接利用大模型生成汇编代码。华为毕昇编译器架构师魏伟的报告《AI for Compiler的技术探索和应用实践》则是介绍了毕昇编译器的自动调优器`Autotuner`，这个一个自动寻找最优化的编译参数组合工具。复旦大学张为华老师的报告《基于学习的编译优化技术》也是一个类似的工作，利用机器学习技术挖掘已有的编译系统中存在的相关知识来指导新的编译优化。
+
+最后一个报告则是字节公司郑思泽研究员的《计算通信融合中的编译器设计》，该报告主要聚焦于如何实现在深度学习算子层的计算通信融合，这个报告主要由搞`MLIR`的同学听，我就摸鱼了。
+
+### 智能终端操作系统OpenHarmony前沿研究
+
+虽然名字叫作OpenHarmony，但是感觉内容实际上和鸿蒙系统没有什么太大的关系。
+
+第一个报告是软件所武延军老师的《万物智联时代基础软件如何驯服碎片化》。报告的标题非常的高大上，但是实际上就讲了两件事情：（1）RISCV架构，或者说RISCV这个可扩展的思想，是解决架构碎片化的思路；（2）`openEular`系统可以作为系统软件适配的一个基线操作系统。总结一下，这其实就是一个广告，希望大家做基础软件的都来和大家一起做。
+
+第二个报告是南京大学冯新宇老师的《基于仓颉语言的嵌入式DSL开发》，同时冯新宇老师也是仓颉语言的首席架构师。冯老师的这个报告主要聚焦于仓颉语言提供的嵌入式DSL能力，而嵌入式DSL这一设计范式已经在前端开发中展现了不俗的潜力。报告中介绍了嵌入式DSL出现的背景，仓颉中为了提供嵌入式DSL而引入的语法糖、仓颉提供的嵌入式DSL工具箱等。虽然仓颉语言是一个主要面向上层应用开发的语言，但是仓颉中丰富的DSL能力还是给异构编程模型的设计提供了不少的启发。而且目前在各种深度学习编译器中DSL的应用也非常广泛，例如`triton`。
+
+![image-20241102212536635](./cncc-2024/image-20241102212536635.png)
+
+第三个报告是在存算一体的芯片上做数据库的加速，第四个报告是OpenHarmony上`ArkTS`程序的静态分析，都没怎么听。
+
+最后一个又是上交杜冬冬老师的报告，《面向下一代智能终端操作系统的渲染服务研究与挑战》。这是一个我感觉还挺有趣的报告，报告中介绍的主要背景是随着终端设备上屏幕刷新率的提高和操作系统动画变得更加精致复杂，用户会发现终端系统上的显示卡顿越来越多、越明显。这是因为目前的终端显示刷新机制是同步的，显示屏会按照当前刷新的频率从操纵系统中读取下一帧的画面，但是操作系统面对这越来越短的刷新时延和越来越复杂的动画常常不能按时把下一帧的画面渲染好。于是我们的杜冬冬老师就提出了一种动态、异步的渲染机制，考虑到系统中显示动画的时间还是占少部分的，于是就可以借用这些系统不繁忙的时间预先渲染（削峰填谷）。但是这种方式需要预知到系统后面会显示的内容，这使得这套技术只能在确定性的场景和部分简单交互场景下使用。
+
+> 这里插入一个杜冬冬老师的八卦，杜老师改过一次名字，之前的名字是杜东（Dong Du），在查找论文的时候使用后面的名字会更好一些（在[IPADS](https://ipads.se.sjtu.edu.cn/zh/members/)和[dblp](https://dblp.org/pid/48/331-3.html)上面都还没有改过来）。
+
+## 其他
+
+首先我要锐评一下浙江省金华市东阳市横店镇。横店镇感觉完全没有为一个旅游目的地做过准备，虽然说镇子上面的酒店还是挺多的，但是不管是吃的还是玩的感觉都非常少。而且镇上的交通简直就是一坨，尤其是我们从酒店到会议举办地圆明新园的一段路，完全被大货车摧残的不成样子，在上面坐车堪比过山车。
+
+然后我要锐评一下会议的举办地横店圆明新园。在去之前听说这里是1:1复刻了被八国联军烧毁的圆明园，结果去了才发现圆明新园分成春苑、夏苑和秋苑，其中春苑是复刻的圆明园，但是会议的举办地是在夏苑和秋苑，感觉有点的被诈骗了。夏苑里面只复刻了圆明园长春园的部分景观，比如海岳开襟、谐奇趣和大水法等，而且还增设了英、法、美、俄、日、德、意和奥等国的特色建筑，而会议就主要在这些特色建筑中进行，属实感觉有点奇怪了。
+
+最后我要锐评一下CNCC会议。名义上看这个会议有涵盖数十个方向的130余场论坛，上万名注册参会者的大型会议，但是这个会议却选在了一个看上去基本上不适合召开大型会议的横店镇圆明新园。同时会议进行的非常寒酸，中午的午餐是横店提供给剧组的盒饭，在主会场发给我们之后只能自己端着吃，下午的茶歇更是少的可怜，除了第三天有好哥们分了我一块蛋挞，三天的茶歇我愣是一点都没见到（有可能是第三天的人最少，提高了我获得茶歇的概率）。
+

YaeBlog/source/posts/heterogeneous-programming-model.md

@@ -1,11 +1,12 @@
 ---
 title: 异构编程模型的昨天、今天与明天
-date: 2024-10-16T13:34:49.0270134+08:00
+date: 2024-11-04T22:20:41.2571467+08:00
-tags: 
+tags:
-    - 编译原理
+- 编译原理
-    - 组会汇报
+- 组会汇报
 ---
 
+
 随着摩尔定律的逐渐失效，将CPU和其他架构的计算设备集成在片上或者通过高速总线互联构建的异构系统成为了高性能计算的主流。但是在这种系统中，上层应用的设计与实现面临着异构系统中各个设备之间体系结构差异过大、缺乏良好的异构抽象以及统一的编程接口和应用程序的优化难度大等困难。
 
 异构并行编程模型便是解决这些编程和执行效率问题的解决方案。
@@ -24,7 +25,7 @@ tags:
 
 首先是异构系统中各个设备之间的并行计算能力不同。在同构的并行计算系统中，比如多核CPU中，虽然同一CPU的不同核之间、同一核的不同SIMD部件之间可以承担不同粒度的并行计算任务，但是其并行计算的能力是完全相同的。但是在一个典型的异构计算系统，例如CPU、GPU和FPGA组成的异构系统，不同设备的微架构具有本质差异，其并行计算的模式和能力都完全不同，设备之间的特长也完全不同。这种设备之间并行计算能力的差异使得系统中的任务划分和任务映射不再是均一的，而是具有显著的特异性。这种特点虽然也有利于表达实际应用的特点，但是却给异构并行计算模型的设计带来了巨大的困难。
 
-![9eb06d8be92ddef3db33e040163c67a7.png](./heterogeneous-programming-model/9eb06d8be92ddef3db33e040163c67a7.png "9eb06d8be92ddef3db33e040163c67a7.png")
+![9eb06d8be92ddef3db33e040163c67a7.png](./heterogeneous-programming-model/9eb06d8be92ddef3db33e040163c67a7.png)
 
 其次是异构系统中加速设备数据分布可配置、设备间数据通信渠道多样性给数据分布和通信带来的困难。在同构并行系统中，CPU片内的存储是对于软件透明的缓存架构，在片外则是一个共享内存模型，因此在这类系统中，数据仅可能分布在片外的共享存储中，具有存储位置单一的特点，也不需要进行显式的通信操作。但是在异构系统中，不仅在单个加速设备内部可能有软件可分配的快速局部存储，设备之间的连接方式差异也很大。目前，大多个加速设备都是通过PCIe总线的方式同CPU进行连接，这使得加速设备无法通过和CPU相同的方式完成地址映射，存在某一设备无法访问另一设备片外存储的问题。这使得异构系统中数据可以分布在CPU、加速设备的片外存储和加速设备的片内多层次局部存储等多个位置，不仅使得编程模型的数据分布问题变得十分复杂，设备间的通信文件也可能需要显式进行。
 
@@ -230,11 +231,11 @@ private:
 作为对比，一个使用CPU单线程计算的例子如下：
 
 ```cpp
-std::vector<std::vector<int>> matrix_multiply(
+inline std::vector<int> cpuMatrixMultiply(
-    const std::vector<std::vector<int>>& a,
+    const std::vector<int>& a,
-    const std::vector<std::vector<int>>& b)
+    const std::vector<int>& b)
 {
-    std::vector result(MATRIX_SIZE, std::vector(MATRIX_SIZE, 0));
+    std::vector result(MATRIX_SIZE * MATRIX_SIZE, 0);
 
     for (int i = 0; i < MATRIX_SIZE; i++)
     {
@@ -243,9 +244,10 @@ std::vector<std::vector<int>> matrix_multiply(
             int temp = 0;
             for (int k = 0; k < MATRIX_SIZE; k++)
             {
-                temp += a[i][k] * b[k][j];
+                // a[i][j] = a[i][k] * b[k][j] where k in (0..MATRIX_SIZE)
+                temp += a[i * MATRIX_SIZE + k] * b[k * MATRIX_SIZE + j];
             }
-            result[i][j] = temp;
+            result[i * MATRIX_SIZE + j] = temp;
         }
     }
 
@@ -255,7 +257,7 @@ std::vector<std::vector<int>> matrix_multiply(
 
 ### OpenMP
 
-OpenMP是`Opem MultiProcessing`的缩写，是一个使用编译器制导（Directives）来进行共享内存平行计算的框架，在C、C++和Fortran语言的并行编程中得到的了广泛的应用。OpenMP提供了一个简单而灵活的接口，让程序员能够充分释放多核和多处理器系统性能。
+OpenMP是`Open MultiProcessing`的缩写，是一个使用编译器制导（Directives）来进行共享内存平行计算的框架，在C、C++和Fortran语言的并行编程中得到的了广泛的应用。OpenMP提供了一个简单而灵活的接口，让程序员能够充分释放多核和多处理器系统性能。
 
 OpenMP从上面的介绍来看似乎并不是一个严格的异步并行编程模型，但是第一，OpenMP作为一个经典的并行编程框架，研究价值还是非常高的，其次在一些较新的OpenMP版本中其宣称也能利用NVIDIA GPU进行加速，似乎也能算是一个异构并行编程模型。
 
@@ -446,7 +448,9 @@ std::vector<std::vector<int>> cudaCalculateMatrix(const std::vector<std::vector<
 > - 首先是换了一台没有大小核异构设计的计算机进行实验，发现这下两次使用CPU计算的时间差异不大；
 > - 加上了热身的阶段之后，计算时间没有发生明显的变化。
 >
-> 综上所述，可以认为此现象和异构CPU之间存在这明显的关联，但是缺乏直接证据。
+> 综上所述，可以认为此现象和异构CPU之间存在着明显的关联，但是缺乏直接证据。
+>
+> 在我们调整了矩阵的数据布局之后，这里提到的实验结果又发生了变化。上面的实验结果是使用二维数据存储矩阵得到的，而在修改为使用一维数组（也就是现在提供的代码）之后，相同的CPU计算代码的计算时间又没有产生明显的变化了。看来这个问题可能和数据布局、CPU缓存等问题相关。
 
 ### OpenCL
 
@@ -772,20 +776,93 @@ OpenACC是作为一个标准的形式提供的，实现了该标准的编译器
 | GCC 12                                                       | 支持到OpenACC 2.6                                            |
 | [Omni Compiler Project](https://github.com/omni-compiler/omni-compiler) | 源到源编译器，将带有制导的源代码翻译到带有运行时调用的平台代码，近两年没有活跃开发 |
 | [OpenUH](https://github.com/uhhpctools/openuh)               | 项目开发者在7年前的最后一次提交了中删除了README中有关OpenACC的内容 |
+| [OpenArc](https://csmd.ornl.gov/project/openarc-open-accelerator-research-compiler) | 是学术界出品的还在活跃开发的编译器，看上去还做了不少工作的样子，就是OpenACC官网上的链接已经失效了找起来比较麻烦，而且宣称是一个开源编译器，但是获取源代码和二进制文件需要联系他们（美国橡树岭国家实验室）创建账户，这看去对于我们这些Foreign Adversary有些抽象了。 |
 
+在试验OpenACC时遇到了巨大的困难，不论是使用gcc还是NVIDIA HPC SDK都没有办法实现明显的并行编程加速，多次实验之后都没有找到的问题的所在。这里还是贴一下实验的代码和实验的数据。
 
+实验中编写的OpenACC加速代码如下：
+
+```cpp
+static std::vector<int> OpenACCCpuCalculateMatrix(const std::vector<int>& a, const std::vector<int>& b)
+{
+    constexpr int length = MATRIX_SIZE * MATRIX_SIZE;
+
+    const auto aBuffer = new int[length];
+    const auto bBuffer = new int[length];
+    const auto cBuffer = new int[length];
+
+    for (int i = 0; i < length; i++)
+    {
+        aBuffer[i] = a[i];
+        bBuffer[i] = b[i];
+        cBuffer[i] = 0;
+    }
+
+#pragma acc enter data copyin(aBuffer[0:length], bBuffer[0:length])
+#pragma acc enter data create(bBuffer[0:length])
+#pragma acc data present(aBuffer[0:length], bBuffer[0:length], cBuffer[0:length])
+    {
+#pragma acc kernels loop independent
+        for (int i = 0; i < MATRIX_SIZE; i++)
+        {
+#pragma acc loop independent
+            for (int j = 0; j < MATRIX_SIZE; j++)
+            {
+                int temp = 0;
+#pragma acc loop independent reduction(+:temp)
+                for (int k = 0; k < MATRIX_SIZE; k++)
+                {
+                    temp += aBuffer[i * MATRIX_SIZE + k] * bBuffer[k * MATRIX_SIZE + j];
+                }
+                cBuffer[i * MATRIX_SIZE + j] = temp;
+            }
+        }
+    }
+#pragma acc exit data copyout(cBuffer[0:length])
+#pragma acc exit data delete(aBuffer[0:length], bBuffer[0:length])
+
+    std::vector result(MATRIX_SIZE * MATRIX_SIZE, 0);
+
+    for (int i = 0; i < length; ++i)
+    {
+        result[i] = cBuffer[i];
+    }
+
+    delete[] aBuffer;
+    delete[] bBuffer;
+    delete[] cBuffer;
+
+    return result;
+}
+```
+
+实验中使用分别使用`NVIDIA HPC SDK`和`GCC`编译运行的结果如下：
+
+| 编译器         | 类型    | 运行时间 |
+| -------------- | ------- | -------- |
+| NVIDIA HPC SDK | OpenACC | 19315ms  |
+| NVIDIA HPC SDK | CPU     | 22942ms  |
+| GCC            | OpenACC | 19999ms  |
+| GCC            | CPU     | 22623ms  |
 
 ### oneAPI
 
+oneAPI是Intel公司提出的一套异构并行编程框架，该框架致力于达成如下几个目标：（1）定义一个跨架构、跨制造商的统一开放软件平台；（2）允许同一套代码可以在不同硬件制造商和加速技术的硬件上运行；（3）提供一套覆盖多个编程领域的库API。为了实现这些目标，oneAPI同上文中已经提到过的开放编程标准SYCL紧密合作，oneAPI也提供了一个SYCL的编译器和运行时；同时oneAPI也提供了一系列API库，包括`oneDPL`、`oneDNN`、`oneTBB`和`oneMKL`等。
 
+![image-20241103162259981](./heterogeneous-programming-model/image-20241103162259981.png)
 
-### Julia
+我对于oneAPI的理解就是Intel用来对标NVIDIA的CUDA的一套高性能编程工具箱。首先为了和NVIDIA完全闭源的CUDA形成鲜明的对比，Intel选择了OpenCL合作同时开发SYCL，当时也有可能是Intel知道自己的显卡技不如人，如果不兼容市面上其他的部件是没有出路的，同时为了和CUDA丰富的生态竞争，Intel再开发并开源了一系列的`oneXXX`。
 
+这里我就把上面SYCL写的例子用Intel提供的`DPC++`编译运行一下，看看在效率上会不会有所变化。
 
+| 类型                          | 运行时间 | 比率  |
+| ----------------------------- | -------- | ----- |
+| Intel UHD Graphics 770 oneAPI | 429ms    | 0.023 |
+| NVIDIA 4060 Ti oneAPI         | 191ms    | 0.010 |
+| Intel i5-13600K oneAPI        | 198ms    | 0.011 |
+| CPU                           | 18643ms  | 1.000 |
 
-### Triton
+在显卡上的计算时间没有明显的变化，但是我们Intel的编译器却在选择到使用Intel CPU进行计算时展现了不俗的实力。
-
-
 
 
 ## 参考文献
@@ -795,5 +872,4 @@ OpenACC是作为一个标准的形式提供的，实现了该标准的编译器
 3. Exploring the performance of SGEMM in OpenCL on NVIDIA GPUs. [https://github.com/CNugteren/myGEMM](https://github.com/CNugteren/myGEMM)
 4. OpenACC Programming and Best Practices Guide. [https://openacc-best-practices-guide.readthedocs.io/en/latest/01-Introduction.html](https://openacc-best-practices-guide.readthedocs.io/en/latest/01-Introduction.html)
 5. oneAPI What is it?. [https://www.intel.com/content/www/us/en/developer/articles/technical/oneapi-what-is-it.html](https://www.intel.com/content/www/us/en/developer/articles/technical/oneapi-what-is-it.html)
-6.