天天看點
傳回

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

說明,

opencl/msvc:kernel因為指針對齊方式(alignment)造成向量類型(vector data type)讀寫異常

說明中有一條很重要的提示就是：

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)

opencl/msvc:kernel因為指針對齊方式(alignment)造成向量類型(vector data type)讀寫異常

解決方案

現在我們知道，

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
上一篇: 圖檔透明
下一篇: 靜态頁面清除浏覽器緩存

繼續閱讀