基于oneAPI的SYCL图像卷积加速实现

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

技术方案
(1)oneAPI、SYCL介绍

oneAPI 是一个由英特尔主导的开放、全面的软件栈,旨在简化多核和异构计算的编程。oneAPI 提供了一个统一的编程模型,可以用于不同类型的处理器,包括 CPU、GPU、FPGA 和其他加速器。通过 oneAPI,开发人员可以更轻松地利用各种处理器的性能优势,同时减少针对特定处理器类型的编程复杂性。

SYCL(pronounced “sickle”)是一种基于 C++ 的编程模型,用于实现异构计算。SYCL 的目标是提供一种简单且高性能的编程模型,以实现 CPU、GPU 和其他加速器之间的无缝协同。SYCL 基于标准的 C++,并引入了许多用于表示并行性和内存关系的新特性。它使开发人员能够使用标准的 C++ 代码编写并行程序,而无需深入了解特定硬件体系结构。

总的来说,oneAPI 为异构计算提供了统一的编程模型和工具集,而 SYCL 则是其中的一部分,专注于提供基于 C++ 的高性能并行编程模型。这些工具和模型使开发人员能够更轻松地利用各种处理器的潜力,并加快异构计算应用程序的开发和部署速度。

(2)图像卷积介绍

图像卷积是一种基本的图像处理技术,用于检测图像中的特征、边缘、纹理等。原理可以简单地描述为在图像的每个像素上应用一个小的矩阵(通常称为卷积核或滤波器),并将卷积核中的元素与图像中对应位置的像素值相乘,然后将所有乘积的和作为结果。其数学公式可以表示如下:

​ G ( i , j ) = ∑ m ∑ n I ( i + m , j + n ) ⋅ K ( m , n )

其中,G ( i , j )是输出图像中位置 ( i , j )处的像素值;I ( i + m , j + n )  是输入图像中位置 ;( i + m , j + n ) 处的像素值 ,K ( m , n )卷积核中位置 ( m , n ) 处的权重值。

(3)基于oneAPI的SYCL图像卷积加速

加速优化的思路是使用SYCL来实现并行计算。分析可知,不加速的卷积操作是通过多重循环实现的,每个像素需要依次遍历滤波器。而使用了SYCL的并行计算能力,可以通过提交到队列中的kernel函数,将卷积操作分配给多个处理单元并行执行,从而加速计算过程。

代码实现
为了验证效果,分别实现了传统的和加速的卷积代码

传统卷积
%%writefile lab/filter_classic.cpp
//==============================================================
// Copyright © Intel Corporation
//
// SPDX-License-Identifier: MIT
// =============================================================
#include <array>
#include <iostream>
#include <chrono>


constexpr int image_size = 100;
constexpr int filter_size = 5;

#include <fstream>
#include <sstream>

// 从文件中加载图像数据
template <size_t N>
std::array<float, N*N> loadImageFromFile(const std::string& filename) {
  std::ifstream file(filename);
  std::string line;
  std::array<float, N*N> data = { 0 };

  if (file.is_open()) {
    for (int i = 0; i < N*N; ++i) {
      if (std::getline(file, line)) {
        std::stringstream ss(line);
        ss >> data[i];
      } else {
        // 处理文件格式错误的情况
      }
    }

    file.close();
  } else {
    std::cerr << "无法打开文件: " << filename << std::endl;
    // 处理文件打开失败的情况
  }

  return data;
}


int main() {
  // 初始化图像和滤波器
  std::array<float, image_size*image_size> image = loadImageFromFile<image_size>("./data/image.txt");
  std::array<float, filter_size*filter_size> filter = loadImageFromFile<filter_size>("./data/filter.txt");
  std::array<float, image_size*image_size> result = {0};
  
  auto start = std::chrono::high_resolution_clock::now(); // 获取开始时间
 
  // 卷积操作
  for (int x = 0; x < image_size; ++x) {
    for (int y = 0; y < image_size; ++y) {
     float sum = 0;
      for (int i = -filter_size/2; i <= filter_size/2; ++i) {
        for (int j = -filter_size/2; j <= filter_size/2; ++j) {
          if (x+i >= 0 && x+i < image_size && y+j >= 0 && y+j < image_size) {
            sum += image[(x+i)*image_size + (y+j)] * filter[(i+filter_size/2)*filter_size + (j+filter_size/2)];
          }
        }
      }
      result[x*image_size + y] = sum;
    }
  }

  auto end = std::chrono::high_resolution_clock::now(); // 获取结束时间
  auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(end - start);
  std::cout << "运行时间: " << duration.count() << " 毫秒" << std::endl;

  return 0;
}

基于oneAPI的SYCL加速卷积
基于上文提到的思路,具体实现如下:

在加速的版本中,首先创建一个名为myQueue的队列对象,并使用数组数据初始化sycl::buffer对象。然后,在队列上提交一个lambda表达式,该表达式包含用于并行计算的kernel函数(parallel_for),其中每个线程块都会为图像中的一个像素点计算卷积值。最后,等待队列中的所有任务完成,即可获取计算结果。

%%writefile lab/filter.cpp
//==============================================================
// Copyright © Intel Corporation
//
// SPDX-License-Identifier: MIT
// =============================================================
#include <CL/sycl.hpp>
#include <array>
#include <iostream>
#include <chrono>

constexpr int image_size = 100;
constexpr int filter_size = 5;

#include <fstream>
#include <sstream>

