专栏首页10km的专栏opencl/msvc:kernel因为指针对齐方式造成向量类型读写异常

opencl/msvc:kernel因为指针对齐方式造成向量类型读写异常

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

kernel中向量数据读写的两种方式

opencl knernel中对全局内存(__global)向量类型(vector data type)数据的读写有两种方式, 一种是直接用=操作符赋值,一种则是通过vstoren,vloadn函数来实现向量数据读写。 示例如下:

#ifdef __OPENCL_VERSION__
// 当为kernel编译器时 cl_int等价于int
typedef int     cl_int;
// 当为kernel编译器时 cl_float4等价于float4
typedef float4      cl_float4;
#endif
typedef struct  _detected_objects_buffer {
    cl_float4 storage[1024];
    cl_int  detected_num;
    kernel_error status;
}detected_objects_buffer;
__kernel void object_cluster(
        __global detected_objects_buffer* global_ptr
         ){ 
    float4 obj;
    int i=0;
    ... // other codes    
    // global_ptr为全局(__global)内存指针
    //向__global指针读写向量数据之方法一:=操作符直接赋值
    global_ptr->storage[i]=obj; // 向__global内存中写入向量数据
    obj=global_ptr->storage[i];// 读取__global内存中向量数据
    //向__global指针读写向量数据之方法二:调用vstoren/vloadn函数  
    vstore4( obj ,i,(__global float*)global_ptr->storage);// 向__global内存中写入向量数据
    obj= vload4( i,(__global float*)global_ptr->storage);// 读取__global内存中向量数据
    ... // other codes
} 

alignment的区别

第一种直接赋值的方式,貌似很简单,第二种则略显复杂,从代码方便性来说,我肯定选择第一种, 但是,请注意,使用两种方式访问__global内存数据,对数据的对齐要求是不一样的: 对于第二种用 vloadn/vstoren读写方式,只要求__global内存指针以向量元素类型的字节长度对齐(参见opencl vloadn/opencl vstoren的opencl原文说明)。 比如上面示例中的float4类型向量,其元素类型为float,float的字节长度为4,所以用vloadn/vstoren读写__global内存指针指向的float4类型向量数据,内存指针只要满足4字节对齐,就可以了。 而第一种直接=操作符赋值的方式,看着写法是简单,但它要求只要求__global内存指针必须以向量总的字节长度对齐。还以float4为例,float4有4个float组成,一共是16个字节,也就是说,用=操作符直接赋值的方式读写__global内存指针指向的float4类型的向量数据的时候,__global内存指针必须是16字节对齐的,否则kernel在运行中可能会抛出异常! 这就是我上一篇博文遇到的问题的根本原因《opencl:一个关于向量赋值的异常》

上一个问题的原因分析

第一种方式对内存地址对齐方式有要求,但从opencl官方的原文档中并没有找到这种提示或说明。是为什么呢?因为OpenCL只是个并行计算标准框架,具体的实现还是由OpenCL设备厂商来完成,每个厂商的OpenCL实现对内存对齐的要求并不一定一样。 我开发用的是AMD APP SDK ,我的电脑并没有gpu显示卡,所以在我的电脑上AMD APP SDK 是在4核的CPU(Core2 Quad Q6600 2.4G)来提供OpenCL计算能力的。Core2 Quad Q6600支持SSE2指令,所以具体的所有OpenCL运算最终都是通过SSE指令来完成的,其中当然包括了内存向量读写指令 ,SSE指令中从内存读取向量数据的函数是_mm_load_ps,参见SSE _mm_load_ps 说明,

说明中有一条很重要的提示就是:

The address must be 16-byte aligned.//地址必须16字节对齐

我们再回头看看这个数据结构定义

#ifdef __OPENCL_VERSION__
// 当为kernel编译器时 cl_int等价于int
typedef int     cl_int;
// 当为kernel编译器时 cl_float4等价于float4
typedef float4      cl_float4;
#endif
typedef struct  _detected_objects_buffer {
    cl_float4 storage[1024]; 
    cl_int  detected_num;
    kernel_error status;
}detected_objects_buffer;

这个结构定义在kernel端编译的时候,因为kernel中的float4是16字节对齐的,所以detected_objects_buffer结构体本身就是16字节对齐的。 但是在主机端cl_float4是这样定义的:

typedef union
{
    cl_float  CL_ALIGNED(16) s[4];// CL_ALIGNED指定16字节对齐
#if __CL_HAS_ANON_STRUCT__
   __CL_ANON_STRUCT__ struct{ cl_float   x, y, z, w; };
   __CL_ANON_STRUCT__ struct{ cl_float   s0, s1, s2, s3; };
   __CL_ANON_STRUCT__ struct{ cl_float2  lo, hi; };
#endif
#if defined( __CL_FLOAT2__) 
    __cl_float2     v2[2];
#endif
#if defined( __CL_FLOAT4__) 
    __cl_float4     v4;
#endif
}cl_float4;
//摘自 cl_platform.h

看上面这个定义,貌似cl_float4也是16字节对齐的,因为明显有CL_ALIGNED(16)嘛! 但是我们再看CL_ALIGNED宏的定义

/* Define alignment keys */
#if defined( __GNUC__ )
    #define CL_ALIGNED(_x)          __attribute__ ((aligned(_x)))
#elif defined( _WIN32) && (_MSC_VER)
    /* Alignment keys neutered on windows because MSVC can't swallow function arguments with alignment requirements     */
    /* http://msdn.microsoft.com/en-us/library/373ak2y1%28VS.71%29.aspx                                                 */
    /* #include <crtdefs.h>                                                                                             */
    /* #define CL_ALIGNED(_x)          _CRT_ALIGN(_x)                                                                   */
    #define CL_ALIGNED(_x)
