在 OpenCL 1.1 中，已经支持对图像的读写。本文描述如何通过 OpenCL API 调用，从一张较大的图像中截取一块矩形区域，供图片浏览器显示。该程序实现的功能和
No.1_YUV420pScissor
类似，不同的是 No.1_YUV420pScissor 中直接对像素数据进行处理，而该程序中通过 OpenCL 内置图像读写函数来对像素数据进行采样。
该程序接收
No.2_2_FreeImage
中由函数
load_image
保存的二进制图像数据，然后通过 OpenCL 来提取其中的部分矩形区域，将该矩形区域保存为二进制图像文件。最后将提取的矩形区域像素使用 No.2_2_FreeImage 中
store_image
函数保存为位图文件，供图片浏览器查看。具体执行流程可参考 No.2_2_FreeImage。
由于截取的矩形区域图像为 256x256，在保存位图文件时需要更新图像格式和尺寸，需要对程序 No.2_2_FreeImage 中对应参数作相关调整，打上如下 patch，内容如下：
+FREE_IMAGE_FORMAT g_format = (FREE_IMAGE_FORMAT)13;
+int g_width = 256, g_height = 256;
完整代码参见
No.1_OpenCLSampler
。
1.创建图像对象
缓冲区对象和图像对象都是内存对象，内存对象是对全局内存区域的引用。不同的是图像对象除了像素数据，还包含了对图像格式及图像属性的描述信息，例如图像宽度、高度，以及深度等。
创建了两个图像对象，分别用来存放原始图像数据和目标图像数据。图像对象和缓冲区对象类似，都是 cl_mem 类型。
cl_image_format image_format;
cl_mem in_buffer, out_buffer;
memset((void *)&image_format, 0, sizeof(cl_image_format));
image_format.image_channel_order = CL_RGBA;
image_format.image_channel_data_type = CL_UNORM_INT8;
in_buffer = clCreateImage2D(context,
CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, &image_format,
orig_width, orig_height, 0, src, &err);
out_buffer = clCreateImage2D(context, CL_MEM_WRITE_ONLY, &image_format,
new_width, new_height, 0, NULL, &err);
其中 cl_image_format 用来定义图像的格式，成员如下：
-
image_channel_order：指定每个像素的通道数目和通道布局；
-
image_channel_data_type：描述每个通道的大小，表示的数据类型。
该程序在华为
荣耀8
上运行。由于
荣耀8
只支持 OpenCL 1.1，该版本创建图像对象对应的函数是 clCreateImage2D，如本文所示。在 OpenCL 1.2 中将函数 clCreateImage2D 和 clCreateImage3D 合并到了 clCreateImage，同时增加了 cl_image_desc 参数，用来描述图像的属性，包含图像类型、宽度、高度和深度等信息。
2.图像读写
图像读写操作通过 OpenCL 内置的图像读写函数
read_imageX
和
write_imageX
来执行，图像读写函数只能作用于图像对象。在主机端，图像对象使用
clSetKernelArg
传递给内核程序。内核代码如下：
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_FILTER_NEAREST | CLK_ADDRESS_CLAMP;
__kernel void rotate_rgba(__read_only image2d_t srcImg,
__write_only image2d_t dstImg)
const int x = get_global_id(0);
const int y = get_global_id(1);
float4 value = read_imagef(srcImg, sampler, (int2)(x, y));
write_imagef(dstImg, (int2)(x, y), value);
2.1采样器
采样器对象描述了读取图像数据时如何对图像进行采样。图像读取函数 read_imageX 包含一个采样器参数，该参数可以在主机端通过调用 OpenCL API 函数创建，然后使用 clSetKernelArg 传递给内核；也可以在内核程序中声明，在内核程序中声明的采样器对象为 sampler_t 类型的常量。采样器对象包含了一些属性，这些属性描述了在读取图像对象的像素时如何采样。分别是规格化浮点坐标，寻址模式和过滤模式。
- 规格化坐标：指定传递的 x、y 和 z 坐标值是规格化浮点坐标还是非规格化坐标值。可以是
CLK_NORMALIZED_COORDS_TRUE 或者 CLK_NORMALIZED_COORDS_FALSE 枚举类型的值； - 寻址模式：指定图像的寻址模式。即，当传递的坐标值超过图像坐标区域时该如何处理。可以是下面的枚举类型的值：
CLK_ADDRESS_MIRRORED_REPEAT：图像区域外的坐标设置为区域内坐标的反射值对应的颜色；
CLK_ADDRESS_REPEAT：图像区域外的坐标重复区域内坐标的颜色，只对规格化坐标有效；
CLK_ADDRESS_CLAMP_TO_EDGE：图像区域外的坐标返回图像边缘的颜色；
CLK_ADDRESS_CLAMP：图像区域外坐标返回的颜色和边框颜色保持一致； - 过滤模式：指定使用的过滤模式。可以是
CLK_FILTER_NEAREST 或 CLK_FILTER_LINEAR 枚举类型值，分别表示最近邻插值和双线性插值。
边框的颜色
当采样器的寻址模式设置为 CLK_ADDRESS_CLAMP 时，超过图像区域的坐标返回边框的颜色。边框颜色的选择依赖图像颜色通道，如下：
- 如果图像颜色通道是 CL_A，CL_INTENSITY，CL_Rx，CL_RA，CL_RGx，CL_RGBx，CL_ARGB，CL_BGRA 或 CL_RGBA，边框颜色为 （0.0f，0.0f， 0.0f，0.0f）；
- 如果图像颜色的通道是 CL_R，CL_RG，CL_RGB 或 CL_LUMINANCE，边框颜色为（0.0f，0.0f，0.0f，1.0f）。
CL_RGBA 的颜色通道中，如果是全 0 则表示黑色，全 0xFF 表示白色。Alpha 通道如果是 1，则不透明，如果是 0 则全透明。
2.2坐标
在图像读写函数 read_imageX 和 write_imageX 中，对于 2D 图像对象，传递的 x 坐标范围是 [0, width-1]，y 坐标范围是 [0, height-1]，坐标原点 [0, 0] 对应图像的左下角。
2.3读写操作
对于不同的图像对象，函数 read_imageX 返回由四个通道组成的颜色值，使用 x、 y、 z 和 w 来表示。x 表示红色，y 表示绿色，z 表示蓝色，w 表示 alpha 通道。函数 write_imageX 将四个通道表示的颜色值写入图像对象指定的坐标位置。图像数据类型 image2d_t，表示 2D 图像，需要使用如下的限定符来修饰：
- 对于内核执行
读操作的图像对象，在声明的时候需要使用 __read_only 限定符修饰； - 对于内核执行
写操作的图像对象，在声明的时候需要使用 __write_only 限定符修饰； - 同一个图像对象不能同时支持
读写操作，限定符使用不当将引发编译错误。
3.读取图像数据
把命令添加到命令队列中，将图像对象指定区域的数据读取到主机内存，以供图片浏览器查看。该区域是一个矩形区域。
size_t origin[3] = {0 , 0, 0};
size_t region[3] = {new_width, new_height, 1};
err = clEnqueueReadImage(queue, out_buffer, CL_TRUE, origin, region, 0,
0, des, 0, NULL, NULL);
部分参数说明如下：
- origin：图像中的偏移（x，y，z），以像素为单位的。如果
image 是 2D 图像对象，origin[2] 必须为 0； - region：定义 1D，2D 或 3D 图像的矩形区域，属性为（width，height，depth）。如果 image 是 2D 图像对象， region[2] 必须是 1。
该程序运行后，经过 No.2_2_FreeImage 处理显示效果如下图所示。
OpenCL内存对象：
OpenCL内存对象就是一些OpenCL数据，这些数据一般在设备内存中，能够被拷入也能够被拷出。OpenCL内存对象包括buffer对象和image对象。
buffer对象：连续的内存块---
// kernel 代码 theta 为旋转角度 如45
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP;
__kernel
void rotation(__read_only image2d_t inputImage, __write_only i
目录自学工具自学用电子书自学用课程自学用平台关于OpenCL平台模型执行模型内核上下文命令队列存储器模型存储器类型存储器对象关于卷积神经网络(CNN)目的与用途构造卷积层(Convolution)池化层(Pooling)全连接层(FC layer)关于将OpenCL与CNN相结合的想法简单程序构建
自学用电子书
OpenCL异构计算
OpenCL编程指南
OpenCL异构并行计算原理、机制与优化
自学用课程
[link]https://www.coursera.org/learn/opencl-f
内存对象可以用来对主机和设备之间所要传输的数据进行包装。内存对象可以分为两种类型：缓存对象和图像对象。缓存对象用来传输通用数据。理论上讲：可以将图像对象保存在缓存对象中，将它的像素作为一般的缓存数据来访问。但选择图像对象有如下四个重要的理由：
（1）在GPU上，图像数据是保存在特殊的全局内存中，这个内存被称为纹理内存，和一般的全局内存不同，纹理内存是被缓存，用于高速访问处理。
（2）用来读...
参考链接：
https://www.zhihu.com/people/wujianming_110117/posts
先以图像旋转的实例，具体介绍OpenCL编程的步骤。 首先给出实现流程，然后给出实现图像旋转的C循环实现和OpenCL C kernel实现。
图像旋转原理
图像旋转是指把定义的图像绕某一点以逆时针或顺时针方向旋转一定的角度， 通常是指绕图像的中心以逆时针方向旋转。假设图像的左上角为(l, t), 右下角为(r, b)，则图
1.设备上的图像对象：image2d_t和image3d_t
当主机将一个图像对象传递到设备上，设备上的内核函数会将它作为一个内核参数来访问。这个参数的数据类型和图像的维度有关。如果接受的是二维图像对象，参数的数据类型就是image2d_t，如果接受的是三维图像对象，参数的数据类型就是image3d_t.
缓存对象参数可以接受__global或__kernel等修饰符来指定在设备上存储缓存对象的位...
img = cv2.imread('image.jpg', cv2.IMREAD_GRAYSCALE)
# Create OpenCL context and queue
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
# Create OpenCL program and kernel
program = cl.Program(ctx, """
__kernel void median_filter(__global const uchar* src, __global uchar* dst,
const int width, const int height, const int window_size) {
const int x = get_global_id(0);
const int y = get_global_id(1);
const int offset = window_size / 2;
const int index = y * width + x;
if (x >= offset && x < width - offset && y >= offset && y < height - offset) {
// Initialize window
uchar window[25];
for (int i = -offset; i <= offset; i++) {
for (int j = -offset; j <= offset; j++) {
window[(i + offset) * window_size + (j + offset)] = src[(y + i) * width + (x + j)];
// Sort window
for (int i = 0; i < window_size * window_size - 1; i++) {
for (int j = i + 1; j < window_size * window_size; j++) {
if (window[i] > window[j]) {
uchar temp = window[i];
window[i] = window[j];
window[j] = temp;
// Set median pixel value
dst[index] = window[window_size * window_size / 2];
else {
// Copy edge pixels
dst[index] = src[index];
""").build()
kernel = program.median_filter
# Create OpenCL buffers
img_buf = cl.Buffer(ctx, cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR, hostbuf=img)
result_buf = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, img.nbytes)
# Execute kernel
window_size = 5
global_size = (img.shape[1], img.shape[0])
local_size = (16, 16) # Must be a multiple of the work group size defined in the kernel
kernel(queue, global_size, local_size, img_buf, result_buf, np.int32(img.shape[1]), np.int32(img.shape[0]), np.int32(window_size))
# Copy result back to host
result = np.empty_like(img)
cl.enqueue_copy(queue, result, result_buf)
# Display result
cv2.imshow('Original Image', img)
cv2.imshow('Filtered Image', result)
cv2.waitKey(0)
cv2.destroyAllWindows()
这个代码首先加载图像，然后创建OpenCL上下文和命令队列。接下来，它创建OpenCL程序和中值滤波内核，并使用它们创建OpenCL缓冲区以存储输入和输出数据。然后，它执行内核，并将结果从OpenCL缓冲区复制回主机内存。最后，它显示原始图像和过滤后的图像。
Yongqiang Cheng:
