问题陈述
使用基于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如何通过并行计算对卷积操作实现加速。