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

当我们需要在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官方说明][1])指定数组的大小
这里写图片描述
请注意,根据上面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

这里写图片描述
[1]:https://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clSetKernelArg.html

评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

10km

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值