#else
   #warning  Need to implement some method to align data here
   #define  CL_ALIGNED(_x)
#endif
//摘自 cl_platform.h

靠!原来在MSVC下CL_ALIGNED定义的空的! 正因为这样,所以我在MSVC下编译的时候,cl_float4仍然是4字节对齐。这就造成我自己定义的结构体detected_objects_buffer也是4字节对齐,当使用CL_MEM_USE_HOST_PTR(即kernel直接使用主机内存地址的数据)模式向kernel传递这个结构体指针后,kernel用=操作符读写其中的float4向量时会抛出异常。

参见 OpenCL Specification(Page367)(http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf#page=231)

解决方案

现在我们知道,vloadn/vstoren读写内存向量数据因为对内存对齐要求低,所以相比是最安全的一种方式,但从性能上来说,=操作符直接赋值这种16字节对齐方式的内存读写却是更快的。如果还是希望在kernel中使用=操作符直接赋值来读写向量数据,该怎么办呢?

方案1: 避免使用CL_MEM_USE_HOST_PTR模式向kernel传递数据。 在向kernel传递数据的时候,不要使用CL_MEM_USE_HOST_PTR(即kernel直接使用主机内存地址的数据),而是CL_MEM_COPY_HOST_PTR(即将主机数据复制到opencl设备内存)这种最安全的方式。因为CL_MEM_COPY_HOST_PTR模式下OpenCL设备会为从主机复制来的数据分配内存,在分配内存的时候,会以根据你的结构定义确定合适的对齐模式,后续kernel对内存向量数据读写与主机端的数据无关。所以CL_MEM_COPY_HOST_PTR这种模式下,对内存对齐的要求比较低。 方案2: 更换编译器,使用gcc编译。 从上面cl_float4的定义可以知道,用gcc下编译的时候,cl_float4确实是16字节对齐的,所以用gcc编译就不会存在这个问题。所以更换gcc编译器也是个解决方法。

方案3: 修改你的数据结构定义,以满足在主机端编译时向量数据对齐的要求。 如果你坚持使用CL_MEM_USE_HOST_PTR模式向kernel传递数据,坚持使用MSVC编译器,可以修改数据结构定义,加上align指令,以满足在MSVC下编译时让自定义的数据结构满足向量数据对齐要求。

还以detected_objects_buffer这个结构体为例,修改后的代码如下:

// 新定义一个_CL_CROSS_ALIGN_宏,只在MSVC下有效
#ifdef _MSC_VER
#define _CL_CROSS_ALIGN_(n) __declspec( align(n) )
#else
#define _CL_CROSS_ALIGN_(n) 
#endif /*_MSC_VER*/
#define _CL_CROSS_ALIGN_16 _CL_CROSS_ALIGN_(16)
#ifdef __OPENCL_VERSION__
typedef int     cl_int;// 当为kernel编译器时 cl_int等价于int
typedef float4      cl_float4;// 当为kernel编译器时 cl_float4等价于float4
#endif
typedef struct  _detected_objects_buffer {
    // must 16-byte aligned,otherwise will be throw exception from kernel
    _CL_CROSS_ALIGN_16 cl_float4 storage[1024];//MSVC下强制storage以16字节对齐
    cl_int  detected_num;
    kernel_error status;
}detected_objects_buffer;

经过上面修改,MSVC编译时detected_objects_buffer就是16字节对齐了,从而问题解决。

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

我来说两句

0 条评论
登录 后参与评论

相关文章

  • opencl: C++ 接口(cl.hpp)创建kernel

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

    用户1148648
  • opencl:改造C++接口增加对内存编译(compile)的支持

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

    用户1148648
  • jface databinding:使用CheckboxTableViewer实现表中(Set)对象与CheckTable中选中条目数据绑定

    上一篇博文《jface databinding:可多选的widget List组件selection项目与java.util.List对象的双向数据绑定》讲述了...

    用户1148648
  • 一个简单的rest_framework demo

    人生不如戏
  • 你被复制了吗?大数据带你看中国人的名字 重名TOP榜

    名字是一个人的代号,起名字也是一门学问。你想知道中国人重名最多的名字是哪个吗?你想知道中国人重名最多的名都有哪些吗?再看看身边“10后”的娃们,有没有叫子涵、子...

    CSDN技术头条
  • 在ASP.NET中,IE与Firefox下载文件带汉字名时乱码的解决方法

    跟着阿笨一起玩NET
  • 教程 | Hinton 机器学习视频中文版:感知器的几何空间解析(2.3)

    本套课程中,Hinton 重点介绍了人工神经网络在语音识别和物体识别、图像分割、建模语言和人类运动等过程中的应用,及其在机器学习中发挥的作用。与吴恩达的《Mac...

    AI研习社
  • python下以api形式调用tesseract识别图片验证码

    之前在博文中介绍在python中如何调用tesseract ocr引擎,当时主要介绍了shell模式,shell模式需要安装tesseract程序,并且效率相对...

    黯然销魂掌
  • spring:如何用代码动态向容器中添加或移除Bean ?

    先来看一张类图: ? 有一个业务接口IFoo,提供了二个实现类:FooA及FooB,默认情况下,FooA使用@Component由Spring自动装配,如果出于...

    菩提树下的杨过
  • 高考啦! JavaScript高考全国卷

    HTML5学堂-码匠:一年一度的高考~走过路过不要错过,做程序的你,来考考前端“高考题”吧! 高考啦!!! 全国卷 第1题 如下代码的运行结果是? [ ] + ...

    HTML5学堂

扫码关注云+社区

领取腾讯云代金券