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