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

时间:2022-06-22
本文章向大家介绍opencl:c++接口(cl.hpp)利用cl::LocalSpaceArg设置__local 参数,主要内容包括其使用实例、应用技巧、基本知识点总结和需要注意事项,具有一定的参考价值,需要的朋友可以参考一下。

版权声明:本文为博主原创文章,转载请注明源地址。 https://blog.csdn.net/10km/article/details/51425981

当我们需要在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