opencl:慎用-cl-opt-disable选项编译kernel(可能会导致一些无法解释的问题)
2016-04-21 15:00
489 查看
在编译opencl kernel代码时,有一个编译选项
参见clBuildProgram
但是今天为了调试kernel代码,实际使用这个选项编译kernel却发现,使用这个选项就是坑。使用之后,kernel参数传递都不正常了。
下面这是个很简单的主机端和kernel共用的数据结构
下面是一个有10个参数的kernel函数,其中第二,第六个参数类型都是
kernel函数中不干别的,只打印传入的参数的值。
在主机端给
当正常编译kernel时(不使用
但是当我使用
请注意红框中
开始我以为是我的定义的数据结构的字节对齐问题(
后来我尝试改变参数顺序,将om_info参数放到第9位(倒数第二),居然正常了!
反复修改参数传递顺序进行尝试,最后得到的的规律是:
把所有传值参数放在前面,把指针参数放在后面,所有的参数传递就正常了
so why?
还是没办法解释。尼玛,快被你搞崩溃了!!!
总之,我认为
在网上找了一下,相关资料很少,stackoverflow有类似与
(我的开发平台是VS2015,gcc下还没有测试)
-cl-opt-disable。根据opencl 官网的原文描述,使用这个选项可以关闭所有的代码优化,便于调试程序。(默认情况下,编译优化选项是打开的)
参见clBuildProgram
但是今天为了调试kernel代码,实际使用这个选项编译kernel却发现,使用这个选项就是坑。使用之后,kernel参数传递都不正常了。
下面这是个很简单的主机端和kernel共用的数据结构
typedef struct _matrix_info_cl { cl_uint width; cl_uint height; cl_uint row_stride; }matrix_info_cl;
下面是一个有10个参数的kernel函数,其中第二,第六个参数类型都是
matrix_info_cl
kernel函数中不干别的,只打印传入的参数的值。
__kernel void object_filter1( __constant detected_objects_buffer* detected_obj_buf_ptr ,__constant feather_cl * const feather ,const matrix_info_cl im_info , __constant INTEG_TYPE *integ_mat ,__constant INTEG_TYPE *integ_sqare_mat ,const matrix_info_cl om_info ,__global uchar* hit_mat ,__global float* mean_mat ,__global float* variance_mat ,const filter1_const_param const_param ){ const int y = get_global_id(1); if(0==y){ DEBUG_LOG("sizeof(detected_obj_buf_ptr)=%d,\nsizeof(im_info)=%d,\nsizeof(integ_mat)=%d,\nsizeof(integ_sqare_mat)=%d\n" ,sizeof(detected_obj_buf_ptr) ,sizeof(im_info) ,sizeof(integ_mat) ,sizeof(integ_sqare_mat)); DEBUG_LOG("sizeof(om_info)=%d\n,sizeof(hit_mat)=%d,\nsizeof(mean_mat)=%d,\nsizeof(variance_mat)=%d,\nsizeof(const_param)=%d\n" ,sizeof(om_info) ,sizeof(hit_mat) ,sizeof(mean_mat) ,sizeof(variance_mat) ,sizeof(const_param)); DEBUG_LOG("sum of all parameters %d\n" ,sizeof(detected_obj_buf_ptr) +sizeof(im_info) +sizeof(integ_mat) +sizeof(integ_sqare_mat) +sizeof(om_info) +sizeof(hit_mat) +sizeof(mean_mat) +sizeof(variance_mat) +sizeof(const_param)); DEBUG_LOG("im_info.width=%d,im_info.row_stride=%d,height=%d,size=%d\n",im_info.width,im_info.row_stride,im_info.height,im_info.row_stride*im_info.height); DEBUG_LOG("om_info.width=%d,om_info.row_stride=%d,height=%d,size=%d\n",om_info.width,om_info.row_stride,om_info.height,om_info.row_stride*om_info.height); DEBUG_LOG("sizeof(om_info)=%d om_info.row_stride=%x\n",sizeof(om_info),om_info.row_stride); DEBUG_LOG("hit_mat PTR =%x\n",hit_mat); DEBUG_LOG("DIFF row_stride-width=%d\n",&om_info.row_stride-&om_info.width); } ......// other code }
在主机端给
im_info传递{738,1024,752},给
om_info传递{714,1000,752}。
当正常编译kernel时(不使用
-cl-opt-disable),结果可以预测,kernel打出来的值跟主机端是一样的。
但是当我使用
-cl-opt-disable编译kernel后,再运行,结果就是下面这样:
请注意红框中
om_info.row_stride的值不对了,是个非常大的数,下一行,是以16进制打印出来的
om_info.row_stride的值0x7866000居然是下一个指针参数
hit_mat的值。
开始我以为是我的定义的数据结构的字节对齐问题(
matrix_info_cl是12个字节),但将
matrix_info_cl对齐到16个字节后问题依旧。而且更不可解释的是同样是
matrix_info_cl,第一个参数
im_info却能正确传递。
后来我尝试改变参数顺序,将om_info参数放到第9位(倒数第二),居然正常了!
反复修改参数传递顺序进行尝试,最后得到的的规律是:
把所有传值参数放在前面,把指针参数放在后面,所有的参数传递就正常了
so why?
还是没办法解释。尼玛,快被你搞崩溃了!!!
总之,我认为
-cl-opt-disable选项编译的kernel代码,参数解析时有问题,但找不到原因。
在网上找了一下,相关资料很少,stackoverflow有类似与
-cl-opt-disable相关的莫名其妙的问题(《OpenCL white space influence private memory usage?》),解决办法就是不用
-cl-opt-disable,却没有人知道原因,不清楚这个问题是具体的OpenCL平台实现有关,还是个通病。
(我的开发平台是VS2015,gcc下还没有测试)
相关文章推荐
- 神器SystemTap
- 解读Linux安全机制之栈溢出保护
- 簡單設定 kernel 選項在使用 iptables 前
- Ubuntu12.04内核升级出了问题
- [Linux学习笔记] Linux系统引导流程(一)
- 更新Debian内核e1000e驱动模块
- Linux Kernel 4.5在3月15日发布最终版
- kernel: printk: 2 messages suppressed.
- linux 内核 hash table 的使用
- Linux Kernel Panic报错解决思路
- kernel: TCP: time wait bucket table overflow错误的解决办法
- nginx编译参数选项详解
- linux安全相关
- 内核的主要配置文件的详细说明
- Linux Kernel 2.6.32 Local Root Exploit (x86_64)
- LINUX ulimit命令详解
- Linux Kernel kNFSd 整数溢出 拒绝服务漏洞 .
- 基于busybox和linux kernel制作小linux
- Linux IPv4代码分析系列(1)
- linux 升级kernel