▶ 上一篇图像卷积 http://www.cnblogs.com/cuancuancuanhao/p/8535569.html。这篇使用了 OpenCV 从文件读取彩色的 jpeg 图像，进行边缘检测以后写回文件。
● 代码（没有使用局部内存优化）
// convolution.cl，核函数，应该和上一篇中无优化内核是一样的
__kernel void convolution(__read_only image2d_t sourceImage, __write_only image2d_t outputImage,
int rows, int cols, __constant float* filter, int filterWidth, sampler_t sampler)
{
const int col = get_global_id(), row = get_global_id(); // 注意工作项的顺序，图像上是先横着数再竖着数
const int halfWidth = (int)(filterWidth / );
float4 sum = { 0.0f, 0.0f, 0.0f, 0.0f }, pixel; // 输出数据类型是四元浮点数，与 image 统一
int i, j, filterIdx; // 传入的卷积窗口是一维的，用一个下标即可遍历
int2 coords;
for (filterIdx = , i = -halfWidth; i <= halfWidth; i++)
{
coords.y = row + i; // 从 work-item 分到的行号偏移 i 行，作为图像坐标的第二分量
for (j = -halfWidth; j <= halfWidth; j++)
{
coords.x = col + j; // 从 work-item 分到的列号偏移 i 列，作为图像坐标的第一分量
pixel = read_imagef(sourceImage, sampler, coords); // 读取源图像上相应位置的值
sum.x += pixel.x * filter[filterIdx++];
}
}
if (row < rows && col < cols) // 将落在有效范围内的计算数据输出
{
coords.x = col;
coords.y = row;
write_imagef(outputImage, coords, sum);
}
return;
}
// convolution.c，主函数
#include <stdio.h>
#include <stdlib.h>
#include <cl.h>
#include <opencv.hpp>
#include <D:\Program\OpenCV\build\include\opencv2\core\cvstd.hpp>// 不知道为什么要加这个，否则报错
#include <iostream>
#include <vector> using namespace std;
using namespace cv; #define CUAN_OPENCL12
#ifndef CUAN_OPENCL12
#pragma warning(disable : 4996)// 部分函数在 OpenCL2.0 中已经被弃用，需要添加该行以支持老版本中的函数，否则报错，信息举例 "'clCreateImage2D': 被声明为已否决 "
#endif const char *sourceProgram = "D:\\Code\\OpenCL\\convolution.cl";// 核函数文件 char* readSource(const char* kernelPath)// 读取文本文件，存储为 char *
{
FILE *fp;
char *source;
long int size;
printf("readSource, Program file: %s\n", kernelPath);
fopen_s(&fp, kernelPath, "rb");
if (!fp)
{
printf("Open kernel file failed\n");
exit(-);
}
if (fseek(fp, , SEEK_END) != )
{
printf("Seek end of file faildd\n");
exit(-);
}
if ((size = ftell(fp)) < )
{
printf("Get file position failed\n");
exit(-);
}
rewind(fp);
if ((source = (char *)malloc(size + )) == NULL)
{
printf("Allocate space failed\n");
exit(-);
}
fread(source, , size, fp);
fclose(fp);
source[size] = '\0';
return source;
} int main()
{
// 卷积窗口相关
const int filterWidth = , filterSize = filterWidth * filterWidth, halfFilterWidth = filterWidth / ;
float filter[filterSize] =
/*
{// 恒等映射
0, 0, 0, 0, 0,
0, 0, 0, 0, 0,
0, 0, 1, 0, 0,
0, 0, 0, 0, 0,
0, 0, 0, 0, 0
};
*/
{// 边缘检测
-, ,-, , ,
,-, , , ,
-, , , ,-,
, , ,-, ,
, ,-, ,-,
}; // 图片相关
Mat image = imread("D:\\1.jpg");
Mat channel[];
split(image, channel);
int imageHeight = image.rows, imageWidth = image.cols;
float *imageData = (float*)malloc(sizeof(float) * imageHeight * imageWidth); // 准备平台，设备，上下文，命令队列部分
cl_int status;
cl_platform_id platform;
clGetPlatformIDs(, &platform, NULL);
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, , &device, NULL);
cl_context_properties props[] = { CL_CONTEXT_PLATFORM,(cl_context_properties)(platform), };
cl_context context;
context = clCreateContext(props, , &device, NULL, NULL, &status);
cl_command_queue queue;
queue = clCreateCommandQueue(context, device, , &status); // 设置 image 数据描述符，仅使用 OpenCL1.2 规范
cl_image_desc desc;
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
desc.image_width = imageWidth;
desc.image_height = imageHeight;
desc.image_depth = ;
desc.image_array_size = ;
desc.image_row_pitch = ;
desc.image_slice_pitch = ;
desc.num_mip_levels = ;
desc.num_samples = ;
desc.buffer = NULL;
cl_image_format format;
format.image_channel_order = CL_R;
format.image_channel_data_type = CL_FLOAT;
cl_mem d_inputImage = clCreateImage(context, CL_MEM_READ_ONLY, &format, &desc, NULL, &status);
cl_mem d_outputImage = clCreateImage(context, CL_MEM_WRITE_ONLY, &format, &desc, NULL, &status); // 卷积窗口缓冲区
cl_mem d_filter = clCreateBuffer(context, , filterSize * sizeof(float), NULL, &status); // 主机数据写入设备
size_t origin[] = { , , }, region[] = { imageWidth, imageHeight, };// 偏移量和每个维度上的尺寸
clEnqueueWriteBuffer(queue, d_filter, CL_TRUE, , filterSize * sizeof(float), filter, , NULL, NULL); // 创建采样器，规定图像坐标系的类型和访问越界时的解决方案，以及插值方式
cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &status); // 程序的运行时编译，创建内核
const char* source = readSource(sourceProgram);
cl_program program = clCreateProgramWithSource(context, , &source, NULL, NULL);
clBuildProgram(program, , &device, NULL, NULL, NULL);
cl_kernel kernel = clCreateKernel(program, "convolution", &status); // 声明内核参数
clSetKernelArg(kernel, , sizeof(cl_mem), &d_inputImage);
clSetKernelArg(kernel, , sizeof(cl_mem), &d_outputImage);
clSetKernelArg(kernel, , sizeof(int), &imageHeight);
clSetKernelArg(kernel, , sizeof(int), &imageWidth);
clSetKernelArg(kernel, , sizeof(cl_mem), &d_filter);
clSetKernelArg(kernel, , sizeof(int), &filterWidth);
clSetKernelArg(kernel, , sizeof(cl_sampler), &sampler); // 内核参数
size_t globalSize[] = { imageWidth, imageHeight }; int i, j;
for (i = ; i < ; i++)// 三个通道，分别为蓝、绿、红
{
// 更新输入缓冲区
for (j = ; j < imageHeight * imageWidth; j++)
imageData[j] = (float)channel[i].data[j];
clEnqueueWriteImage(queue, d_inputImage, CL_TRUE, origin, region, , , imageData, , NULL, NULL); // 执行内核
clEnqueueNDRangeKernel(queue, kernel, , NULL, globalSize, NULL, , NULL, NULL); // 向文件中写入结果
clEnqueueReadImage(queue, d_outputImage, CL_TRUE, origin, region, , , imageData, , NULL, NULL);
for (j = ; j < imageHeight * imageWidth; j++)
channel[i].data[j] = (imageData[j] < ? : (unsigned char)int(imageData[j]));
} merge(channel, , image); // 三个通道合成
imwrite("D:\\2.jpg", image, vector<int>{IMWRITE_JPEG_QUALITY, });// 最后一个参数为输出图片的选项，95%质量
imshow("merge", image); // 在窗口中展示图片
waitKey(); // 等待键盘输入 free(imageData);
getchar();
return ; }
● 输出结果，感谢助教提供的图 
，原图大小 1440 * 900
readSource, Program file: D:\Code\OpenCL\convolution.cl
● 用到 OpenCV 的关键部分
#include <opencv.hpp>
#include <D:\Program\OpenCV\build\include\opencv2\core\cvstd.hpp>// 不知道为什么要加这个，否则报错
#include <iostream>
#include <vector> {
Mat image = imread("D:\\1.jpg"); // 读取图片，OpenCV 自动识别文件类型，返回一个 Mat 类
Mat channel[]; // 分别存放图像的三个通道
split(image, channel); // 将原图像拆分为三个通道，分别为蓝色、绿色、红色
int imageHeight = image.rows, imageWidth = image.cols;// 获取图像的行数和列数 float value = (float)channel[].data[]; // 获取图像中某一像素的值，注意格式为 unsigned char，注意与 float 之间的转换
channel[].data[] = (unsigned char)int(value); // float 转回 unsigned char，注意数据范围和四舍五入方向 merge(channel, , image); // 三个通道合成
imwrite("D:\\2.jpg", image, vector<int>{IMWRITE_JPEG_QUALITY, });// 最后一个参数为输出图片的选项，95% 压缩质量
imshow("merge", image); // 在窗口中展示图片
waitKey(); // 等待键盘输入
}
● 一个玄学错误：在使用图像写入函数 imwrite() 时报错
0x00007FFBB8A086C2 (opencv_world320.dll)处(位于 OpenCLProjectTemp.exe 中)引发的异常: 0xC0000005: 读取位置 0x000001DB4743F000 时发生访问冲突。
找了很多方法，大致有：
① 改文件途径，把路径从 "XX.jpg" 变成 "./XX.jpg" 或 ".\\XX.jpg"。无效，我用的绝对路径仍然会报错。
② OpenCV 不支持中文路径。我把输入和输出文件放到 "D:\\中文" 目录下，程序仍然能够正确执行。
③ 使用 release 模式而不是 debug 模式，仍然会出错。
④ 在调用函数 imwrite() 时显式规定第三个参数。这个我信了Orz，虽然该函数的原型中该参数有默认值 vector<int>{0}，但是不显式规定第三个参数的情况下一次都没有通过。可能还与我用的保存格式（jpg）有关，保存为 jpg、png 的时候可能必须规定该参数，而保存为 bmp 时没有改参数还是通过了。
⑤ 我之前觉得这个是内存泄漏了，因为该问题没有可重现性，直接用 .exe 执行程序会有以一定概率出错（注意动态库 opencv_world320.dll 的位置，拿不准环境变量的就放一份到 .exe 目录下）；在 VS 里 debug 时完全不改的代码仍有一定概率报错。
⑥ 关于 lib 的配置，这个解决了我的问题（不知道其他人是不是这样）。一开始 配置属性 -> 链接器-> 输入 -> 附加依赖项 的时候我用的是 opencv_world320.lib 而不是 opencv_world320d.lib，且运行程序用的是 debug 模式，所以会报错。其实这里 d 是 debug 的意思，opencv_world320.lib 是用于 release 模式的，opencv_world320d.lib 才是用于 debug 模式的，所以 ③ 说的是对的，应该使用与库相应的模式，不再有报错。