OpenCL开发.docx
《OpenCL开发.docx》由会员分享，可在线阅读，更多相关《OpenCL开发.docx（12页珍藏版）》请在冰豆网上搜索。
OpenCL开发
这是第一篇真正的OpenCL教程。
这篇文章不会从GPU结构的技术概念和性能指标入手。
我们将会从OpenCL的基础API开始，使用一个小的kernel作为例子来讲解基本的计算管理。
首先我们需要明白的是，OpenCL程序是分成两部分的：
一部分是在设备上执行的（对于我们，是GPU），另一部分是在主机上运行的（对于我们，是CPU）。
在设备上执行的程序或许是你比较关注的。
它是OpenCL产生神奇力量的地方。
为了能在设备上执行代码，程序员需要写一个特殊的函数（kernel函数）。
这个函数需要使用OpenCL语言编写。
OpenCL语言采用了C语言的一部分加上一些约束、关键字和数据类型。
在主机上运行的程序提供了API，所以i可以管理你在设备上运行的程序。
主机程序可以用C或者C++编写，它控制OpenCL的环境（上下文，指令队列…）。
设备（Device）
我们来简单的说一下设备。
设备，像上文介绍的一样，OpenCL编程最给力的地方。
我们必须了解一些基本概念：
Kernel：
你可以把它想像成一个可以在设备上执行的函数。
当然也会有其他可以在设备上执行的函数，但是他们之间是有一些区别的。
Kernel是设备程序执行的入口点。
换言之，Kernel是唯一可以从主机上调用执行的函数。
现在的问题是：
我们如何来编写一个Kernel？
在Kernel中如何表达并行性？
它的执行模型是怎样的？
解决这些问题，我们需要引入下面的概念：
SIMT：
单指令多线程（SINGLEINSTRUCTIONMULTITHREAD）的简写。
就像这名字一样，相同的代码在不同线程中并行执行，每个线程使用不同的数据来执行同一段代码。
Work-item（工作项）：
Work-item与CUDAThreads是一样的，是最小的执行单元。
每次一个Kernel开始执行，很多（程序员定义数量）的Work-item就开始运行，每个都执行同样的代码。
每个work-item有一个ID，这个ID在kernel中是可以访问的，每个运行在work-item上的kernel通过这个ID来找出work-item需要处理的数据。
Work-group（工作组）：
work-group的存在是为了允许work-item之间的通信和协作。
它反映出work-item的组织形式（work-group是以N维网格形式组织的，N=1，2或3）。
Work-group等价于CUDAthreadblocks。
像work-items一样，work-groups也有一个kernel可以读取的唯一的ID。
ND-Range:
ND-Range是下一个组织级别，定义了work-group的组织形式（ND-Rang以N维网格形式组织的，N=1，2或3）；
这是ND-Range组织形式的例子
Kernel
现在该写我们的第一个kernel了。
我们写一个小的kernel将两个向量相加。
这个kernel需要四个参数：
两个要相加的向量，一个存储结果的向量，和向量个数。
如果你写一个程序在cpu上解决这个问题，将会是下面这个样子：
voidvector_add_cpu（constfloat*src_a,
constfloat*src_b,
float*res,
constintnum）
{
for（inti=0;ires[i]=src_a[i]+src_b[i];
}
在GPU上，逻辑就会有一些不同。
我们使每个线程计算一个元素的方法来代替cpu程序中的循环计算。
每个线程的index与要计算的向量的index相同。
我们来看一下代码实现：
__kernelvoidvector_add_gpu（__globalconstfloat*src_a,
__globalconstfloat*src_b,
__globalfloat*res,
constintnum）
{
/*get_global_id（0）返回正在执行的这个线程的ID。
许多线程会在同一时间开始执行同一个kernel，
每个线程都会收到一个不同的ID，所以必然会执行一个不同的计算。
*/
constintidx=get_global_id（0）;
/*每个work-item都会检查自己的id是否在向量数组的区间内。
如果在，work-item就会执行相应的计算。
*/
if（idxres[idx]=src_a[idx]+src_b[idx];
}
有一些需要注意的地方：
1.Kernel关键字定义了一个函数是kernel函数。
Kernel函数必须返回void。
2.Global关键字位于参数前面。
它定义了参数内存的存放位置。
另外，所有kernel都必须写在“.cl”文件中，“.cl”文件必须只包含OpenCL代码。
主机（Host）
我们的kernel已经写好了，现在我们来写host程序。
建立基本OpenCL运行环境
有一些东西我们必须要弄清楚：
Plantform（平台）：
主机加上OpenCL框架管理下的若干设备构成了这个平台，通过这个平台，应用程序可以与设备共享资源并在设备上执行kernel。
平台通过cl_plantform来展现，可以使用下面的代码来初始化平台：
//Returnstheerrorcode
cl_intoclGetPlatformID（cl_platform_id*platforms）//Pointertotheplatformobject
Device（设备）：
通过cl_device来表现，使用下面的代码：
//Returnstheerrorcode
cl_intclGetDeviceIDs（cl_platform_idplatform,
cl_device_typedevice_type,//Bitfieldidentifyingthetype.FortheGPUweuseCL_DEVICE_TYPE_GPU
cl_uintnum_entries,//Numberofdevices,typically1
cl_device_id*devices,//Pointertothedeviceobject
cl_uint*num_devices）//Putsherethenumberofdevicesmatchingthedevice_type
Context（上下文）：
定义了整个OpenCL化境，包括OpenCLkernel、设备、内存管理、命令队列等。
上下文使用cl_context来表现。
使用以下代码初始化：
//Retursthecontext
cl_contextclCreateContext（constcl_context_properties*properties,//Bitwisewiththeproperties（seespecification）
cl_uintnum_devices,//Numberofdevices
constcl_device_id*devices,//Pointertothedevicesobject
void（*pfn_notify）（constchar*errinfo,constvoid*private_info,size_tcb,void*user_data）,//（don'tworryaboutthis）
void*user_data,//（don'tworryaboutthis）
cl_int*errcode_ret）//errorcoderesult
Command-Queue（指令队列）：
就像它的名字一样，他是一个存储需要在设备上执行的OpenCL指令的队列。
“指令队列建立在一个上下文中的指定设备上。
多个指令队列允许应用程序在不需要同步的情况下执行多条无关联的指令。
”
cl_command_queueclCreateCommandQueue（cl_contextcontext,
cl_device_iddevice,
cl_command_queue_propertiesproperties,//Bitwisewiththeproperties
cl_int*errcode_ret）//errorcoderesult
下面的例子展示了这些元素的使用方法：
cl_interror=0;//Usedtohandleerrorcodes
cl_platform_idplatform;
cl_contextcontext;
cl_command_queuequeue;
cl_device_iddevice;
//Platform
error=oclGetPlatformID（&platform）;
if（error!
=CL_SUCCESS）{
cout<<"Errorgettingplatformid:
"<exit（error）;
}
//Device
error=clGetDeviceIDs（platform,CL_DEVICE_TYPE_GPU,1,&device,NULL）;
if（err!
=CL_SUCCESS）{
cout<<"Errorgettingdeviceids:
"<exit（error）;
}
//Context
context=clCreateContext（0,1,&device,NULL,NULL,&error）;
if（error!
=CL_SUCCESS）{
cout<<"Errorcreatingcontext:
"<exit（error）;
}
//Command-queue
queue=clCreateCommandQueue（context,device,0,&error）;
if（error!
=CL_SUCCESS）{
cout<<"Errorcreatingcommandqueue:
"<exit（error）;
}
分配内存
主机的基本环境已经配置好了，为了可以执行我们的写的小kernel，我们需要分配3个向量的内存空间，然后至少初始化它们其中的两个。
在主机环境下执行这些操作，我们需要像下面的代码这样去做：
constintsize=1234567
float*src_a_h=newfloat[size];
float*src_b_h=newfloat[size];
float*res_h=newfloat[size];
//Initializebothvectors
for（inti=0;isrc_a_h=src_b_h=（float）i;
}
在设备上分配内存，我们需要使用cl_mem类型，像下面这样：
//Returnsthecl_memobjectreferencingthememoryallocatedonthedevice
cl_memclCreateBuffer（cl_contextcontext,//Thecontextwherethememorywillbeallocated
cl_mem_flagsflags,
size_tsize,//Thesizeinbytes
void*host_ptr,
cl_int*errcode_ret）
flags是逐位的，选项如下：
CL_MEM_READ_WRITE
CL_MEM_WRITE_ONLY
CL_MEM_READ_ONLY
CL_MEM_USE_HOST_PTR
CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR–从host_ptr处拷贝数据
我们通过下面的代码使用这个函数：
constintmem_size=sizeof（float）*size;
//Allocatesabufferofsizemem_sizeandcopiesmem_sizebytesfromsrc_a_h
cl_memsrc_a_d=clCreateBuffer（context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,mem_size,src_a_h,&error）;
cl_memsrc_b_d=clCreateBuffer（context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,mem_size,src_b_h,&error）;
cl_memres_d=clCreateBuffer（context,CL_MEM_WRITE_ONLY,mem_size,NULL,&error）;
程序和kernel
到现在为止，你可能会问自己一些问题，比如：
我们怎么调用kernel？
编译器怎么知道如何将代码放到设备上？
我们怎么编译kernel？
下面是我们在对比OpenCL程序和OpenCLkernel时的一些容易混乱的概念：
Kernel：
你应该已经知道了，像在上文中描述的一样，kernel本质上是一个我们可以从主机上调用的，运行在设备上的函数。
你或许不知道kernel是在运行的时候编译的！
更一般的讲，所有运行在设备上的代码，包括kernel和kernel调用的其他的函数，都是在运行的时候编译的。
这涉及到下一个概念，Program。
Program：
OpenCLProgram由kernel函数、其他函数和声明组成。
它通过cl_program表示。
当创建一个program时，你必须指定它是由哪些文件组成的，然后编译它。
你需要用到下面的函数来建立一个Program：
//ReturnstheOpenCLprogram
cl_programclCreateProgramWithSource（cl_contextcontext,
cl_uintcount,//numberoffiles
constchar**strings,//arrayofstrings,eachoneisafile
constsize_t*lengths,//arrayspecifyingthefilelengths
cl_int*errcode_ret）//errorcodetobereturned
当我们创建了Program我们可以用下面的函数执行编译操作：
cl_intclBuildProgram（cl_programprogram,
cl_uintnum_devices,
constcl_device_id*device_list,
constchar*options,//Compileroptions,seethespecificationsformoredetails
void（*pfn_notify）（cl_program,void*user_data）,
void*user_data）
查看编译log，必须使用下面的函数：
cl_intclGetProgramBuildInfo（cl_programprogram,
cl_device_iddevice,
cl_program_build_infoparam_name,//Theparameterwewanttoknow
size_tparam_value_size,
void*param_value,//Theanswer
size_t*param_value_size_ret）
最后，我们需要“提取”program的入口点。
使用cl_kernel：
cl_kernelclCreateKernel（cl_programprogram,//Theprogramwherethekernelis
constchar*kernel_name,//Thenameofthekernel,i.e.thenameofthekernelfunctionasit'sdeclaredinthecode
cl_int*errcode_ret）
注意我们可以创建多个OpenCLprogram，每个program可以拥有多个kernel。
以下是这一章节的代码：
//Createstheprogram
//UsesNVIDIAhelperfunctionstogetthecodestringandit'ssize（inbytes）
size_tsrc_size=0;
constchar*path=shrFindFilePath（"vector_add_gpu.cl",NULL）;
constchar*source=oclLoadProgSource（path,"",&src_size）;
cl_programprogram=clCreateProgramWithSource（context,1,&source,&src_size,&error）;
assert（error==CL_SUCCESS）;
//Buildstheprogram
error=clBuildProgram（program,1,&device,NULL,NULL,NULL）;
assert（error==CL_SUCCESS）;
//Showsthelog
char*build_log;
size_tlog_size;
//Firstcalltoknowthepropersize
clGetProgramBuildInfo（program,device,CL_PROGRAM_BUILD_LOG,0,NULL,&log_size）;
build_log=newchar[log_size+1];
//Secondcalltogetthelog
clGetProgramBuildInfo（program,device,CL_PROGRAM_BUILD_LOG,log_size,build_log,NULL）;
build_log[log_size]='\0';
cout<delete[]build_log;
//Extractingthekernel
cl_kernelvector_add_kernel=clCreateKernel（program,"vector_add_gpu",&error）;
assert（error==CL_SUCCESS）;
运行kernel
一旦我们的kernel建立好，我们就可以运行它。
首先，我们必须设置kernel的参数：
cl_intclSetKernelArg（cl_kernelkernel,//Whichkernel
cl_uintarg_index,//Whichargument
size_targ_size,//Sizeofthenextargument（notofthevaluepointedbyit!
）
constvoid*arg_value）//Value
每个参数都需要调用一次这个函数。
当所有参数设置完毕，我们就可以调用这个kernel：
cl_intclEnqueueNDRangeKernel（cl_command_queuecommand_queue,
cl_kernelkernel,
cl_uintwork_dim,//Chooseifweareusing1D,2Dor3Dwork-itemsandwork-groups
constsize_t*global_work_offset,
constsize_t*global_work_size,//Thetotalnumberofwork-items（musthavework_dimdimensions）
constsize_t*local_work_size,//Thenumberofwork-itemsperwork-group（musthavework_dimdimensions）
cl_uintnum_events_in_wait_list,
constcl_event*event_wait_list,
cl_event*event）
下面是这一章节的代码：
//Enqueuingparameters
//Notethatweinformthesizeofthecl_memobject,notthesizeofthememorypointedbyit
error=clSetKernelArg（vector_add_k,0,sizeof（cl_mem）,&src_a_d）;
error|=clSetKernelArg（vector_add_k,1,sizeof（cl_mem）,&src_b_d）;
error|=clSetKernelArg（vector_add_k,2,sizeof（cl_mem）,&res_d）;
error|=clSetKernelArg（vector_add_k,3,sizeof（size_t）,&size）;
assert（error==CL_SUCCESS）;
//Launchingkernel
constsize_tlocal_ws=512;//Numberofwork-itemsperwork-group
//shrRoundUpreturnsthesmallestmultipleoflocal_wsbiggerthansize
constsize_tglobal_ws=sh
升级会员 