使用基于oneAPI的C++/SYCL实现用于计算图像的卷积操作

描述

使用基于oneAPI的C++/SYCL实现一个用于计算图像的卷积操作。输⼊为一个图像矩阵和一个卷积核矩阵，输出为卷积后的图像。

分析

图像卷积是一种常见的图像处理操作，用于应用各种滤波器和特征检测器。其原理可以简单地描述为在图像的每个像素上应用一个小的矩阵（通常称为卷积核或滤波器），并将卷积核中的元素与图像中对应位置的像素值相乘，然后将所有乘积的和作为结果。这个过程可以看作是对图像进行了平滑、锐化、边缘检测等操作。
假设有⼀个大小为M × N 的输入图像 I 和一个大小为 m × n 的卷积核 K 。图像卷积操作可以用下面的数学公式来表示：
$$S(i,j)={\Sigma }_k{\Sigma }_{l}I(i+k,j+l)\cdot K(k,l)$$
其中，
S(i, j)是卷积操作的结果图像中位置 (i, j) 处的像素值,
I(i+k, j+l) 是图像中位置 (i+k, j+l) 处的像素值,
K(k, l) 是卷积核中位置 (k, l) 处的权重。
卷积核通常是一个小的⼆维矩阵，用于捕捉图像中的特定特征。在异构计算编程中，可以使用并行计算来加速图像卷积操作。通过将图像分割成小块，然后在GPU上并行处理这些块，可以实现高效的图像卷积计算。通过合理的块大小和线程组织方式，可以最大限度地利用GPU的并行计算能力来加速图像处理过程。
基于GPU的图像卷积操作的原理基于并行处理和矩阵乘法的基本原理，通过将图像数据和卷积核数据分配给不同的线程块和线程，利用GPU的并行计算能力实现对图像的快速处理。

代码

#include <CL/sycl.hpp>
#include <iostream>
#include <vector>

constexpr size_t KERNEL_SIZE = 5;   // 卷积核尺寸
constexpr size_t IMAGE_WIDTH = 100;   // 图像宽度
constexpr size_t IMAGE_HEIGHT = 100;  // 图像高度

// 定义卷积核
const float kernel[9]={1,0,1,2,0,-2,1,0,1};


// 定义输入图像
const std::vector<float> inputImage = {1, 2, 3, 4,
                                       5, 6, 7, 8,
                                       9, 10, 11, 12,
                                       13, 14, 15, 16};

// 定义输出图像
std::vector<float> outputImage(IMAGE_WIDTH * IMAGE_HEIGHT);

int main() {
    
  // 创建一个队列来执行DPC++任务
  cl::sycl::queue queue;

  // 创建输入和输出缓冲区
  cl::sycl::buffer<float, 2> inputBuffer(inputImage.data(),
                                         cl::sycl::range<2>(IMAGE_HEIGHT, IMAGE_WIDTH));
  cl::sycl::buffer<float, 2> outputBuffer(outputImage.data(),
                                          cl::sycl::range<2>(IMAGE_HEIGHT, IMAGE_WIDTH));

  // 提交DPC++任务
  queue.submit([&](cl::sycl::handler& cgh) {
    // 获取输入和输出访问器
    auto inputAccessor = inputBuffer.get_access<cl::sycl::access::mode::read>(cgh);
    auto outputAccessor = outputBuffer.get_access<cl::sycl::access::mode::write>(cgh);

    // 定义内核函数
    cgh.parallel_for<class ConvolutionKernel>(cl::sycl::range<2>(IMAGE_HEIGHT, IMAGE_WIDTH),
                                              [=](cl::sycl::item<2> item) {
      size_t row = item[0];
      size_t col = item[1];

      // 执行卷积操作
      float result = 0.0f;
      for (int i = -1; i <= 1; ++i) {
        for (int j = -1; j <= 1; ++j) {
          if (row + i >= 0 && row + i < IMAGE_HEIGHT &&
              col + j >= 0 && col + j < IMAGE_WIDTH) {
            float pixel = inputAccessor[row + i][col + j];
            float kernelValue = kernel[(i + 1) * KERNEL_SIZE + (j + 1)];
            result += pixel * kernelValue;
          }
        }
      }
      outputAccessor[item] = result;
    });
  });

  // 将输出图像从设备内存复制到主机内存
  queue.wait_and_throw();
  queue.submit([&](cl::sycl::handler& cgh) {
    auto outputAccessor = outputBuffer.get_access<cl::sycl::access::mode::read>(cgh);
    cgh.copy(outputAccessor, outputImage.data());
  });
  queue.wait_and_throw();

  // 打印输出图像
  std::cout << "Output Image:\n";
  for (size_t i = 0; i < IMAGE_HEIGHT; ++i) {
    for (size_t j = 0; j < IMAGE_WIDTH; ++j) {
      std::cout << outputImage[i * IMAGE_WIDTH + j] << " ";
    }
    std::cout << "\n";
  }

  return 0;
}

验证

针对黑客松基准数据集在oneapi dev cloud跑了结果；

为验证正确性，构造了一个小的测试样例如下：

// 定义卷积核
const float kernel[9]={1,0,1,2,0,-2,1,0,1};


// 定义输入图像
const std::vector<float> inputImage = {1, 2, 3, 4,
                                       5, 6, 7, 8,
                                       9, 10, 11, 12,
                                       13, 14, 15, 16};

结果如下，与预期相符：

黑客松基准数据集的结果如下：

声明

参考：https://zhuanlan.zhihu.com/p/634628468