oneAPI GPU 优化指南 - 使用性能库

利用性能库提升计算效率：矩阵乘法与oneAPI示例,

翻译已于 2024-01-05 16:33:22 修改 · 427 阅读

文章标签：

#oneapi

于 2024-01-05 16:22:27 首次发布

oneAPI GPU 优化指南专栏收录该内容

46 篇文章

订阅专栏

本文介绍了如何通过使用性能库如onemkl和onednn，而非自定义实现，来提高矩阵乘法等计算密集型任务的性能。作者展示了SYCL和USM在加速器上的应用，并强调了oneAPI库如oneMKL、oneDAL和oneDNN在封装优化算法和分离应用程序开发者与优化专家关注点的重要性。

本章节翻译by chenchensmail@163.com 原文：Using Performance Libraries (intel.com)

使用性能库

本节讨论使用来自库的高效函数，例如|onemkl|或|onednn|，而不是手工编写的替代方案。除非你是研究特定数学运算的专家，否则编写自己版本的该运算通常是一个坏主意。例如，矩阵乘法是一种常见的、直接的数学运算。

$C_{m,n} = A_{m,k} \times B_{k,n} = \sum_{}^{k}A_{m,k} \times B_{k,n}$

只需几行代码即可轻松实现：

  // Multiply matrices A and B
  for (m = 0; m < M; m++) {
    for (n = 0; n < N; n++) {
      C[m][n] = 0.0;
      for (k = 0; k < K; k++) {
        C[m][n] += A[m][k] * B[k][n];
      }
    }
  } // End matrix multiplication

然而，这种初级的实现不会给出最佳性能。内部循环的简单视觉检查显示矩阵B的内存访问是不连续的。缓存重用，因此性能将很差。

将初级算法移植到 SYCL 并不难，以将矩阵乘法 kernel 部署到加速器上。以下代码初始化队列以提交工作到默认设备，并在统一共享内存 (USM) 中为矩阵分配空间:

   // Initialize SYCL queue
   sycl::queue Q(sycl::default_selector_v);
   auto sycl_device = Q.get_device();
   auto sycl_context = Q.get_context();
   std::cout << "Running on: "
             << Q.get_device().get_info<sycl::info::device::name>() << std::endl;
 
   // Allocate matrices A, B, and C in USM
   auto A = sycl::malloc_shared<float *>(M, sycl_device, sycl_context);
   for (m = 0; m < M; m++)
     A[m] = sycl::malloc_shared<float>(K, sycl_device, sycl_context);
 
   auto B = sycl::malloc_shared<float *>(K, sycl_device, sycl_context);
   for (k = 0; k < K; k++)
     B[k] = sycl::malloc_shared<float>(N, sycl_device, sycl_context);
 
   auto C = sycl::malloc_shared<float *>(M, sycl_device, sycl_context);
   for (m = 0; m < M; m++)
     C[m] = sycl::malloc_shared<float>(N, sycl_device, sycl_context);
 
   // Initialize matrices A, B, and C

USM 中的数据可以通过 SYCL runtime 在主机和设备内存之间移动。不需要显式缓冲。为了将计算部署到默认加速器，它被转换为 SYCL kernel 并提交到队列：

   // Offload matrix multiplication kernel
   Q.parallel_for(sycl::range<2>{M, N}, [=](sycl::id<2> id) {
      unsigned int m = id[0];
      unsigned int n = id[1];
 
      float sum = 0.0;
      for (unsigned int k = 0; k < K; k++)
        sum += A[m][k] * B[k][n];
 
      C[m][n] = sum;
    }).wait(); // End matrix multiplication

然而，仅仅将这样的代码部署到加速器上不太可能恢复性能。实际上，性能可能会变得更差。糟糕的编写代码无论是在主机上还是在设备上运行，都是糟糕的。

常见的计算密集型操作，如矩阵乘法，已经得到了广泛的研究。专家们设计了许多算法，这些算法比基本数学公式的初级实现具有更好的性能。他们还使用调优技术，如缓存阻塞和循环展开，以在不管矩阵A和B的形状如何的情况下实现性能提升。

oneMKL 提供了一个优化的通用矩阵乘法函数 (oneapi::mkl::blas::gemm)，可在主机处理器或各种加速器设备上获得高性能。与以前一样，在 USM 中分配矩阵，并将其与设备队列、矩阵维度和各种其他选项一起传递给 gemm 函数：

  // Offload matrix multiplication
  float alpha = 1.0, beta = 0.0;
  oneapi::mkl::transpose transA = oneapi::mkl::transpose::nontrans;
  oneapi::mkl::transpose transB = oneapi::mkl::transpose::nontrans;
  sycl::event gemm_done;
  std::vector<sycl::event> gemm_dependencies;
  gemm_done = oneapi::mkl::blas::gemm(Q, transA, transB, M, N, K, alpha, A, M,
                                      B, K, beta, C, M, gemm_dependencies);
  gemm_done.wait();

库函数比初级的实现更加通用，预计会提供更好的性能。例如，如果需要，库函数可以在乘法之前转置一个或两个矩阵。这说明了应用程序开发人员和调优专家之间的关注点分离。前者应该依赖后者将常见计算封装在高度优化的库中。 oneAPI 规范定义了许多库，以帮助创建加速应用程序，例如：

用于数学运算的 oneMKL 用于数据分析和机器学习的 oneDAL 用于深度学习框架开发的 oneDNN 用于视频处理的 oneVPL 在创建自己的实现之前，请检查所需操作是否已在oneAPI库中可用。