opencl::kernel中获取local memory size

在OpenCL设备中一个workgroup中的所有work-item可以共用本地内存(local memory),在OpenCL kernal编程中,合理的利用local memory,可以提升系统的整体效率。

但是，根据OpenCL的标准，不论在kernel代码的编译期还是运行时，kernel程序在不借助主机端程序的帮助下，是无法知道当前设备(device)的local memory容量的。也就是说，不论是local memory的容量还是其他类型的设备信息，都必须由主机端程序在编译期或运行时告诉kernel。

我们知道，主机程序可以通过clGetDeviceInfo(点击打开函数说明)函数获取local memory size。

那么问题来了：主机如何告诉kernel这些信息？

编译期 -D name=value

在编译kernel的时候，不论你是用clBuildProgram还是用clCompileProgram+clLinkProgram

都可以提供编译选项，而编译选项中的-D 参数允许定义宏。可以通过这个途径，将local memory size定义成一个常量提供给kernel代码。

下面是OpenCL编译器选项的部分说明,参见clBuildProgram函数说明

Compiler Options

The compiler options are categorized as pre-processor options, options for math intrinsics, options that control optimization and miscellaneous options. This specification defines a standard set of options that must be supported by the OpenCL C compiler when building program executables online or offline. These may be extended by a set of vendor- or platform specific options.

Preprocessor Options

These options control the OpenCL C preprocessor which is run on each program source before actual compilation.

-D options are processed in the order they are given in the options argument to clBuildProgram or or clCompileProgram.

-D name
Predefine name as a macro, with definition 1.

-D name=definition
The contents of definition are tokenized and processed as if they appeared during translation phase three in a `#define' directive. In particular, the definition will be truncated by embedded newline characters.

-I dir
Add the directory dir to the list of directories to be searched for header files.

以下是我的C++代码片段

// 当OpenCL设备只有1个时，定义CL_DEVICE_LOCAL_MEM_SIZE
if (1 == _devices.size()) {
// 如果设备不支持local memory,抛出异常
throw_if(CL_NONE == _devices[0].getInfo<CL_DEVICE_LOCAL_MEM_TYPE>(),
_devices[0].getInfo<CL_DEVICE_NAME>().append("not support local memory"))
// 获取设备的local memory size
auto local_mem_size=_devices[0].getInfo<CL_DEVICE_LOCAL_MEM_SIZE>();
// 在编译选项中加入 -D CL_DEVICE_LOCAL_MEM_SIZE=%local_mem_size%
add_define("CL_DEVICE_LOCAL_MEM_SIZE", local_mem_size);
}

上面代码中add_define模板函数的实现

template<typename T>
builder &add_define(const std::string &def,const T &v){
throw_if(def.empty(),"def is empty")
std::stringstream stream;
stream<<def<<"="<<v;
_compiler_options+="-D "+stream.str()+" ";
return *this;
}

运行时clSetKernelArg

如果要在kernel运行时,告诉kernel local memory size,就要在kernel代码中增加参数

下面是kernel代码：

__kernel void local_test(__local char*p,int local_size){
for(int i=0;i<local_size;++i){
p[i]=i;
}
}

下面是主机端代码：

cl_int local_mem_size;
..... //调用clGetDeviceInfo获取local memory size赋值给local_mem_size
//设置kernel的第一个参数，
//因为local_test的参数p定义为__local，所以不需要指定参数地址，
//opencl设备会根据第三个参数的值分配相应字节数的local memory.
//所以clSetKernelArg中最后一个参数只需要填NULL
clSetKernelArg(kernel,0,local_mem_size,NULL);
//设置kernel的第二个参数，告诉kernel 数组p的长度
clSetKernelArg(kernel,1,size(local_mem_size),&local_mem_size);
....//调用 clEnqueueNDRangeKernel执行kernel

总结

以上两种办法，各有优劣，所以具体使用哪种方法更合适，这真的根据你的需要，就我个人而言我采用第一种方法，因为第一种办法，直接在编译期就可以根据local memory大小来分配数组大小，这样以来kernel代码不需要带过多参数，代码维护性好一点点。

第二种办法因为在kernel运行时才确定local memory size,所以它的优点是一个设备的kernel代码可以不需要重新编译，就能在另一台设备上运行。