OneAPI 矩阵乘法

OneAPI 矩阵乘法

OneAPI 是一个由英特尔（Intel）推动的跨架构编程模型和开发工具的倡议。该倡议的目标是使开发人员能够在不同类型的处理器架构上编写性能高效的代码，包括 CPU、GPU、FPGA 等。OneAPI 的设计理念是实现统一的编程模型，以便开发人员能够更容易地利用异构计算资源，而不必为每个硬件架构单独编写不同的代码。

在SYCL编程模型中，parallel_for 函数用于指示SYCL运行时在指定的工作组内的可用工作项之间分发计算。每个工作组内的每个工作项都独立且并发地执行 lambda 表达式。

Work Group:

• 概念： 将工作组视为并行的工作单元集合。工作组中的每个工作单元（工作项）都可以执行独立的计算。
• 实际执行： 工作组内的工作项可以并行执行，每个工作项独立完成其任务。

parallel_for:

• 概念： 类似于告诉整个工作组同时执行特定任务。工作组内的每个工作项执行其任务的一部分，但它们在整个工作组范围内协同工作。
• 实际执行： 每个工作组内的工作项独立执行，并共同完成任务。

nd_range 和 work_group 在代码中的使用及参数解释:

• nd_range： 定义了并行计算的全局和局部维度。它指定了全局工作项的范围和局部工作组的大小。
  • 在代码片段中使用了 nd_range<2>{{M, N}, {1, tile_size}}，其中 {M, N} 定义了二维全局工作项范围，{1, tile_size} 指定了局部工作组的大小。
• work_group： 是并行性的概念单元，将多个工作项组合在一起。它代表一组可以在本地上下文中协作和共享数据的工作项。
  • 工作组的大小在 nd_range 中定义。在提供的代码中，{1, tile_size} 指定了工作组在第二维上的大小。

全局 ND-range 和 局部 ND-range

1. 全局 ND-range：
   • 全局 ND-range 表示整个问题的分布范围，即在全局范围内有多少个工作项。
   • 通常，全局 ND-range 用于划分任务的总体规模，例如，表示计算矩阵乘法的总大小，其中包含了所有需要计算的元素。
   • 在二维矩阵乘法的例子中，全局 ND-range 的第一个维度可能表示矩阵的行数，第二个维度表示列数。
2. 局部 ND-range：
   • 局部 ND-range 表示每个工作组（work-group）的分布范围，即在每个工作组内有多少个工作项。
   • 工作组是由多个工作项组成的，这些工作项可以在同一设备上的计算单元（例如处理器核心）上并行执行。
   • 局部 ND-range 的设置通常与硬件架构和任务的性质有关。例如，可以选择一定数量的工作项组成一个工作组，以适应硬件上的并行性。
   • 局部 ND-range 通常用于定义每个工作组的工作项的分布，以便它们可以协同工作并共享局部内存

下面是一个使用 SYCL（DPC++）实现的矩阵乘法并行算法。该算法旨在充分利用异构计算资源，例如 GPU。我们将对其进行解析和说明：

#include <chrono>
#include <iostream>
#include <sycl/sycl.hpp>

#define RANDOM_FLOAT() (rand() / static_cast<double>(RAND_MAX))

using namespace sycl;

const int TILE_DIM_X = 8;
const int TILE_DIM_Y = 8;

double matrix_mult_parallel(float *matA, float *matB, float *matC, 
                             int dimM, int dimN, int dimK, 
                             int BLOCK_SIZE, sycl::queue &device_queue) {

  auto grid_rows = dimM / TILE_DIM_Y;
  auto grid_cols = dimN / TILE_DIM_X;
  auto local_range  = range<2>(BLOCK_SIZE, BLOCK_SIZE);
  auto global_range = range<2>(grid_rows, grid_cols);

  double elapsed_time = 0.0;

  // 提交 SYCL 任务
  auto event = device_queue.submit([&](sycl::handler &handler) {
    handler.parallel_for(
      sycl::nd_range<2>(global_range, local_range), [=](sycl::nd_item<2> item_idx) {

        int row = TILE_DIM_Y * item_idx.get_global_id(0);
        int col = TILE_DIM_X * item_idx.get_global_id(1);

        float partial_sum[TILE_DIM_Y][TILE_DIM_X] = {0.0f};
        float localA[TILE_DIM_Y] = {0.0f};
        float localB[TILE_DIM_X] = {0.0f};

        // 进行矩阵乘法
        for (int k = 0; k < dimN; k++) {

          // 将数据加载到局部内存
          for(int m = 0; m < TILE_DIM_Y; m++) {
            localA[m] = matA[(row + m) * dimN + k];
          } 

          for(int p = 0; p < TILE_DIM_X; p++) {
            localB[p] = matB[k * dimN + p + col];
          } 

          // 计算局部乘积
          for (int m = 0; m < TILE_DIM_Y; m++) {
            for (int p = 0; p < TILE_DIM_X; p++) {
              partial_sum[m][p] += localA[m] * localB[p];
            }
          }
        }

        // 将局部结果写回全局内存
        for (int m = 0; m < TILE_DIM_Y; m++) {
          for (int p = 0; p < TILE_DIM_X; p++) {
            matC[(row + m) * dimN + col + p] = partial_sum[m][p];
          }
        }

      });
  });

  // 等待任务完成
  event.wait();

  // 计算任务执行时间
  elapsed_time += (event.get_profiling_info<info::event_profiling::command_end>() -
                   event.get_profiling_info<info::event_profiling::command_start>()) / 1000.0f / 1000.0f;

  return elapsed_time;
}

1. 全局 ND-range 和 局部 ND-range： 使用 nd_range 定义了全局和局部的工作项范围，以便在全局范围内有多少个工作项，每个工作组内有多少个工作项。

2. parallel_for： 利用 SYCL 的 parallel_for 函数，指示在全局 ND-range 内的工作组上执行特定任务。

3. 数据加载和计算： 在每个工作组内，使用局部内存加载矩阵的部分数据，以及计算局部乘积。

4. 局部结果写回： 将局部乘积的结果写回全局内存。

5. 任务执行时间： 使用 SYCL 事件计算任务执行的时间。

这个算法的目标是在异构计算平台上实现高效的矩阵乘法，通过分块和使用局部内存，以及并行计算的方式来提高计算性能。