Opencl API解释(一)
首先注明:我用的AMD的opencl,它有很多sample代码,结合代码来解释这些API
Opencl 常用的API 汇总总结:
信息查询函数
1.
cl_int clGetDeviceInfo(cl_device_id device,
cl_device_info param_name,
size_t param_value_size,
void * param_value,
size_t *param_value_size_ret )
参数说明
此函数用来查询OpenCL设备信息。首先介绍其参数:
第一个参数device是clGetDeviceIDs的返回值。
第二个参数param_name是一个枚举常量,标识要查询的设备信息,具体有哪些信息稍后详述。
第三个参数param_value_size声明下一个参数param_value所指向的存储空间的字节大小。这个大小要>=查询参数大小。
第四个参数param_value指向要查询参数返回到的存储空间的地址。NULL时表示忽略。
第五个参数返回查询到的参数的实际的大小。设为NULL则忽略。
下面我们来具体介绍一下可以查询的几个常用参数。
CL_DEVICE_TYPE:OpenCL设备类型。目前支持CL_DEVICE_TYPE_CPU,CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_ACCELERATOR, CL_DEVICE_TYPE_DEFAULT或以上联合。
CL_DEVICE_VENDOR_ID:一个唯一的供应商识别码。一个唯一设备识别码的例子,PCIe ID。
CL_DEVICE_MAX_COMPUTE_UNITS:OpenCL设备上并行计算单元数目。一个work-group只在一个compute unit上执行。该参数最小为1.
CL_DEVICE_MAX_WORK_ITEM_DEMENSIONS:数据并行执行模块用来声明global和localwork-item IDS的最大维度。该参数最小为3。参考clEnqueueNDRangeKernel,该函数第三个参数work_dim是声明global work-item和work-group中的work-items(specify the global work-items and work items in the work-group.)中使用的维度的数目。work-item应该比0大,且<=CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS.
CL_DEVICE_MAX_WORK_ITEM_SIZES:数组类型,指能被在work-group的每一个维度声明的work-item的最大数目。最小值(1,1,1)。
CL_DEVICE_MAX_WORK_GROUP_SIZE:在一个computeunit中执行一个kernel的work-group中work-item的最大数目。最小值为1。Maximum number of work-items in a work-group executing a kernel ona single compute unit, using the data parallelexecution model. (Refer to clEnqueueNDRangeKernel ). The minimum value is1.
举例:
//Get max compute units
status = clGetDeviceInfo(
deviceId,
CL_DEVICE_MAX_COMPUTE_UNITS,
sizeof(cl_uint),
&maxComputeUnits,
NULL);
CHECK_OPENCL_ERROR(status, "clGetDeviceIDs(CL_DEVICE_MAX_COMPUTE_UNITS) failed");
2.
cl_int clGetKernelWorkGroupInfo(cl_kernelkernel ,
cl_device_iddevice ,
cl_kernel_work_group_info param_name,
size_t param_value_size,
void* param_value,
size_t*param_value_size_ret )
参数说明:
此函数返回指定到某一device上的kernel对象信息。同样,先来看参数:
第一个参数要查询的kernel对象。
第二个参数指定与kernel绑定的设备列表中的某一个设备。这个列表就是与kernel绑定的context对应的kernel列表。如果列表中只有一个device,此处device参数可以为NULL。
第三个参数指定要查询参数名称,这也是个枚举值。
第三个参数是要指定的要查询的参数返回的字节数,要>=返回值。
第四个参数指返回值指向内存空间的地址,若设为NULL则忽略。
第五个参数返回实际查询到的参数的大小。
下面来说几个重要的可查询的参数:
CL_KERNEL_WORK_GROUP_SIZE:查询在某一指定设备上执行一个kernel可以使用的最大work-group的size。OpenCL实现会使用资源,这就要求kernel确定work-group大小。
This provides a mechanism for the applicationto query the maximum work-group size that can be used to execute a kernel on aspecificdevice given by device. The OpenCL implementationuses the resource requirements of the kernel (register usage etc.) to determinewhat this work -group size should be.
举例:
status = clGetKernelWorkGroupInfo(kernel,
device,
CL_KERNEL_LOCAL_MEM_SIZE,
sizeof(cl_ulong),
&localMemoryUsed,
NULL);
if(checkVal(status, CL_SUCCESS, "clGetKernelWorkGroupInfo failed(CL_KERNEL_LOCAL_MEM_SIZE)"))
return SDK_FAILURE;
两个例子均来自SDKCommon.cpp.
work-group/work-item/size等关系说明
为执行一个数据并行kernel,除work-items的数目外也要指定work-groups的数目。这也就是为什么两个参数都必须传递给clEnqueueNDRangeKernel。例如:
size_t global_item_size = 4;//总的线程数
size_t local_item_size = 1;//每一个group的线程数
/* Execute OpenCL kernel as data parallel*/
ret = clEnqueueNDRangeKernel(command_queue,kernel, 1, NULL,
&global_item_size,&local_item_size, 0, NULL, NULL);
这个就表示上面这个数据并行计算的kernel中每一个work-group由1个work-item组成,而共有4个work-items要被处理,即总的work-items要被分成4个work-group。
另外work-item对应硬件上的一个PE(processing element),而一个work-group对应硬件上的一个CU(computing unit)。这种对应可以理解为,一个work-item不能被拆分到多个PE上处理;同样,一个work-group也不能拆分到多个CU上同时处理(忘了哪里看到的信息)。当映射到OpenCL硬件模型上时,每一个work-item运行在一个被称为处理基元(processing element)的抽象硬件单元上,其中每个处理基元可以处理多个work-item(注:摘自《OpenCL异构计算》P87)。(如此而言,是不是说对于二维的globalx必须是localx的整数倍,globaly必须是localy的整数倍?那么如果我数据很大,work-item所能数量很多,如果一个group中中work-item的数量不超过CU中PE的个数,那么group的数量就可能很多;如果我想让group数量小点,那work-item的数目就会很多,还能不能处理了呢?以当前这个示例是能的,但是对于多的work-item,这涉及到如何确定work-item数目的问题。
结合Cuda的概念进行解释:因为实际上,一个 SM 可以允许的 block 数量,还要另外考虑到他所用到 SM 的资源:shared memory、registers 等。在 G80 中,每个 SM 有 16KB 的 shared memory 和 8192 个 register。而在同一个 SM 里的 block 和 thread,则要共享这些资源;如果资源不够多个 block 使用的话,那 CUDA 就会减少 Block 的量,来让资源够用。在这种情形下,也会因此让 SM 的 thread 数量变少,而不到最多的 768 个。
比如说如果一个 thread 要用到 16 个 register 的话(在 kernel 中宣告的变量),那一个 SM 的 8192 个 register 实际上只能让 512 个 thread 来使用;而如果一个 thread 要用 32 个 register,那一个 SM 就只能有 256 个 thread 了~而 shared memory 由于是 thread block 共享的,因此变成是要看一个 block 要用多少的 shread memory、一个 SM 的 16KB 能分给多少个 block 了。
所以虽然说当一个 SM 里的 thread 越多时,越能隐藏 latency,但是也会让每个 thread 能使用的资源更少。因此,这点也就是在优化时要做取舍的了
继续向下解释work-group,work-item,size的关系:
每一个work-group中work-item的数目是不能改变的,始终如一。如果work-item的数目不能在work-groups中均分,clEnqueueNDRangeKernel失败,返回错误码CL_INVALID_WORK_GROUP_SIZE。此处要注意,自己在尝试检测GPU处理能力的时候给出的work-item和work-group的数目不能整除时不一定是数量超限,有可能只是不能整除。
global work-item ID、localwork-item ID,和work-group ID之间的关系如下图所示。
图1 work-group ID和work-item ID
表1 获取ID的函数
函数 |
返回值 |
get_group_id |
Work-group ID |
get_global_id |
Global work-item ID |
get_local_id |
Local work-item ID |
因为要处理2D图像或3D空间,work-items和work-groups可以被指定为2或3维。图2给出一个work-group和work-item被定义为2D的例子。
图2 work-group和work-item定义为2D
因为work-group和work-item可至3维,get_group_id(), get_global_id(), get_local_id()每一个的参数可以是0~2。
注意,空间维度指数和每个work-group中work-item的数目能够依据设备而变化。最大维度指数可以通过clGetDeviceInfo()来获取CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,每个work-group中work-items的最大值可以通过CL_DEVICE_MAX_WORK_ITEM_SIZES获取。前者是cl_uint型,后者是size_t的数组。
在kernel函数中,我们能够通过API调用得到global id以及其他信息:
get_global_id(dim)
get_global_size(dim)
这两个函数能得到每个维度上的global id。
get_group_id(dim)
get_num_groups(dim)
get_local_id(dim)
get_local_size(dim)
这几个函数用来计算group id以及在group内的local id。
get_global_id(0) = column, get_global_id(1) = row
get_num_groups(0) * get_local_size(0) == get_global_size(0)
CL_DEVICE_MAX_WORK_ITEM_SIZES,CL_DEVICE_MAX_WORK_GROUP_SIZE(clGetDeivceInfo获取)它跟CL_KERNEL_WORK_GROUP_SIZE(clGetKernelWorkGroupInfo获取)有什么区别?
CL_DEVICE_MAX_WORK_ITEM_SIZES : Max work-items sizes in each dimensions, 每一个维度允许的最大的work-item数
CL_DEVICE_MAX_WORK_GROUP_SIZE: Max allowed work-items in a group,一个workgroup所允许的最多work-item数。
CL_KERNEL_WORK_GROUP_SIZE: Group size returned by kernel 实际在kernel中执行的workgroup数目。
执行cinfo,可以检测硬件信息
【推荐】FFA 2024大会视频回放:Apache Flink 的过去、现在及未来
【推荐】中国电信天翼云云端翼购节,2核2G云服务器一口价38元/年
【推荐】抖音旗下AI助手豆包,你的智能百科全书,全免费不限次数
【推荐】轻量又高性能的 SSH 工具 IShell:AI 加持,快人一步