您的位置：首页 > 运维架构

opencl/msvc:kernel因为指针对齐方式(alignment)造成向量类型(vector data type)读写异常

2016-04-19 12:03 531 查看

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字节对齐了，从而问题解决。
内容来自用户分享和网络整理，不保证内容的准确性，如有侵权内容，可联系管理员处理 点击这里给我发消息
标签：  opencl kernel vector align msvc
相关文章推荐
新的分享
章节导航