基于Intel oneAPI的并⾏矩阵乘法

基于Intel oneAPI的并⾏矩阵乘法

问题描述

编写⼀个基于oneAPI的C++/SYCL程序来执行矩阵乘法操作。需要考虑大尺寸矩阵的乘法操作以及不同线程之间的数据依赖关系。通常在实现矩阵乘法时，可以使用块矩阵乘法以及共享内存来提高计算效率。

技术方案

我的作业中一共使用了3种不同的矩阵乘算法：朴素串行矩阵乘法、朴素并行矩阵乘法、分块并行矩阵乘法。

根据下一节的测试，并行朴素算法相对串行算法加速比约为10，而并行分块算法相对串行算法加速比约为67。

• 朴素串行矩阵乘法

  $O(n^3)$时间复杂度的朴素矩阵乘法，串行执行，作为baseline。

  void sequentialMatrixMultiply(const std::vector<float>& A, const std::vector<float>& B, std::vector<float>& C) {
      for (size_t i = 0; i < N; ++i) {
          for (size_t j = 0; j < N; ++j) {
              float sum = 0.0f;
              for (size_t k = 0; k < N; ++k) {
                  sum += A[i * N + k] * B[k * N + j];
              }
              C[i * N + j] = sum;
          }
      }
  }
  

• 朴素并行矩阵乘法

  创建三个二位buffer，并获取迭代器，使用Intel oneAPI提供的parallel_for，展开矩阵乘法内的两层循环，代码如下：

  sycl::queue myQueue;
  sycl::range<2> size(N, N);
  sycl::buffer<float, 2> bufferA(matrixA.data(), size);
  sycl::buffer<float, 2> bufferB(matrixB.data(), size);
  sycl::buffer<float, 2> bufferC(matrixC.data(), size);
  myQueue.submit([&](sycl::handler& cgh) {
    auto accessorA = bufferA.get_access<sycl::access::mode::read>(cgh);
    auto accessorB = bufferB.get_access<sycl::access::mode::read>(cgh);
    auto accessorC = bufferC.get_access<sycl::access::mode::write>(cgh);
  
    cgh.parallel_for<class MatrixMultiply>(size, [=](sycl::id<2> idx) {
      float sum = 0.0f;
      for (int k = 0; k < N; ++k) {
        sum += accessorA[idx[0]][k] * accessorB[k][idx[1]];
      }
      accessorC[idx] = sum;
    });
  });
  

• 分块并行矩阵乘法

  基于Intel oneAPI提供的并行计算接口，将原矩阵分成若干块，每块大小作为一个超参数指定，对于每个小块都可以并行处理，进行分块矩阵乘算法。

  constexpr size_t N = 1024;
  
  constexpr size_t blockSize = 32;
  
  void blockMatrixMultiply(sycl::queue& queue, const std::vector<float>& A, const std::vector<float>& B, std::vector<float>& C) {
      sycl::range<2> globalSize(N, N);
      sycl::range<2> localSize(blockSize, blockSize);
  
      sycl::buffer<float, 2> bufferA(A.data(), sycl::range<2>(N, N));
      sycl::buffer<float, 2> bufferB(B.data(), sycl::range<2>(N, N));
      sycl::buffer<float, 2> bufferC(C.data(), sycl::range<2>(N, N));
  
      queue.submit([&](sycl::handler& cgh) {
          auto accessorA = bufferA.get_access<sycl::access::mode::read>(cgh);
          auto accessorB = bufferB.get_access<sycl::access::mode::read>(cgh);
          auto accessorC = bufferC.get_access<sycl::access::mode::write>(cgh);
  
          cgh.parallel_for<class BlockMatrixMultiply>(sycl::nd_range<2>(globalSize, localSize), [=](sycl::nd_item<2> item) {
              size_t row = item.get_global_id(0);
              size_t col = item.get_global_id(1);
  
              float sum = 0.0f;
              for (size_t i = 0; i < N; i += blockSize) {
                  for (size_t j = 0; j < blockSize; ++j) {
                      sum += accessorA[row][i + j] * accessorB[i + j][col];
                  }
              }
              accessorC[row][col] = sum;
          });
      });
      queue.wait();
  }
  

性能测试与验证

对于作业的运行环境，我使用了Intel DevCloud免费提供的具有oneAPI环境的Jupyter Lab(https://jupyter.oneapi.devcloud.intel.com/)。

Intel Devcloud Jupyter Lab提供持久化存储，并且集成了Intel提供的各种开发套件环境，和许多示例教程，亦提供免费的计算队列资源，可以直接在浏览器中访问开发环境完成开发，包括Intel CPU、GPU等均可免费使用，无需本地配置任何环境。

下面测试三种算法的运行效率（用加速比来衡量），以及正确性（通过与朴素算法比较计算结果确定）。

首先，使用随机浮点数初始化A、B两个矩阵，对于验证结果正确性，只要inconsistent_cnt为0，说明计算结果一致：

inconsistent_cnt = 0;
for(int i = 0; i < N; ++i) {
  for(int j = 0 ; j < N; ++j) {
    if(fabs(matrixC[i * N + j] - matrixC_std[i * N + j]) > 1e-3) {
      inconsistent_cnt ++;
    }
  }
}

接下来，对于运行效率，使用std::chrono::high_resolution_clock::now()进行计时，作差比较运行时间。

测试的矩阵大小为1024*1024，分块算法的块大小为32。

通过如下脚本编译并运行作业程序：

#!/bin/bash
source /opt/intel/oneapi/setvars.sh > /dev/null 2>&1
/bin/echo "##" $(whoami) is compiling SYCL_Essentials Module1 -- oneAPI Intro sample - 1 of 1 homework1.cpp
icpx -fsycl lab/homework1.cpp
if [ $? -eq 0 ]; then ./a.out; fi

运行结果如下：

可见在较大的矩阵下，并行朴素算法相对串行算法加速比为10.6487，而并行分块算法相对串行算法加速比为67.0414，并行计算对性能优化显得尤为重要。

学习心得

在这门课程中，我通过实际操作深入理解了Intel oneAPI的核心概念和优势。oneAPI提供了一个统一的编程模型，用于跨多种硬件（如CPU、GPU、FPGA）构建高性能应用，这一点在目前的高性能计算框架中难能可贵。

oneAPI强大之处在于其支持异构计算。在此作业中，我利用了oneAPI的SYCL扩展，它提供了一个标准的C++编程模型，简化了在不同类型的处理器上编程的复杂性，且算法本身可以在不进行任何更改的情况下，支持异构计算，这种跨平台兼容性对于降本增效、优化性能至关重要。我使用了oneAPI中的并行编程工具来加速矩阵乘法。这个过程让我体会到了并行化对于提高大规模计算任务效率的重要性，且即便是Intel Devcloud Jupyter Notebook中提供的CPU资源，在应用并行加速后加速比也十分优秀。

通过使用SYCL和一系列API，我更深入地理解了如何将问题分解为可以并行处理的小任务。这种思维方式不仅对于使用oneAPI，而且对于现代高性能计算领域都是极其重要的。在开发过程中，我也学会了使用Intel DevCloud的工具进行调试和性能分析。这些工具帮助我识别并优化了代码中的瓶颈，这对于编写高效的并行程序至关重要。

最后，感谢Intel Devcloud提供的免费计算资源和环境，可以让我快速上手直接进行实验，无需进行任何环境配置，或购买任何特定硬件即可在真实的应用场景中进行实践，欢迎大家也来体验。