学习OpenCL与卷积神经网络 Day1——对OpenCL与CNN的简单认识

自学工具

自学用电子书

OpenCL异构计算
OpenCL编程指南
OpenCL异构并行计算原理、机制与优化

自学用课程

链接: link.
链接: link.

自学用平台

Vs2017
Matlab
Quartus Prime 16

关于OpenCL

平台模型

OpenCL平台模型包含一个主机或多个OpenCL设备，每个OpenCL设备包含一个或多个计算单元，每个计算单元至少包含一个或者多个处理单元。

执行模型

OpenCL程序包含主机端程序和设备端内核(Kernel)程序。其中主机端程序以命令方式将内核程序从主机提交到OpenCL设备上，OpenCL设备在处理单元上执行计算。

内核

内核(Kernel)程序一般都是计算量大、逻辑比较简单的函数。
当主机发出一个命令，提交一个内核到OpenCL设备上执行，OpenCL运行时将会创建一个整数索引空间(N维)，称为NDRange。其中N为1或2或3。三个长度为N的数据确定了NDRange的以下特征:
1.每个维度索引的范围。
2.一个偏移指数F表明每个维度的初始索引值(默认为0)。
3.一个工作组的维度大小。
work-item

上下文

主机使用OpenCL API来创建和管理上下文，内核在此上下文中执行。
上下文定义内核执行环境:
设备:OpenCL平台包含一个或多个设备
内核对象:在OpenCL设备上运行的OpenCL内核函数
程序对象:实现整个内核程序的源代码和目标二进制码
存储对象:对主机和OpenCL设备可见的对象，内核执行时操作这些对象的实例

命令队列

通过命令队列描述主机与OpenCL设备进行交互。一般由主机或是运行在设备中的内核提交给命令队列。

存储器模型

存储器区域包含主机与设备的内存。

存储器类型

主机内存(host memory)
全局存储器(global memory)
常量存储器(constant memory)
局部存储器(local memory)
私有存储器(private memory)

存储器对象

缓冲(buffer):可以将内检数据类型、矢量类型数据或用户自定义数据结构映射到缓冲区，内核通过指针来访问缓冲区。
图像(image):使用OpenCL API函数管理。通常不允许OpenCL内核对单个图像同时进行读和写。
管道(pipe):是数据项有序的队列。其有两个端点，一个是读端点，一个是写端点。

关于卷积神经网络(CNN)

目的与用途

主要用于计算机视觉，即图像识别一类。比如果园中自动采摘西红柿机器人，图像识别起到判断该物体"是否是西红柿"的功能。

构造

卷积层(Convolution)

卷积层的作用就是提取图片中的有用信息(特征)
假定我们有一个尺寸为5x5的图像，每一个像素点里都存储着图像的信息。

我们再定义一个3x3的卷积核[滤波器]（filter），用来从图像中提取一定的特征。

我们可以通过其与不同的卷积核来进行测试，得出不同的输出结果。而卷积层输出值越大，则与原图的匹配度就越高，识别就越精确。

但如上图我们边缘信息可能会丢失，所以我们可以采取一个简单的办法，保留其信息。

就比如在周围加上一圈"信息保护层"。

池化层(Pooling)

在一张图片中，相邻像素点的阈值往往极为相似，对它们做了卷积之后，卷积层(Convolution)相邻的值会有很大的一个冗余，此时我们便对其进行一次池化(Pooling)，作用是对卷积后的值进行一次(猜测)处理。一般会有max、min、average三种操作。
以下是一个平均池的例子:

但同样的，虽然池化层不需要训练，但会损失不少的特征值，这时候我们还需要进行误差反传。

全连接层(FC layer)

经过多轮卷积层与池化层的处理过后，一般会由1到2个全连接层来给出最后的分类结果。

关于将OpenCL与CNN相结合的想法

通过异构并行体系，将一些简单而又具有庞大计算量的工作交给OpenCL设备，例如FPGA来完成，这样可以大大缩短CPU的工作负荷与运作速率。

简单程序构建

OpenCL与简单卷积的结合：
main.cpp

#include<stdio.h>
#include<stdlib.h>
#include<CL/cl.h>
#include<string.h>
#include <iostream>  
#include <ctime>

