1摘要
在devcloud云服务器上利用SYCL实现了sobel算子的边沿检测算法并利用并发算法对sobel算子进行加速,算法在Intel® Xeon® Gold 6128 CPU上运行。收集了5类数据,测试cpu多核和单核执行sobel算子的时间,并计算了加速比,也测试了调用不同cpu核数对类函数执行sobel算子和执行用函数执行sobel算子的影响。也测试了不同编译器对单核代码的影响。
2运行环境
作者在devcloud[1]上用sycl[2]实现sobel算子[3],[4].[5],[6],[7],[8]的边沿检测[9],[10],[11],[12],[13],[14],[15],需要的硬件有cpu、内存和硬盘。表1展示了cpu的信息,其中可以同时运行12个核数,因此可以创建12个sobel算子并行检测图片边沿。
表1.CPU详细信息:
型号: Intel® Xeon® Gold 6128 CPU @ 3.40GHz
核数: 6
同时运行线程数: 12
L1d cache: 384 KiB
浮点数运算能力: 652.8Gflops
L1i cache :384 KiB
L2 cache :12 MiB
L3 cache :38.5 MiB
CPU的最大主频 :3700 MHz
CPU的最小主频 :1200 MHz
浮点数计算能力 :652.8Gflops
内存详细信息:
图1.内存信息
free -h 命令用于显示系统内存的使用情况。在图1中中,我们可以看到以下几个部分:
总内存(total):表示系统总内存,单位是 Gi(Gigabytes)。如图所示系统总内存为 188 Gi。
已使用内存(used):这部分显示的是已经被应用程序使用的内存。如图可知,已使用的内存为 4.2 Gi。
空闲内存(free):这部分显示的是当前未被使用的内存。如图可知,空闲内存为 181 Gi。
共享内存(shared):这是被两个或更多进程共享的内存。如图可知,共享内存为 2.0 Mi (Mebibytes)。
缓冲区和缓存(buff/cache):这部分内存被用作文件系统缓存和用于存储最近打开的文件数据。如图可知,缓冲区和缓存的内存为 2.9 Gi。
可用的内存(available):这部分显示的是在不损害正在运行的进程的前提下,可以立即分配的内存。如图可知,可用的内存为 182 Gi。
随机存储器(Mem)是指计算机的物理内存,也称为RAM(Random Access Memory)。它用于存储当前正在运行的应用程序和操作系统数据。
交换空间(Swap):这是系统的交换空间,即当物理内存用尽时,系统将使用的备用内存。在这个例子中,交换空间为 1.9 Gi。
硬盘相信信息:
图2硬盘信息
NAME: 这是设备的名称,例如 sda、sda1、sda2 等。
MAJ:MIN: 这表示设备号码。
RM: 表示设备是否可移除,这里所有的设备都是固定的,所以这一列都是0。
SIZE: 这是设备的总大小,以GB为单位。
RO: 表示设备是否只读,这里所有的设备都是可写的,所以这一列都是0。
TYPE: 这是设备的类型,例如 disk、part、lvm等。
MOUNTPOINT: 这是设备挂载的点,如果设备没有挂载,这一列就是空。
系统中有一个名为 sda 的磁盘,它的大小是447.1G。这个磁盘被分成了几个分区:sda1、sda2、sda5和sda6。其中sda6又被分成了几个逻辑分区:main-root、main-var和main-local。
/boot 分区(sda1)的大小是976M,它是用来存放操作系统的引导文件的。
[SWAP] 分区(sda5)是用于交换的分区,大小为1.9G。在大多数情况下,操作系统的运行需要用到交换分区。当物理内存用完时,一些暂时不用的进程就会被移到交换分区上,从而释放物理内存的空间。
/ 分区(main-root)的大小是119.2G,这是操作系统的根分区,几乎所有的文件都存放在这里。
/var 分区(main-var)的大小是37.3G,这个分区主要用于存放经常变动的文件,例如日志文件、临时文件等。
/local 分区(main-local)的大小是287.8G,这个分区可以用来存放一些较大的、不经常变动的文件,例如备份文件、大型程序等。
软件环境:
操作系统:Ubuntu,版本Ubuntu 20.04.6 LTS。
编译器:Intel® oneAPI DPC++/C++ Compiler,版本2023.2.0。
开发语言:C++,版本号C++17。
辅助软件
cmath:用于调用函数进行计算如开方函数以及平方函数。
sycl:编写cpu多核并行的编程模型。
fstream:读取存储图片路径的txt文件。
string:存储图片路径。
3.DPC++实现sobel算子的边沿检测
SYCL介绍
SYCL 源于 OpenCL 的高级编程模型,是一种单源的、高级的编程模型与规范。单源即它允许使用同一个 C++源文件编写主机 CPU 和设备代码,但仍需要分别 对主机和设备代码进行编译。其高级在于它提供 API 和抽象层:用于查找可以执行代码的设备、缓冲区创建和数据移动、内核函数编译、依赖管理和调度以及事件管理等功能。标准在于它允许使用现代的 ISO C++标准,和 CUDA,OpenMP 等编程模型不同,它的学习 成本最低,同时能使用 C++生态工具进行扩展。 它允许开发者使用多种异构设备,支持 CPU、GPU、FPGA、AI 和定制芯片等。
DPC++介绍
为了应对异构计算中的挑战,Intel 推出了其新开发框架 oneAPI。具体来说,oneAPI旨在推进异构计算技术的发展,并消除专有编程模型所带来的经济和技术障碍的一套开发框架。DPC++是 Data Parallel C++的首字母缩写,它是 Intel 为了将 SYCL 引入 LLVM 和oneAPI 所开发的开源项目。该框架基于标准的 ISO C++语言,并在 SYCL 规范的基础上进行了扩展,其中包括统一共享内存等特性已被纳入 SYCL2020 规范。
Sobel算子的边沿检测算法
Sobel算子是计算机视觉领域[16],[17],[18]的一种重要处理方法。主要用于获得数字图像的一阶梯度,常见的应用和物理意义是边缘检测。索贝尔算子是把图像中每个像素的上下左右四领域的灰度值加权差,在边缘处达到极值从而检测边缘。
图3.原始图片
图4.边沿检测结果
该算子包含两组3x3的卷积核,分别为横向及纵向,将之与图像作卷积,即可分别得出横向及纵向的亮度差分近似值。如果以A代表原始图像,Gx及Gy分别代表经横向及纵向边缘检测的图像,其公式如下:
图5.横向和纵向边沿检测公式
Gxx,y表示Gx图像在x行,y列点的像素值,Gyx,y表示Gy图像在x行,y列点的像素值,f(x,y)表示A图像x行y列像素值。
利用sobel水平算子对下图进行检测。
图6.原始图片
图7.水平边沿检测结果
利用sobel垂直算子对下图进行检测。
图8.原始图片
图9.垂直边沿检测结果
图像的每一个像素的横向及纵向梯度近似值可用以下的公式结合,来计算梯度的大小。
Gxx,y是Gx的第x行第j列像素值,Gyx,y是Gy第x行第j列像素值,Gx,y表示G图像的第x行第j列像素值。
图10.sobel算子的最终检测结果
Sobel算子串行程序设计的代码
ApplyFilter函数的伪代码如下:
y=sobel卷积核中心位置所处行数
x=sobel卷积核中心位置所处列数
temp =
temp1 =
temp2=(temptemp+temp1temp1)0.5;
dst_image[i] = temp2 > 255 ? 255 : temp2;}
ApplyFilter是单次sobel算子提取边沿的过程,其中参数src_image表示被提取边沿的图片,dst_image表示存储边沿特征的图片,i表示提取位置(图片以一维数组存储的索引值),将其转化为i对应的行值和列值,获取Gxx,y和Gyx,y所需像素的值进行卷积操作。
for(int i=0;i<G宽G高;i++){ApplyFilter(image, image_ref, i);}
利用for循环不断调用ApplyFilter对image图像进行单次sobel算子的卷积操作。
并发设计图
sobel算子的边沿检测算法是用sobel算子串行检测图片边沿,串行处理导致该算法的处理效率较低,为了提高sobel算子的边沿检测效率,创建了多个独立的sobel算子并行进行图片的边沿检测算法。图展示了横向边沿检测Gx算子并发对图像的处理,图展示了纵向边沿检测Gy算子并发对图像处理。
图11.并行实现水平边沿检测原理图
图12.并行实现垂直边沿检测原理图
并发程序设计流程图
1.利用stbi_load函数读取图片,获取图片的长和宽。2.利用new uint8_t[(图像长-2)(图像宽-2)]创建存储空间,并将初始值设置为0,用来存储边沿特征信息。3.利用queue创建队列并为其分配cpu。4.利用buffer将数据放入缓冲区中,以便在执行并行操作时能争取访问数据。5.利用queue.submit函数创建多个工作项实现sobel算子的并行边沿检测。6.利用stbi_write_png保存图片。7.利用get_profiling_infoinfo::event_profiling::command_end和get_profiling_infoinfo::event_profiling::command_start计算结束并发结束时间。
图13.并行程序流程图
实验结果
数据集的准备
为了测试并发算法的效果,作者收集39张图片水果和山水图构建了像素512512数据集,像素10241024数据集,像素2048像素2048数据集,像素4096像素4096数据集。收集663张图片水果和山水图作为宽6144像素和高6144像素的图片。
像素512512:39张宽为512像素,高为512像素的图片。
像素10241024:39张宽为1024像素,高为1024像素的图片。
像素20482048:39张宽为2048像素,高为2048像素的图片。
像素40964096:39张宽为4096像素,高为4096像素的图片。
像素61446144:663张宽为6144像素,高为6144像素的图片。
并行算法的准备
作者也编写了俩种调用算法,并测试他们的并发效率。
算法一h.parallel_for(range<1>(img_widthimg_high), [=](auto i) {
ApplyFilter(image_acc.get_pointer(), image_exp_acc.get_pointer(), i);});
算法二Sobel kernel(image_acc, image_exp_acc);
h.parallel_for(range<1>(img_width*img_high), kernel);
算法一执行调用sobel滤波函数进行并发操作,算法二先创建sobel类,然后用parallel_for函数调用类函数实现并发算法。
实验结果
表2.单核时间,算法一和算法二在数据集的运行时间及加速比:
由表可知并发算法的加速比随数据规模的增大而增大,而调用类函数进行并发的加速比比直接调用函数进行并发的加速比有更好的效果。
表3.算法一不同并发数的执行时间
表4.算法二不同并发数的执行时间
由表可知创建的线程越多,并发算法所需的时间越少,而作者的服务器的cpu最多只能支持12核也就是12线程。
表5.不同编译器对单核算法的影响
作者发现icpx对单核算法具有优化性能,因此作者用icpx和g++对相同的单核代码进行编译,得到结果由表5所示,由表5可知icpx编译的代码比g++编译代码具有更快的运行时间。但通过查阅资料[19]可知icpx编译的代码不如g++,缺点是icpx编译方法不是很成熟例如某些特定情况下icpx编译的代码会出现bug,而g++编译的代码不会出现bug。
结论
本文在devcloud上实现了基于cpu多核的sobel算子的并行边沿检测算法。使边沿检测算法在Intel® Xeon® Gold 6128 CPU
@ 3.40GHz上运行,测试了函数并行算法(算法一)和类函数并行算法(算法二)在5种数据集下的加速比,发现算法二比算法一有更高的加速比,也测试了不同核数对算法一和算法二性能的影响,发现核数在12以内时,核数越大,算法处理数据集越快。实验发现icx编译器编译的代码比g++编译器编译的代码有更快的运行效率。
参考文献
[1]李斌.基于华为DevCloud的计算机软件项目实训课程创新性实践[J].信息技术与信息化,2019(12):30-32.
[2]丁越,徐传福,邱昊中等.基于SYCL的多相流LBM模拟跨平台异构并行计算研究[J].计算机科学,2023,50(11):
[3]张婧,郭风成,左泽丹等.融合多方向Sobel算子的相干斑各向异性扩散抑制[J].遥感技术与应用,2023,38(05):
[4]朴思儒,李彬,张赫等.基于多方向扩展Sobel加权算法的合作目标检焦方法研究[J].机电工程技术,2023,52(05):
[5]朴思儒,李彬,张赫等.基于多方向扩展Sobel加权算法的合作目标检焦方法研究[J].机电工程技术,2023,52(05):
[6]薛文格,邝天福.基于Sobel算子和灰色关联分析的图像边缘检测[J].楚雄师范学院学报,2023,38(03):
[7]程伟涛,郭来功,汪强.基于FPGA的多方向阈值自适应Sobel算法实现[J].无线互联科技,2023,20(09):
[8]徐乐意,黄鑫,陈玲玲等.基于Sobel算子级联地震属性的三角洲沉积相边缘识别[J].长江大学学报(自然科学版),2023,20(03):
[9]王辉,王晓红,周润民等.改进Canny的无人机影像边缘检测算法研究[J].智能计算机与应用,2023,13(09):
[10]张文喆.FPGA优化的快速边缘检测方法研究[J].单片机与嵌入式系统应用,2023,23(10):
[11]杨榆樟,王碧珺,梁洪健等.基于Pony边缘检测算法的隧道车辆识别研究[J].武汉理工大学学报(交通科学与工程版).
[12]刘超超,司亚超.引入注意力机制的整体嵌套边缘检测网络[J].河北建筑工程学院学报,2023,41(02):
[13]于晓,林世基,庄光耀等.基于多梯度融合的污水域污染物边缘提取算法研究[J].黑龙江工业学院学报(综合版),2023,23(06):
[14]吴陈易.基于Hadamard门和双链量子遗传算法的量子图像边缘检测方案[D].南昌大学,2023.
[15]王珊珊.基于FPGA的集装箱边缘检测算法研究[D].北方工业大学,2023.
[16]陆钟超,邱月,张安强等.机器视觉判别牛肉新鲜度的多模型定量分析[J].食品与发酵工业.
[17]曾飞,刘欣,王涛等.基于螺栓间弹性线曲率变化的螺栓松动视觉检测[J].传感器与微系统,2023,42(11):
[18]金涛,黄俊波,蔡澍雨等.基于机器视觉的无人机自动巡检定位控制技术[J].电子设计工程,2023,31(21):
[19]知乎链接https://www.zhihu.com/question/21675828/answers/updated
附录
```cpp
// =============================================================
#include <chrono>
#include <cmath>
#include <iostream>
#include <sycl/sycl.hpp>
#include "device_selector.hpp"
#include<fstream>
#include<string>
// dpc_common.hpp can be found in the dev-utilities include folder.
// e.g., $ONEAPI_ROOT/dev-utilities/<version>/include/dpc_common.hpp
#include "dpc_common.hpp"
// stb/*.h files can be found in the dev-utilities include folder.
// e.g., $ONEAPI_ROOT/dev-utilities/<version>/include/stb/*.h
#define STB_IMAGE_IMPLEMENTATION
#include "stb/stb_image.h"
#define STB_IMAGE_WRITE_IMPLEMENTATION
#include "stb/stb_image_write.h"
#include<iostream>
#include <vector>
using namespace std;
using namespace sycl;
// Few useful acronyms.
constexpr auto sycl_read = access::mode::read;
constexpr auto sycl_write = access::mode::write;
constexpr auto sycl_device = access::target::device;
static void ReportTime(const string &msg, event e) {
cl_ulong time_start =
e.get_profiling_info<info::event_profiling::command_start>();
cl_ulong time_end =
e.get_profiling_info<info::event_profiling::command_end>();
double elapsed = (time_end - time_start) / 1e6;
cout << msg << elapsed << " milliseconds\n";
}
// SYCL does not need any special mark-up for functions which are called from
// SYCL kernel and defined in the same compilation unit. SYCL compiler must be
// able to find the full call graph automatically.
// always_inline as calls are expensive on Gen GPU.
// Notes:
// - coeffs can be declared outside of the function, but still must be constant
// - SYCL compiler will automatically deduce the address space for the two
// pointers; sycl::multi_ptr specialization for particular address space
// can used for more control
__attribute__((always_inline)) static void ApplyFilter(uint8_t *src_image,
uint8_t *dst_image,
int i) {
int row,col;
row=int(i/6142);
col=i%6142;
float temp,temp1,temp2;
temp = (1.0f * src_image[row*6144+col]) + (2.0f * src_image[row*6144+col + 1]) +
(1.0f * src_image[row*6144+col + 2])+(0 * src_image[(row+1)*6144+col]) + (0 * src_image[(row+1)*6144+col + 1]) +
(0 * src_image[(row+1)*6144+col + 2])+(-1.0f * src_image[(row+2)*6144+col]) + (-2.0f * src_image[(row+2)*6144+col + 1]) +
(-1.0f * src_image[(row+2)*6144+col + 2]);//垂直
temp1 = (-1.0f * src_image[row*6144+col]) + (0.0f * src_image[row*6144+col + 1]) +
(1.0f * src_image[row*6144+col + 2])+(-2.0f * src_image[(row+1)*6144+col]) + (0 * src_image[(row+1)*6144+col + 1]) +
(2.0f * src_image[(row+1)*6144+col + 2])+(-1.0f * src_image[(row+2)*6144+col]) + (-0.0f * src_image[(row+2)*6144+col + 1]) +
(1.0f * src_image[(row+2)*6144+col + 2]);
//dst_image[i] = temp;
temp2=sqrt(temp*temp+temp1*temp1);
dst_image[i] = temp2 > 255 ? 255 : temp2;
}
void ApplyFilter2(uint8_t *src_image,
uint8_t *dst_image,
int i) {
int row,col;
row=int(i/6142);
col=i%6142;
float temp,temp1,temp2;
temp = (1.0f * src_image[row*6144+col]) + (2.0f * src_image[row*6144+col + 1]) +
(1.0f * src_image[row*6144+col + 2])+(0 * src_image[(row+1)*6144+col]) + (0 * src_image[(row+1)*6144+col + 1]) +
(0 * src_image[(row+1)*6144+col + 2])+(-1.0f * src_image[(row+2)*6144+col]) + (-2.0f * src_image[(row+2)*6144+col + 1]) +
(-1.0f * src_image[(row+2)*6144+col + 2]);//垂直
temp1 = (-1.0f * src_image[row*6144+col]) + (0.0f * src_image[row*6144+col + 1]) +
(1.0f * src_image[row*6144+col + 2])+(-2.0f * src_image[(row+1)*6144+col]) + (0 * src_image[(row+1)*6144+col + 1]) +
(2.0f * src_image[(row+1)*6144+col + 2])+(-1.0f * src_image[(row+2)*6144+col]) + (-0.0f * src_image[(row+2)*6144+col + 1]) +
(1.0f * src_image[(row+2)*6144+col + 2]);
//dst_image[i] = temp;
temp2=sqrt(temp*temp+temp1*temp1);
dst_image[i] = temp2 > 255 ? 255 : temp2;
}
// This is alternative (to a lambda) representation of a SYCL kernel.
// Internally, compiler transforms lambdas into instances of a very simlar
// class. With functors, capturing kernel parameters is done manually via the
// constructor, unlike automatic capturing with lambdas.
class SepiaFunctor {
public:
// Constructor captures needed data into fields
SepiaFunctor(
accessor<uint8_t, 1, sycl_read, sycl_device> &image_acc_,
accessor<uint8_t, 1, sycl_write, sycl_device> &image_exp_acc_)
: image_acc(image_acc_), image_exp_acc(image_exp_acc_) {}
// The '()' operator is the actual kernel
void operator()(id<1> i) const {
ApplyFilter(image_acc.get_pointer(), image_exp_acc.get_pointer(), i.get(0));
}
private:
// Captured values:
accessor<uint8_t, 1, sycl_read, sycl_device> image_acc;
accessor<uint8_t, 1, sycl_write, sycl_device> image_exp_acc;
};
int main(int argc, char **argv) {
clock_t allTime=0;
double sum_e1=0;
double sum_e2=0;
ifstream readFile;
readFile.open("tu6144.txt", ios::in);
string str;
while (getline(readFile,str)){
char lastChar = str[str.length() - 1];
if (lastChar != 'g') {
cout<<"已去除"<<str<<std::endl;
str.pop_back(); // 去除最后一个字符
}
cout<<str<<std::endl;
cout<<str.c_str()<<std::endl;
// loading the input image
int img_width, img_height, channels;
uint8_t *image = stbi_load(str.c_str(), &img_width, &img_height, &channels, 0);
//uint8_t *image = stbi_load("hnew.png", &img_width, &img_height, &channels, 0);
if (image == NULL) {
cout << "Error in loading the image\n";
exit(1);
}
cout << "Loaded image with a width of " << img_width << ", a height of "
<< img_height << " and " << channels << " channels\n";
img_height=img_height-2;
img_width=img_width-2;
size_t num_pixels = (img_width) * (img_height);
size_t img_size = (img_width) * (img_height) * channels;
// allocating memory for output images
uint8_t *image_ref = new uint8_t[img_size];
uint8_t *image_exp1 = new uint8_t[img_size];
uint8_t *image_exp2 = new uint8_t[img_size];
cout << img_size << " channels\n";
memset(image_ref, 0, img_size * sizeof(uint8_t));
memset(image_exp1, 0, img_size * sizeof(uint8_t));
memset(image_exp2, 0, img_size * sizeof(uint8_t));
// Create a device selector which rates available devices in the preferred
// order for the runtime to select the highest rated device
// Note: This is only to illustrate the usage of a custom device selector.
// default_selector can be used if no customization is required.
MyDeviceSelector sel;
// Using these events to time command group execution
event e1, e2;
// Wrap main SYCL API calls into a try/catch to diagnose potential errors
try {
// Create a command queue using the device selector and request profiling
auto prop_list = property_list{property::queue::enable_profiling()};
queue q(sel, dpc_common::exception_handler, prop_list);
// See what device was actually selected for this queue.
cout << "Running on " << q.get_device().get_info<info::device::name>()
<< "\n";
// Create SYCL buffer representing source data .
// By default, this buffers will be created with global_buffer access
// target, which means the buffer "projection" to the device (actual
// device memory chunk allocated or mapped on the device to reflect
// buffer's data) will belong to the SYCL global address space - this
// is what host data usually maps to. Other address spaces are:
// private, local and constant.
// Notes:
// - access type (read/write) is not specified when creating a buffer -
// this is done when actual accessor is created
// - there can be multiple accessors to the same buffer in multiple command
// groups
// - 'image' pointer was passed to the constructor, so this host memory
// will be used for "host projection", no allocation will happen on host
buffer image_buf(image, range(img_size));
// This is the output buffer device writes to
buffer image_buf_exp1(image_exp1, range(img_size));
cout << "Submitting lambda kernel...\n";
// Submit a command group for execution. Returns immediately, not waiting
// for command group completion.
e1 = q.submit([&](auto &h) {
// This lambda defines a "command group" - a set of commands for the
// device sharing some state and executed in-order - i.e. creation of
// accessors may lead to on-device memory allocation, only after that
// the kernel will be enqueued.
// A command group can contain at most one parallel_for, single_task or
// parallel_for_workgroup construct.
accessor image_acc(image_buf, h, read_only);
accessor image_exp_acc(image_buf_exp1, h, write_only);
// This is the simplest form sycl::handler::parallel_for -
// - it specifies "flat" 1D ND range(num_pixels), runtime will select
// local size
// - kernel lambda accepts single sycl::id argument, which has very
// limited API; see the spec for more complex forms
// the lambda parameter of the parallel_for is the kernel, which
// actually executes on device
h.parallel_for(range<1>(num_pixels), [=](auto i) {
ApplyFilter(image_acc.get_pointer(), image_exp_acc.get_pointer(), i);
});
});
q.wait_and_throw();
cout << "Submitting functor kernel...\n";
buffer image_buf_exp2(image_exp2, range(img_size));
// Submit another command group. This time kernel is represented as a
// functor object.
e2 = q.submit([&](auto &h) {
accessor image_acc(image_buf, h, read_only);
accessor image_exp_acc(image_buf_exp2, h, write_only);
SepiaFunctor kernel(image_acc, image_exp_acc);
h.parallel_for(range<1>(num_pixels), kernel);
});
cout << "Waiting for execution to complete...\n";
q.wait_and_throw();
} catch (sycl::exception e) {
// This catches only synchronous exceptions that happened in current thread
// during execution. The asynchronous exceptions caused by execution of the
// command group are caught by the asynchronous exception handler
// registered. Synchronous exceptions are usually those which are thrown
// from the SYCL runtime code, such as on invalid constructor arguments. An
// example of asynchronous exceptions is error occurred during execution of
// a kernel. Make sure sycl::exception is caught, not std::exception.
cout << "SYCL exception caught: " << e.what() << "\n";
return 1;
}
cout << "Execution completed\n";
// report execution times:
ReportTime("Lambda kernel time: ", e1);
ReportTime("Functor kernel time: ", e2);
sum_e1=sum_e1+(e1.get_profiling_info<info::event_profiling::command_end>()- e1.get_profiling_info<info::event_profiling::command_start>())/1e6;
sum_e2=sum_e2+(e2.get_profiling_info<info::event_profiling::command_end>()- e2.get_profiling_info<info::event_profiling::command_start>())/1e6;
clock_t startTime,endTime;
startTime = clock();//计时开始
// get reference result
for (size_t i = 0; i < num_pixels; i++) {
ApplyFilter2(image, image_ref, i);
}
endTime = clock();//计时结束
string str1="save6144dan";
string str2="save6144lambda";
string str3="save6144functor";
string str13=str1+str.erase(0,9);
cout<<"字符"<<str13<< std::endl;
stbi_write_png(str13.c_str(), img_width, img_height, channels, image_ref,
img_width * channels);
stbi_write_png((str2+str).c_str(), img_width, img_height, channels,
image_exp1, img_width * channels);
stbi_write_png((str3+str).c_str(), img_width, img_height, channels,
image_exp2, img_width * channels);
stbi_image_free(image);
delete[] image_ref;
delete[] image_exp1;
delete[] image_exp2;
cout << "The run time is: " <<(double)(endTime - startTime) / CLOCKS_PER_SEC << "s" << std::endl;
allTime=allTime+(endTime - startTime);
}
cout << "The a singlle kernel run time is: " <<(double)allTime / CLOCKS_PER_SEC << "s" << std::endl;
cout << sum_e1 << " milliseconds\n";
cout << sum_e2 << " milliseconds\n";
return 0;
}