opencl:慎用-cl-opt-disable选项编译kernel(可能会导致一些无法解释的问题)

在编译opencl kernel代码时，有一个编译选项-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)=%dn"
            ,sizeof(detected_obj_buf_ptr)
            ,sizeof(im_info)
            ,sizeof(integ_mat)
            ,sizeof(integ_sqare_mat));
        DEBUG_LOG("sizeof(om_info)=%dn,sizeof(hit_mat)=%d,nsizeof(mean_mat)=%d,nsizeof(variance_mat)=%d,nsizeof(const_param)=%dn"
            ,sizeof(om_info)
            ,sizeof(hit_mat)
            ,sizeof(mean_mat)
            ,sizeof(variance_mat)
            ,sizeof(const_param));
        DEBUG_LOG("sum of all parameters %dn"
                                    ,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=%dn",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=%dn",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=%xn",sizeof(om_info),om_info.row_stride);
        DEBUG_LOG("hit_mat PTR =%xn",hit_mat);
        DEBUG_LOG("DIFF row_stride-width=%dn",&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下还没有测试）