基于英特尔oneAPI的图像卷积并行加速项目

 问题陈述

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

项目简介

Intel oneAPI是Intel推出的一套工具和库的集合，旨在支持异构计算和高性能计算 。它旨在简化和加速开发者在不同硬件架构上的应用程序，并提供对多种处理器和加速器的支持，包括CPU、GPU、FPGA等。oneAPI的目标是提供一个统一的编程模型，使开发者能够更轻松地利用不同类型的硬件来实现更高的性能。

oneAPI所提供的一系列工具、库和编程模型其中之一就是SYCL。SYCL是一种用于编写并行C++代码的开放标准，旨在简化和统一针对异构计算的编程。SYCL的目标是使开发者能够轻松地将C++代码用于不同类型的处理器，如CPU、GPU、FPGA等。

图像卷积是一种常见的图像处理操作，用于应用各种滤波器和特征检测器。其原理可以简单地描述为在图像的每个像素上应用一个小的矩阵（通常称为卷积核或滤波器），并将卷积核中的元素与图像中对应位置的像素值相乘，然后将所有乘积的和作为结果。这个过程可以看作是对图像进行了平滑、锐化、边缘检测等操作。假设有⼀个大小为M × N 的输入图像I 和一个大小为m × n 的卷积核 K 。图像卷积操作可以用下面的数学公式来表示：

其中， S(i, j)是卷积操作的结果图像中位置 (i, j) 处的像素值。 I(i + k, j + l) 是图像中位置 (i + k, j + l) 处的像素值， K(k, l) 是卷积核中位置 (k, l) 处的权重。卷积核通常是一个小的⼆维矩阵，用于捕捉图像中的特定特征。在异构计算编程中，可以使用并行计算来加速图像卷积操作。SYCL可以将各个像素的卷积计算交给不同计算单元并行执行，来加速图像处理过程。

相关技术

SYCL编程模型

oneAPI中支持SYCL编程模型的C++编译器

英特尔oneAPI Developer Cloud 服务

代码

#include <sycl/sycl.hpp>
#include <iomanip>
#include <iostream>
#include <fstream>
#include <string.h>
#include <math.h>

using namespace sycl ;
using namespace std;

// 定义卷积核大小
const int KERNEL_SIZE = 5;

// 定义图像大小
const int IMAGE_WIDTH = 100;
const int IMAGE_HEIGHT = 100;

// 定义图像矩阵和卷积核矩阵
std::vector<double> img(IMAGE_WIDTH*IMAGE_HEIGHT, 0.0f);
std::vector<double> krn={370.53,898.7,597.57,537.95,691.58,
552.51,995.54,208.58,445.72,261.22,
106.74,435.67,909.39,849.86,457.39,
299.15,856.77,32.85,893.19,874.23,
82.89,509.09,659.11,877.06,959.67};
//定义结果矩阵
std::vector<double> output((IMAGE_WIDTH - KERNEL_SIZE + 1)*(IMAGE_HEIGHT - KERNEL_SIZE + 1), 0.0f);

int main() {   
    // 创建设备队列
    queue queue;

    // 定义缓冲区并复制数据到设备
    sycl::buffer<double, 2> imageBuffer(img.data(), sycl::range<2>(IMAGE_WIDTH, IMAGE_HEIGHT));
    sycl::buffer<double, 2> kernelBuffer(krn.data(), sycl::range<2>(KERNEL_SIZE, KERNEL_SIZE));
    sycl::buffer<double, 2> outputBuffer(output.data(), sycl::range<2>(IMAGE_WIDTH - KERNEL_SIZE + 1, IMAGE_HEIGHT - KERNEL_SIZE + 1));

    // 执行卷积操作的内核函数
    queue.submit([&](handler& cgh) {
        auto inputImage = imageBuffer.get_access<sycl::access::mode::read>(cgh);
        auto convolutionKernel = kernelBuffer.get_access<sycl::access::mode::read>(cgh);
        auto outputImage = outputBuffer.get_access<sycl::access::mode::write>(cgh);

        cgh.parallel_for<class ConvolutionKernel>(range<2>(IMAGE_WIDTH - KERNEL_SIZE + 1, IMAGE_HEIGHT - KERNEL_SIZE + 1), [=](id<2> id) {
                double sum = 0;
                for (int i = -ceil(KERNEL_SIZE/2); i < KERNEL_SIZE/2; ++i) {
                    for (int j = -ceil(KERNEL_SIZE/2); j < KERNEL_SIZE/2; ++j) {
                        sum += inputImage[id[0] + i][id[1] + j] * convolutionKernel[i+KERNEL_SIZE/2][j+KERNEL_SIZE/2];
                    }
                }
                outputImage[id] = sum;
        });
    });

    // 等待队列执行完成
    queue.wait();

    string filename = R"(output.txt)";
    ofstream fout;
    fout.open(filename,ios::app);
    // 保存卷积后的图像
    for (int i = 0; i < IMAGE_WIDTH - KERNEL_SIZE + 1; ++i) {
        for (int j = 0; j < IMAGE_HEIGHT - KERNEL_SIZE + 1; ++j) {
            fout << output[i*(IMAGE_WIDTH - KERNEL_SIZE + 1)+j] <<std::endl;
        }
    }

    return 0;
}

运行

使用英特尔oneAPI Developer Cloud 服务运行，不需要额外配置运行环境。在英特尔oneAPI Developer Cloud 服务中，启动Jupyter服务，可以直接在平台上编写C++/SYCL程序并运行。我们可以编写一个shell脚本记录运行时间、编译并运行程序。

#!/bin/bash
echo
echo start: $(date "+%y%m%d.%H%M%S.%3N")
echo

source /opt/intel/oneapi/setvars.sh > /dev/null 2>&1
icpx -fsycl ./convolution.cpp 
if [ $? -eq 0 ]; then ./a.out; fi

echo
echo stop:  $(date "+%y%m%d.%H%M%S.%3N")
echo

通过另外编写一个未加速的卷积程序，对比得出oneAPI为给定的黑客松基准数据集加速约30%。

总结

在本次项目中，我们学到了Intel oneAPI的概念，以及oneAPI所提供的其中一种编程模型SYCL的概念、编程方法以及底层原理。我们还了解了SYCL如何通过并行计算对卷积操作实现加速。