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

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字节对齐了，从而问题解决。