目录
自学工具
自学用电子书
OpenCL异构计算
OpenCL编程指南
OpenCL异构并行计算原理、机制与优化
自学用课程
自学用平台
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函数将输入矩阵设为了随机值。
卷积的方法有很多 明天尝试下