自學工具

自學用電子書

OpenCL異構計算
OpenCL程式設計指南
OpenCL異構平行計算原理、機制與優化

自學用課程

[link]https://www.coursera.org/learn/opencl-fpga-introduction/lecture/6HxML/introduction-and-parallel-computing
[link]https://www.youtube.com/watch?v=YU_pRT-Be0c&list=PLzy5q1NUJKCJocUKsRxZ0IPz29p38xeM-

自學用平臺

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的工作負荷與運作速率。

簡單程式構建

main.cpp

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


#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函式將輸出矩陣設為了隨機值。

但這只是在圖片上加了一層卷積（濾波器），並沒有進行Pooling
明天想辦法加一個4x4的Pooling加到cl檔案中去