generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Use vload to read unaligned data instead of dereference operator.

Browse Source 
According to opencl 1.2 spec 6.1.5:
    For arguments to a __kernel function declared to be a pointer to a
    data type, the OpenCL compiler can assume that the pointee is always
    appropriately aligned as required by the data type. The behavior of
    an unaligned load or store is undefined, except for the
    vloadn, vload_halfn, vstoren, and vstore_halfn functions defined in
    section 6.12.7.

Original code read data of type T from address not aligned by multiple
of sizeof(T), so the result is incorrect. With this patch, the cases
./opencv_perf_imgproc
--gtest_filter=OCL_ImgSize_TmplSize_Method_MatType_MatchTemplate.MatchTemplate/*
could work well with beignet 0.9.3.

Signed-off-by: Chuanbo Weng <chuanbo.weng@intel.com>
This commit is contained in:
Chuanbo Weng 2014-09-17 19:28:07 +08:00
parent c445ce6125
commit c5552788c5
1 changed files with 1 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
modules/imgproc/src/opencl/match_template.cl
View File

@ -161,7 +161,7 @@ __kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_s
for (int j = 0; j < template_cols; ++j)
{
T temp = (T)(template[j]);
T src = *(__global const T*)(srcptr + ind + j*(int)sizeof(T1));
T src = vload4(0, (__global const T1*)(srcptr + ind + j*(int)sizeof(T1)));
sum = mad(convertToWT(src), convertToWT(temp), sum);