opencl/msvc:kernel因为指针对齐方式造成向量类型读写异常
opencl/msvc:kernel因为指针对齐方式造成向量类型读写异常
版权声明：本文为博主原创文章，转载请注明源地址。 https://cloud.tencent.com/developer/article/1433761
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字节对齐了，从而问题解决。