opencl:c++接口(cl.hpp)利用cl::LocalSpaceArg设置__local 参数

10km

发布于 2019-05-25 22:04:18

9860

发布于 2019-05-25 22:04:18

文章被收录于专栏：10km的专栏

当我们需要在kernel中使用local memory数组的时候，有两种方式定义local 数组

第一种,编译期静态定义，这是比较普通的使用方式，如下代码，这种方式，在编译期就分配了local 数组的大小。

#define LOCAL_ARRAY_SIZE 64 // LOCAL_ARRAY_SIZE 可以通过编译选项-D在编译的时候定义
__kernel void test_kernel(

 ){
     __local int local_buf[LOCAL_ARRAY_SIZE];
    DEBUG_LOG("(%d,%d)\n",get_local_size(0),get_local_size(1));
}

第二种，kernel运行时指定，如下代码，：

__kernel void test_kernel(
        __local int *local_buf
 ){
    DEBUG_LOG("(%d,%d)\n",get_local_size(0),get_local_size(1));
}

在调用kernel的时候，通过clSetKernelArg(参见 clSetKernelArg官方说明)指定数组的大小

请注意，根据上面clSetKernelArg的参数说明(红线标记部分)，当对于地址修饰符为__local的参数，arg_value指针必须为NULL。

使用opencl的C接口时，这都不是事儿。但是如果使用opencl的C++接口，如何用cl::Kernel::setArg成员函数，设置一个有长度却指针为nullptr的参数呢？这是个不可能完成的任务嘛。

下面是cl::Kernel::setArg的代码

    template <typename T>
    cl_int setArg(cl_uint index, const T &value)
    {
        return detail::errHandler(
            ::clSetKernelArg(
                object_,
                index,
                detail::KernelArgumentHandler<T>::size(value),
                detail::KernelArgumentHandler<T>::ptr(value)),
            __SET_KERNEL_ARGS_ERR);
    }

opencl在设计c++接口的时候已经考虑到了这一点，所以提供了一个LocalSpaceArg结构对象用于local地址空间指针参数的设置。

下面代码是LocalSpaceArg的定义，非常简单，就只有一个size_t,指定要分配的local memory字节数。

//! \brief Local address wrapper for use with Kernel::setArg
struct LocalSpaceArg
{
    ::size_t size_;
};
/*! Local
 * \brief Helper function for generating LocalSpaceArg objects.
 */
inline LocalSpaceArg
Local(::size_t size)
{
    LocalSpaceArg ret = { size };
    return ret;
}
// Local函数用于返回一个LocalSpaceArg对象

所以使用opencl C++接口时，设置__local参数，

只需要将要分配的local memory的长度值，封装在LocalSpaceArg结构中再调用cl::Kernel::setArg就成了，

如下：

cl::Kernel kernel;
kernel.setArg(0,cl::LocalSpaceArg{512});//分配512字节的local memory
//也可以使用cl::Local创建cl::LocalSpaceArg对象
kernel.setArg(0,cl::Local(512));//分配512字节的local memory

注意：

当使用这种方式动态分配local memory的时候，因为无法确定local memory的使用量，所以在使用CodeXL进行kernel代码静态分析的时候，只能假设使用了全部local memory，所以有效并发约束(Effective concurrency constraint)永远显示为1