学习OpenCL与卷积神经网络 Day4——OpenCL内核(Kernel)程序构建

内核程序修饰符

函数修饰符

函数修饰符用来修饰OpenCL内核函数及一般函数的一些特性，以帮助编译器确定一些信息。

kernel修饰符

__kernel修饰符声明一个函数为一个内核函数，这个内核函数将会在OpenCL设备上执行，这也是我们最常用的修饰符。内核函数返回类型必须是void类型，且主机端可以调用这个函数。

__kernel void trial(...)
__kernel int trial(...)//(x)kernel函数的返回值必须是空

地址空间修饰符

__global int * filter_in,
__local //
__constant//
__private//

以上四种空间修饰符分别对应全局存储器、局部存储器、常量存储器与私人存储器。若一个变量由一个地址空间修饰符限定，那么这个变量就在指定的地址空间中分配。
程序中的函数参数或者函数中缺略地址修饰符的局部变量，它们的地址修饰符为private（所有函数参数一定在私有地址空间）；在程序范围内的一个变量，或者程序内的一个static变量，它们在全局或者常量地址空间；如果没有地址修饰符指定，默认为全局。
内核参数声明的指针类型必须指向global、local、constant三种类型之一。
为了进一步了解，在这里举个例子：
在私有地址空间中声明一个指针p,指针指向的是一个全局地址空间对象；

__global int *p;

在私有地址空间中声明一个整数数组；

float a[4];

在函数参数中：

int code_b(constant float*scr, int4 v) {
	//code_b小括号中的都为函数参数故都在私有地址空间
	//scr:在私有空间中声明了一个指向常量地址空间的指针
	float temp;
	//temp分配在私有空间
}

全局地址空间

可以用来指示缓冲区或者图像。一个缓冲区对象在内核参数中声明为指针，可以指向标量、矢量和自定义结构体。内核可以读写缓冲区的任何位置。存储器大小由主机端调用API分配时所决定。
一般在全局地址空间分配图像对象，但全局地址修饰符不能用于图像类型，对图像对象不能直接读写，OpenCL提供了内建函数来支持对图像对象的读写。
一般使用情况：

__kernel void trial(__global int * image_in,  //image input
	__global int * filter_in, //filter input
	__global int * image_out) //feature map output

局部地址空间

局部地址空间用来描述需要在局部存储器中分配的变量，这些变量被一个工作组中的所有工作项所共享。局部地址空间的指针可以作为函数的参数和函数内声明的变量。内核函数中声明的变量可以在局部地址空间中分配，但是有一些限制：这些变量声明必须出现在内核函数作用域；这些变量不能被初始化。
一般使用情况：

__kernel void trial(__global....)
{
__local float a;
__local float b[10];
__local float c=1;//(X)变量c不能会在声明时初始化
//__local float c;
//c=1;
if(...){
__local float d;//(X)变量d的作用域不在内核函数作用域
  }
}

常量地址空间

常量地址空间用来描述全局存储器中分配的一些变量。这些变量在内核中作为只读变量访问。这些只读变量在内核执行时允许所有工作项访问。常量地址空间的指针可以作为函数参数和函数内声明的变量。
一般使用情况：

constant float A[] = {0,1,2,3};
kernel void trial(constant float*A,constant float*B)
{
	constant float4 *p = A;
	constant float a = 1.0f;
	constant float b;//(x)未初始化
	char *c = "OpenCL";
}

私有地址空间

内核函数中未使用地址空间修饰符的变量、非内核函数中声明的所有变量，及其函数参数，都在私有地址空间中。
一般使用情况：

int code_b(constant float*scr, int4 v) {
	float temp;
	//temp分配在私有空间
}

对象访问修饰符

访问修饰符可以用于指定图像类型的参数。内核参数中的图像对象可以声明为只读、只写或者读写。如果内核读或写图像对象，那么使用__read_only和__write_only来修饰图像对象参数（这两者只能在内核中）
一般使用情况：

__kernel void trial(__read_only_image2d_t imageA,
	__write_only_image2d_t imageB
) {
	...
}

主程序中内核

创建内核对象

创建内核对象一般有两种方式：

clCreateKernel(cl_program program,
               const char *kernel_name,
               cl_int *errcode_ret)
                         

clCreateKernelsInProgram(cl_program program,
                         cl_unit num_kernels,
                         cl_kernel *kernels,
                         cl_unit *num_kernels_ret)

详情: link.
值得注意的一点是：如果程序对象有多个内核时，若采用clCreateKernelsInProgram()函数来创建内核对象，内核列表中内核函数名称顺序并不是根据cl文件中内核书写顺序，而是依赖于现实。我们可以通过传递param_name参数CL_KERNEL_FUNCTION_NAME值来调用clGetKernelInfo函数来查询每个内核对象的函数名。

设置内核参数

为了执行一个具体内核，必须能够向内核函数传递参数。

cl_int clSetKernelArg(cl_kernel      kernel,
      				  cl_uint        arg_index,
            		  size_t         arg_size, 
     				  const void     *arg_value)	 

详情: link.

执行内核

利用命令队列使将在设备上执行的内核排队：

clEnqueueNDRangeKernel(cl_command_queue queue,
					   cl_kernel kernel,
                       cl_uint work_dims,
                       const size_t *global_work_offset,
                       const size_t *global_work_size, 
                       const size_t *local_work_size,
                       cl_uint num_events,
                       const cl_event *wait_list,
                       cl_event *event)

详情: link.

内核程序样例

__constant sampler_t sampler=
	CLK_NORMALIZED_COORDS_FALSE|
	CLK_ADDRESS_CLAMP|
	CLK_FILTER_NEAREST;
	
typedef struct ImageScaleNode{
	float X;
	float Y;
	}ImageScale;
	
__kernel void ADL(__read_only image2d_t original ,__global ImageScale * scale, __write_only image2d_t output){
	int2 outputcoord = (int2){get_global_id(0),get_global_id(1)};
	float2 Qcoord = (float2){outputcoord.x/scale->X,outputcoord.y/scale->Y};
	int srcH= get_image_height(original);
	int srcW= get_image_width(original);
	uint4 P;
	int2 tlp=convert_int2(Qcoord);/*(int2){outputcoord.x/(*scale),outputcoord.y/(*scale)};*/
	if(Qcoord.x<=0||Qcoord.y<=0||Qcoord.x>=srcW||Qcoord.y>=srcH){
		if(tlp.x<0 ){
			tlp.x=0;
		}
		if( tlp.y<0 ){
			tlp.y=0;
		}
		if( tlp.x>=srcW ){
			tlp.x=srcW-1;
		}
		if(tlp.y >=srcH){
			tlp.y=srcH-1;	
		}
		P=read_imageui(original,tlp);
	}
	else{
		
		uint4 tl=read_imageui(original,tlp);
		uint4 tr=read_imageui(original,tlp+(int2)(-1,0));
		
		uint4 bl=read_imageui(original,tlp+(int2)(0,-1));
		uint4 br=read_imageui(original,tlp+(int2)(-1,-1));
		
		float dx=Qcoord.x-tlp.x;
		float dy=Qcoord.y-tlp.y;
		
		float4 Qt=(float4)(dx*convert_float4(tl)+(1-dx)*convert_float4(tr));
		
		float4 Qb=(float4)(dx*convert_float4(bl)+(1-dx)*convert_float4(br));
		
		P=convert_uint4(dy*Qt+(1-dy)*Qb);
	}
	write_imageui(output,outputcoord,P);
	
}