书看的有点慢,还没有上机运行。用了一个下午和一个晚上才把第三单元看完。对opencl的并行编程有了一个初步的印象。下面有贴点阵相乘的例子。
Chapter3:主机编程:数据传输和数据划分
S主机需要设备完成复杂的计算任务,至少需要提供三个东西:
-
所需执行的指令2、包含有待处理数据的缓存3、保存处理结果的缓存。
A:设置内核参数:clSetKernelArg
clSetKernelArg(cl_kernel kernel,cl_uint index, size_t size,const void *value)
其中,index为0表示,内核参数优先被访问。
为1表示,被第二个访问。
最后一个指针参数,指向的是传输给内核函数的数据。指针参数可指向基本的数据类型也可以指向内存对象的指针。或者采样器对象的指针或者NULL。
内存对象——主机和设备间数据传输的标准载体。
内存对象有两种类型:缓存对象和图像对象。
内存对象由cl_mem 结构表示。Eg:处理的数据为像素点,则对应的内存对象就是图像对象。
B缓存对象: 由clCreatBuffer创建。 缓存对象可以作为任何与图像处理无关的数据的载体。
clCreatBuffer(cl_context context,cl_mem_flags options,size_t size,void *host_ptr,cl_int *error)
返回的是cl_mem。options用来配置缓存对象的可访问性以及缓存数据在主机内存中的分配。一定是两个标志值的组合。
内存对象的性质(cl_mem_flags)中的可访问性(只读、只写、可读可写)都是限制的是设备(对缓存对象的访问)而非主机。如果试图修改这种访问性,则会产生未定义的结果。
B.1 分配缓存对象
主机可访问内内存的真正初衷应当是固定内存(pinned memory)——在操作系统创建和释放物理内存的时候,不参与页交换,如此可以提高系统的性能。
创建一个缓存内存的例子:
vec_buff=clCreateBuffer(context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,sizeof(float)*32,vec,&error);//创建了一个缓存内存,来打包一个由32个浮点数组成的数组vec。vec_buff是一个只读型的,它的内存分配由CL_MEM_COPY_HOST_PTR 来控制。这个缓存对象所包装的数据由vec来表示,由这些数据已经在主机上进行了初始化,所以vec也被称为主机指针。
再例如:我要传建一个缓存对象来保存内核的输出数据。一般做法:将缓存对象设置为只写模式。将主机指针设为Null【因为这里,只写缓存指的是设备能够非配内存,但主机不能】.即:
clCreatBuffer(cl_context context , CL_MEM_WRITE_ONLY , size_t size,NULL,&error)
如果是主机向设备传输数据,此时主机指针不能为NULL,需要设定缓存对象的数据分配地址。
例子:创建了两个缓存对象,一个是输入数组(只读),一个是输出数组(只写),然后两次调用设置内核函数参数的函数,分别将输入缓存和输出缓存设为内核参数。
Input_buffer=clCreateBuffer(context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,sizeof(input_vector),input_vector,&error);
Output_buffer=clCreateBuffer(context,CL_MEM_WRITE_ONLY,sizeof(input_vector),NULL,&error);//这里只调用了一个只写的标志值,是因为数据输出所在的内存空间并不在意主机上,所以不需要对缓存对象在主机上的分配方式进行设置了。
clSetKernelArg(kernel,0,sizeof(cl_mem),&input_buffer);
clSetKernelArg(kernel,1,sizeof(cl_mem),&output_buffer);
PS:尽管只读内存对象可以向主机指针一样设为NULL,但是它们也有大小,也需要进行内存分配。
B.2创建子缓存对象:由clCreatSubBuffer创建
当一个内核需要另一个内核的一部分数据时,需要创建一个子缓存对象。
数据类型 Const void* 指的是一个_cl_buffer_region结构,定义如下:
Typedef struct _cl_buffer_region{
Size_t origin;
Size_t size;
cl_buffer_region;//origin域指的是缓存对象内子对象的数据的起始位置,size域表示的是子对象的大小。例子:创建了一个子缓存对象,其所包含的40个浮点数来自包含100个浮点数的缓存对象。子缓存对象的数据起始位置是主缓存对象的第50个浮点数。
Cl_buffer_region region;
Region.size=40*sizeof(float);
Region.origin=50*sizeof(float);
sub_buffer=clCreatSubBuffer(main_buffer,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,CL_BUFFER_CREATE_TYPE_REGION,®ion,&error);
C图像对象:由cl_mem结构表示
Opencl 是一种少数能够面向显卡编程的语言中的一种。
C.1创建图像对象
创建2维图像:clCreatImage2D;创建3维图像:clCreatImage3D
clCreatImage2D(cl_context context,cl_mem_flags opts , const cl_image_format *format,size_t width,size_t height , size_t row_pitch,void *data , cl_int *error)//const cl_image_format *format 是图像中像素点的存储格式。这个参数的数据类型是cl_inmage_format结构。
Typedef struce _cl_image_format{
Cl_channel_order image_channel_order;//描述的是像素点所需要的色彩通道及各个通道的存储顺序。其枚举类型包括:CL_RGBx (其中,x表示为位填充)CL_LUMINANCE(表示灰度图像)
Cl_channel_type image_channel_data_type;//图像通道在比特水平方向的呈现方式,包括了通道的数值格式以及每个通道的比特数。
}cl_image_format;
Eg:完成对 cl_image_format结构初始化并将它的像素点模型设置为24位RGB格式。
cl_image_format rgb_format;
rgb_format.image_channel_order=CL_RGB;
rgb_format.image_data_type=CL_UNSIGNED_INT8;
Eg:创建一个三维图像对象,包含了四个图像截面,每个界面都包含了64*80个像素点。
#DEFINE NUM_ROWS 64
#DEFINE NUM_COLS 80
#DEFINE NUM_SLICES 4
unsigned charimage_data[NUM_SLICES][NUM_ROWS][NUM_COLS];
Cl_mem image_object=clCreatImage3D(context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,rgb_format,NUM_COLS,NUM_ROWS,NUM_SLICES,NULL,NULL,image_data,&error);
alpha透明度
C.2 获取图像对象信息:getImageInfo (提供唯独和像素点格式的信息)
D获取缓存对象的相关信息:clGetMemObjectInfo(可以反悔图像对象和缓存对象两种内存对象的信息)
CL_MEM_FLAGS返回用于配置内存对象的可访问性和分配的标志位。
CL_MEM_HOST_PTR返回指向内存对象数据的主机指针。
E内存对象的传输命令:
入列——主机在访问设备之前先创建一个命令队列,然后主机将命令发送到这个队列中,以此来和设备通信,我们将这个发送命令到命令队列的过程称为入列。
E.1 读写数据传输。
clEnqueueReadBuffer来完成设备到主机的数据传输(是读写内存对向的六个函数之一)。
offset表示所读写的缓存数据的起始位置。
Data_size表示从offset开始的待传数据的大小。
这六个读写内存对象的函数,都包含一个名为blocking的布尔型向量。,若其值位CL_TRUE,则函数只有在完成读写操作之后才可返回。反之为:CL_FALSE,則可在结束读写操作之前才可返回。
读写——都是主机的操作。READ是将缓存或者图像对象中的数据读到主机内存。而WRITE则是将主机内存中的数据写到缓存或者图像对象中去。
Origin[3]/region[3]——表示的是传入或传出图像所含图像数据的矩形区域。origin表示待访问的第一个像素点的位置,数组中的三个size_t表示像素点的行列以及截面。
Region中三个size_t分别表示的是待读写操作的图像数据的维度信息(宽高以及深度)。如果是二维的图像,则origin的最后一个元素必须为0,region的最后一个元素必须为1。
Eg:已经有一个矩形保存在缓存对象中,而你需要将其子矩阵读到主机内存中。
:首先创建并初始化一个缓存对象——>然后将这个缓存对象设置为内核参数——>再将内核命令入列——>再向这个缓存写数据并将数据以矩形区域的格式从缓存的指定位置读取数据。
E.2映射内存对象
内存映射——一般的c、c++应用程序访问文件的做法是将文件的内容先放到进程内存中,然后通过内存操作的方法读取、修改文件,该过程称为内存映射
Opencl也提供了相似的机制来访问对象。
映射和解映射内存对象的函数(三个):
clEnqueueMapBuffer(cl_command_queue queue,cl_men buffer,cl_bool blocking,cl_map_flags map_flags,size_t offset,size_t data_size,cl_uint num_events,const cl_event *wait_list,cl_event *event,cl_int *errcode_ret)//将缓存对象中的区域映射到主机内存中。
有两个映射函数返回值是void指针,这个指针两个作用:
-
表示主机上映射的起始位置2、表示整个映射
映射/解映射,区别于读/写函数的另一个区别是:其中的map_flag 参数是用来配置主机上的映射内存的可访问性。
用opencl来操作映射内存的数据常分为三步:
-
调用函数clEnqueueMapBuffer或函数clEnqueueMapImage将内存映射命令入列。然后是使用诸如memcpy之类的函数对内存中的数据进行传输操作。最后调用clEnqueueUnmapMemObject函数解映射内存。
内存映射的运行性能远高于普通的读/写函数
E.3内存对象之间的数据复制
clEnqueueCopyBuffer\clEnqueueCopyImage\clEnqueueCopyBufferToImage等等。
F数据划分
在单个设备上进行数据划分的函数只有一个clEnqueueNDRangeKernel,它和clEnqueueTask不同,因为前者允许在设备上的不同处理资源间分配、执行内核。
F.1循环和工作项
工作项——最内层的一次循环称为一个工作项。
opencl中,你不许奥在内核中配置循环语句,因为内核只会执行最内侧的循环。
内核和工作项的区别:
内核表示的是一系列处理数据的任务所组成的集合。而工作项只是针对某个数据集的内核实现。不同的内核,工作项的数量也不同。
全局ID唯一标识某个工作项,允许其访问所要处理的数据。如process(point[i][j][k])称作内核,则
process(point[1][2][3])称作一个工作项
元组{i,j,k}被称作工作项的全局ID。
维度——一个全局ID中的元素数量,最小值为1
F.2 工作项的大小和偏移量
索引空间
F.3例子:
图片贴不了,,
F.4工作组和计算单元
工作组——由访问相同处理资源的工作项所组成。
它有两个优势:
-
工作组中的工作项可以访问高速内存(也称局部内存)的同一块内存
-
工作组中的工作项可以通过栅栏和障碍的方式来进行同步
处理单元——在opencl中,能够支持工作组的处理资源
各个工作组都可以在单个处理单元上执行,而各个处理单元一次则只能执行一个工作组。创建工作组并非是必须的。
本章总结:
本章目的是讨论Opencl如何分配和处理数据的。
Opencl提供了一个内存对象,cl_mem结构,作为主机和设备之间的数据传输的载体。内存对象的传输过程很简单:先是通过存在的数据创建内存对象,然后调用clSetKernelArg将这个内存对象设置为内核的参数。如此,内核在执行过程中,便可以将其作为函数参数来访问。如果设备还需要传输数据回主机,或将数据复制到另外的缓存对象中,主机可以另外再发相应的命令。
内存对象共两种。缓存对象以一维的形式保存普通数据;图像对象以二维或三维的形式保存格式化的像素点数据。针对这俩不同的对象,Opencl分别提供了相应的函数入列数据传输命令。尽管读/写函数针对的都是设备上的内存对象和主机之间的数据传输,但也可以通过将设备上的内存对象映射到主机内存中,来提高计算性能。
数据划分是每个高性能opencl应用程序所不可回避的问题。分割操作的基本单元是工作项,对应对的是c/c++程序中的每次迭代。每个工作项都会有一个全局ID,用来标识代码所访问的数据位置。如果工作项之间还需要进行同步化处理,那么就需要对这些工作组划分。各个工作组将会分配到设备上相应的计算单元来执行。