专栏首页10km的专栏opencl::kernel中获取local memory size

opencl::kernel中获取local memory size

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

在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代码可以不需要重新编译,就能在另一台设备上运行。

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

我来说两句

0 条评论
登录 后参与评论

相关文章

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

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

    用户1148648
  • cmake:Cannot specify include directories for imported target

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

    用户1148648
  • opencl:异步复制函数的注意事项

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

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

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

    用户1148648
  • 你的微信号,修改成功了吗?

    前段时间安利的微信内测版可以改微信号的朋友不信,没想到今天安卓版微信全开放可以改咯。

    IT小侠公社
  • 精确统计所有库表的大小

    统计每个库每个表的大小是数据治理工作的最基本内容,本文将从抽样统计结果及精确统计结果两方面来统计MySQL的每个库每个表的数据量情况。

    July
  • kubernetes中常用网络插件之Flannel

    Kubernetes中解决网络跨主机通信的一个经典插件就是Flannel。Flannel实质上只是一个框架,真正为我们提供网络功能的是后端的Flannel实现,...

    极客运维圈
  • IIS6+Rewrite3重载隐藏index.php

    [ISAPI_Rewrite] RewriteEngine on RewriteCond %{HTTP:Host} ^phpfs.com$ RewriteRul...

    苦咖啡
  • Kali Linux Web渗透测试手册(第二版) - 9.4 - 绕过web服务器的CORS限制

    跨源资源共享(Cross-OriginResource Sharing, CORS)是在服务器端配置的一组策略,它告诉浏览器服务器是否允许在外部站点(跨源请求)...

    7089bAt@PowerLi
  • Windows命令新解

    在我看来,Windows和Linux是相通的,都是通过命令实现控制,只不过Windows图形化做得好,很多命令不用记,直接点按钮就实现了,而在日常Windows...

    我爱你的一诺

扫码关注云+社区

领取腾讯云代金券