#define MAX_SOURCE_SIZE (0x10000)
#pragma warning(disable : 4996)
int main(void)
{
	srand(time(NULL));
	cl_platform_id	platform_id = NULL;
	cl_uint			ret_num_platforms;
	cl_device_id	device_id = NULL;
	cl_uint			ret_num_devices;
	cl_context		context = NULL;
	cl_command_queue command_queue = NULL;
	cl_mem			data_in = NULL;
	cl_mem			data_out = NULL;
	cl_mem			filter_in = NULL;
	cl_program		program = NULL;
	cl_kernel		kernel = NULL;
	size_t			kernel_code_size;
	char			*kernel_str;
	int				*result;
	cl_int			ret;
	FILE			*fp;
	cl_uint			work_dim;
	size_t			global_item_size[2];
	size_t			local_item_size[2];
	

	int const W = 8;			//image width
	int const H = 8;			//image height
	int const K = 3;			//filter kernel size
	int const Wn = (W + K - 1); //padded image width
	int const Hn = (H + K - 1); //padded image height

	int point_num = Wn * Hn;
	int data_vecs[Wn*Hn];
	int filter_coe[K*K] = { -1,0,1,-2,0,2,-1,0,1 }; //sobel filter: horizontal gradient
	int i, j;

	for (i = 0; i < point_num; i++)
	{
		data_vecs[i] = rand() % 20;
	}

	//display input data
	printf("\n");
	printf("Array data_in:\n");
	for (i = 0; i < Hn; i++) {
		printf("row[%d]:\t", i);
		for (j = 0; j < Wn; j++) {
			printf("%d,\t", data_vecs[i*Wn + j]);
		}
		printf("\n");
	}
	printf("\n");



	kernel_str = (char *)malloc(MAX_SOURCE_SIZE);
	result = (int *)malloc(W*H * sizeof(int));

	ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
	ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id,
		&ret_num_devices);
	context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
	command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
	fp = fopen("Conv2D.cl", "r");
	kernel_code_size = fread(kernel_str, 1, MAX_SOURCE_SIZE, fp);
	fclose(fp);
	program = clCreateProgramWithSource(context, 1, (const char **)&kernel_str,
		(const size_t *)&kernel_code_size, &ret);
	ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
	kernel = clCreateKernel(program, "Conv2D", &ret);
	data_in = clCreateBuffer(context, CL_MEM_READ_WRITE, Wn*Hn * sizeof(int), NULL, &ret);
	data_out = clCreateBuffer(context, CL_MEM_READ_WRITE, W*H * sizeof(int), NULL, &ret);
	filter_in = clCreateBuffer(context, CL_MEM_READ_WRITE, K*K * sizeof(int), NULL, &ret);

	//write image data into data_in buffer
	ret = clEnqueueWriteBuffer(command_queue, data_in, CL_TRUE, 0, Wn*Hn * sizeof(int), data_vecs, 0, NULL, NULL);

	//write filter data into filter_in buffer
	ret = clEnqueueWriteBuffer(command_queue, filter_in, CL_TRUE, 0, K*K * sizeof(int), filter_coe, 0, NULL, NULL);

	//set kernel arguments
	ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&data_in);
	ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&filter_in);
	ret = clSetKernelArg(kernel, 2, sizeof(int), (void *)&K);
	ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *)&data_out);

	work_dim = 2;
	global_item_size[0] = { W };
	global_item_size[1] = { H };
	local_item_size[0] = { 1 };
	local_item_size[1] = { 1 };

	//execute data parallel kernel */
	ret = clEnqueueNDRangeKernel(command_queue, kernel, work_dim, NULL,
		global_item_size, local_item_size, 0, NULL, NULL);


	// read data_out to host
	ret = clEnqueueReadBuffer(command_queue, data_out, CL_TRUE, 0,
		W*H * sizeof(int), result, 0, NULL, NULL);

	//display output data
	FILE *f_img_out = fopen("image_out.txt", "w+");
	printf("Array data_out: \n");
	for (i = 0; i < H; i++) {
		printf("row[%d]:\t", i);
		for (j = 0; j < W; j++) {
			printf("%d,\t", result[i*W + j]);
			fprintf(f_img_out, "%d,\t", result[i*W + j]);
		}
		printf("\n");
		fprintf(f_img_out, "\n");
	}
	printf("\n");
	fclose(f_img_out);
	

	ret = clReleaseKernel(kernel);
	ret = clReleaseProgram(program);
	ret = clReleaseMemObject(data_in);
	ret = clReleaseMemObject(data_out);
	ret = clReleaseMemObject(filter_in);
	ret = clReleaseCommandQueue(command_queue);
	ret = clReleaseContext(context);
	free(result);
	free(kernel_str);

	system("pause");
	return 0;
}

cl文件

__kernel void Conv2D(__global int * image_in,  //image input
	__global int * filter_in, //filter input
	int K,            //filter kernel size
	__global int * image_out) //feature map output
{

	int W;       //work group global size
	int Wn;      //padded image width
	int x;       //global id x
	int y;       //global id y
	int ki, kj;  //filter coordinate,(kj, ki)

	int sum = 0; //multiply and sum of filter and data
	W = get_global_size(0);
	x = get_global_id(0);
	y = get_global_id(1);
	Wn = W + (K - 1);

	for (ki = 0; ki < K; ki++)
		for (kj = 0; kj < K; kj++)
		{
			sum = sum + filter_in[ki*K + kj] * image_in[Wn*(y + ki) + x + kj];
		}

	image_out[y*W + x] = sum;
}

如上代码参考自github 链接: link.

运行结果

通过strand函数将输入矩阵设为了随机值。

卷积的方法有很多 明天尝试下