write: heterogeneous programming model

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:
 - 编译原理
 - 组会汇报
 ---
 
+
 随着摩尔定律的逐渐失效，将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(
-    const std::vector<std::vector<int>>& a,
-    const std::vector<std::vector<int>>& b)
+inline std::vector<int> cpuMatrixMultiply(
+    const std::vector<int>& a,
+    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.