// 从文件中加载图像数据
template <size_t N>
std::array<float, N*N> loadImageFromFile(const std::string& filename) {
  std::ifstream file(filename);
  std::string line;
  std::array<float, N*N> data = { 0 };

  if (file.is_open()) {
    for (int i = 0; i < N*N; ++i) {
      if (std::getline(file, line)) {
        std::stringstream ss(line);
        ss >> data[i];
      } else {
        // 处理文件格式错误的情况
      }
    }

    file.close();
  } else {
    std::cerr << "无法打开文件: " << filename << std::endl;
    // 处理文件打开失败的情况
  }

  return data;
}


int main() {
  // 初始化图像和滤波器
  std::array<float, image_size*image_size> image = loadImageFromFile<image_size>("./data/image.txt");
  std::array<float, filter_size*filter_size> filter = loadImageFromFile<filter_size>("./data/filter.txt");
  std::array<float, image_size*image_size> result = {0};

    
    
  auto start = std::chrono::high_resolution_clock::now(); // 获取开始时间
    
  try {
    sycl::queue myQueue;

    sycl::buffer<float, 1> image_buffer(image.data(), image.size());
    sycl::buffer<float, 1> filter_buffer(filter.data(), filter.size());
    sycl::buffer<float, 1> result_buffer(result.data(), result.size());

    myQueue.submit([&](sycl::handler& cgh) {
      auto image_accessor = image_buffer.get_access<sycl::access::mode::read>(cgh);
      auto filter_accessor = filter_buffer.get_access<sycl::access::mode::read>(cgh);
      auto result_accessor = result_buffer.get_access<sycl::access::mode::write>(cgh);

      cgh.parallel_for<class Convolution>(sycl::range<2>(image_size, image_size), [=](sycl::id<2> idx) {
        int x = idx[0];
        int y = idx[1];
        float sum = 0;
        for (int i = -filter_size/2; i <= filter_size/2; ++i) {
          for (int j = -filter_size/2; j <= filter_size/2; ++j) {
            if (x+i >= 0 && x+i < image_size && y+j >= 0 && y+j < image_size) {
              sum += image_accessor[(x+i)*image_size + (y+j)] * filter_accessor[(i+filter_size/2)*filter_size + (j+filter_size/2)];
            }
          }
        }
        result_accessor[x*image_size + y] = sum;
      });
    });

    myQueue.wait();
  } catch (sycl::exception const& e) {
    std::cerr << "An exception occurred: " << e.what() << std::endl;
    return 1;
  }
  
  auto end = std::chrono::high_resolution_clock::now(); // 获取结束时间
  auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(end - start);
  std::cout << "运行时间: " << duration.count() << " 毫秒" << std::endl;

  return 0;
}


实验对比
(1)运行截图

共设计4组随机生成的不同尺寸的图像(n X n)和卷积核(m X m),加速前后均得到一致结果。

解释: 这里没有使用黑客松的数据,因为黑客松的数据尺寸太小,无法体现oneapi的优势。

以下展示加速卷积在n=1000,m=5的运行截图:

(2)效率对比

输入矩阵尺寸(n、m)    传统卷积    基于oneAPI的SYCL加速卷积
100,5    5ms    1543ms
200,5    18ms    1797ms
1000,5    287ms    1806ms
200,50    1059ms    1932ms
1000,50    28618ms    2112ms
运行平台:Intel® DevCloud

注:

1、由于使用Devcloud平台,计算资源为动态分配,表中每组数据均为重复5次取平均数

2、由于平台差异,传统卷积得到的计算时间或与其他平台(如本机)可能存在较大的差距,但不影响同平台下的加速效果比较。

实验总结
根据给出的数据,可以看出在不同的图像大小和卷积核大小下,传统卷积和基于oneAPI的SYCL加速卷积的表现是不同的。在这些数据中,我们可以得出一些初步的结论:

当图像尺寸较小时,传统卷积比基于oneAPI的SYCL加速卷积更快。例如,在100x5和200x5的情况下,传统卷积的运行时间远远小于基于SYCL的加速卷积。
当图像尺寸增大时,基于oneAPI的SYCL加速卷积的运行时间相对稳定,而传统卷积的运行时间大幅增加。例如,在1000x5和1000x50的情况下,传统卷积的运行时间显著增加,而基于SYCL的加速卷积的运行时间变化不大。
卷积核大小对结果也有影响。在300x5和300x50的情况下,传统卷积和基于oneAPI的SYCL加速卷积的运行时间都有所增加,在给定的代码中,卷积核大小对结果的影响是通过循环实现的。具体地,对于每个像素点,卷积操作都会在以该像素为中心的filter_size x filter_size 的窗口内进行。因此,当卷积核大小变大时,窗口内的像素数量也随之增加。随着卷积核大小的增加,需要计算的像素数量也随之增加,这会导致计算时间的增加。另外,在窗口内进行卷积操作也需要更多的内存访问,这可能会导致更多的缓存未命中和访问延迟,从而影响性能。因此,当卷积核大小增加时,计算时间和内存访问量都会增加,从而影响性能。但是在不同的卷积核大小下,基于SYCL的加速卷积的运行时间仍然相对稳定。
根据以上分析,可以得出以下结论:

在图像尺寸较小时,传统卷积可能更快,因为启动并配置并行计算所需的开销可能会超过串行计算的性能优势。
随着图像尺寸增大,基于oneAPI的SYCL加速卷积开始展现出明显的优势,因为并行计算能够更好地发挥在大规模问题上的优势,而传统的串行计算则受到了严重的性能瓶颈。
卷积核大小对结果也有影响,但是在给定的数据中并没有明显的规律可循。
综上所述,相比于不加速的版本,加速版本使用oneAPI,通过利用GPU等硬件平台的并行计算能力,可以显著提高计算性能,特别是对于大规模的图像数据和复杂的卷积核来说,加速的优势更为明显。
 

  • 0
    点赞
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值