opencl:C++实现双线性插值图像缩放

用OpenCL实现图像缩放代码是比较简单的，因为OpenCL本身就支持双线性插值
下面是kernel代码(从Mali OpenCL SDK 抄来的：/samples/image_scaling/assets/image_scaling.cl)
非常简单只有4行

// 定义采样器 
// CLK_NORMALIZED_COORDS_TRUE指定使用归一化坐标
// CLK_ADDRESS_CLAMP 指定超出图像范围的颜色为黑色
// CLK_FILTER_LINEAR指定使用双线性插值
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR;
__kernel void image_scaling(__read_only image2d_t sourceImage,
                            __write_only image2d_t destinationImage,
                            const float widthNormalizationFactor,
                            const float heightNormalizationFactor)
{
//从glob_id中获取目标像素坐标
    int2 coordinate = (int2)(get_global_id(0), get_global_id(1));
//计算归一化浮点坐标    
    float2 normalizedCoordinate = convert_float2(coordinate) * (float2)(widthNormalizationFactor, heightNormalizationFactor);
//根据归一化坐标从原图中读取像素数据
    float4 colour = read_imagef(sourceImage, sampler, normalizedCoordinate);
//将像素数据写入目标图像    
    write_imagef(destinationImage, coordinate, colour);
}

下面是主机代码片段

/* 向OpenCL设备中上传原始图像数据 */
void gray_matrix_cl::upload(const cl::CommandQueue& command_queue) const {
	cl::size_t<3> region;
	region[0] = width, region[1] = height, region[2] = 1;
	command_queue.enqueueWriteImage(cl_img,
	CL_TRUE, ZeroOffset, region, get_row_stride(), 0, (void*) pixels.data());
}
/* 从OpenCL设备中下载结果数据 */
void gray_matrix_cl::download(const cl::CommandQueue& command_queue) {
	cl::size_t<3> region;
	region[0] = width, region[1] = height, region[2] = 1;
	// 分配目标图像空间
	if (pixels.size() != get_row_stride() * height)
		pixels = std::vector<uint8_t>(get_row_stride() * height);
	command_queue.enqueueReadImage(cl_img,
	CL_TRUE, ZeroOffset, region, get_row_stride(), 0, (void*) pixels.data());
}
/* 缩放图像(双线性插值) 返回缩放后的图像数据*/
// facecl_context参见上一篇博客(http://blog.csdn.net/10km/article/details/50755251)
gray_matrix_cl gray_matrix_cl::zoom(size_t dst_width, size_t dst_height, const facecl_context& context)const {
	gray_matrix_cl dst_matrix(dst_width, dst_height);
	auto kernel = context.getKernel(KERNEL_NAME(image_scaling));// 获取已经编译好的cl::Kernel
	// 参见上一篇博客(http://blog.csdn.net/10km/article/details/50755251)
	auto command_queue = global_facecl_context.getCommandQueue();// 获取cl::CommandQueue
	this->upload(command_queue);//向OpenCL设备中上传原始图像数据
	cl_float widthNormalizationFactor = 1.0f / dst_width;
	cl_float heightNormalizationFactor = 1.0f / dst_height;
	// 设置kernel参数
	kernel.setArg(0, cl_img);
	kernel.setArg(1, dst_matrix.cl_img);
	kernel.setArg(2, widthNormalizationFactor);
	kernel.setArg(3, heightNormalizationFactor);
	const cl::NDRange global(dst_width, dst_height);
	// 执行 kernel
	command_queue.enqueueNDRangeKernel(kernel, gray_matrix_cl::NullRange, global);
	command_queue.finish();// 等待kernel执行结束
	dst_matrix.download(command_queue);// 从OpenCL设备中下载结果数据
	return std::move(dst_matrix);
}