目录
- 内核程序修饰符
- 函数修饰符
- 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);
}
