<commit_before>//                           License Agreement
//                For Open Source Computer Vision Library
//
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
//   * Redistribution's of source code must retain the above copyright notice,
//     this list of conditions and the following disclaimer.
//
//   * Redistribution's in binary form must reproduce the above copyright notice,
//     this list of conditions and the following disclaimer in the documentation
//     and/or other materials provided with the distribution.
//
//   * The name of the copyright holders may not be used to endorse or promote products
//     derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors as is and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.

#if cn != 3
#define loadpix(addr) *(__global const T *)(addr)
#define TSIZE (int)sizeof(T)
#else
#define loadpix(addr) vload3(0, (__global const T1 *)(addr))
#define TSIZE ((int)sizeof(T1)*3)
#endif

#define SQSUMS_PTR(ox, oy) mad24(y + oy, src_sqsums_step, mad24(x + ox, cn, src_sqsums_offset))
#define SUMS_PTR(ox, oy) mad24(y + oy, src_sums_step, mad24(x + ox, cn, src_sums_offset))
#define SUMS(ox, oy)    mad24(y+oy, src_sums_step, mad24(x+ox, (int)sizeof(T1)*cn, src_sums_offset))
#define SQ_SUMS(ox, oy) mad24(y+oy, src_sqsums_step, mad24(x+ox, (int)sizeof(T1)*cn, src_sqsums_offset))

inline float normAcc(float num, float denum)
{
    if (fabs(num) < denum)
        return num / denum;
    if (fabs(num) < denum * 1.125f)
        return num > 0 ? 1 : -1;
    return 0;
}

inline float normAcc_SQDIFF(float num, float denum)
{
    if (fabs(num) < denum)
        return num / denum;
    if (fabs(num) < denum * 1.125f)
        return num > 0 ? 1 : -1;
    return 1;
}

#define noconvert

#if cn == 1
#define convertToDT(value) (float)(value)
#elif cn == 2
#define convertToDT(value) (float)(value.x + value.y)
#elif cn == 3
#define convertToDT(value) (float)(value.x + value.y + value.z)
#elif cn == 4
#define convertToDT(value) (float)(value.x + value.y + value.z + value.w)
#else
#error "cn should be 1-4"
#endif

#ifdef CALC_SUM

__kernel void calcSum(__global const uchar * srcptr, int src_step, int src_offset,
                      int cols, int total, __global float * dst)
{
    int lid = get_local_id(0), id = get_global_id(0);

    __local WT localmem[WGS2_ALIGNED];
    WT accumulator = (WT)(0), tmp;

    for ( ; id < total; id += WGS)
    {
        int src_index = mad24(id / cols, src_step, mad24(id % cols, TSIZE, src_offset));
        T src = loadpix(srcptr + src_index);

        tmp = convertToWT(src);

        accumulator = mad(tmp, tmp, accumulator);
    }

    if (lid < WGS2_ALIGNED)
        localmem[lid] = accumulator;
    barrier(CLK_LOCAL_MEM_FENCE);

    if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
        localmem[lid - WGS2_ALIGNED] += accumulator;
    barrier(CLK_LOCAL_MEM_FENCE);

    for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
    {
        if (lid < lsize)
        {
            int lid2 = lsize + lid;
            localmem[lid] += localmem[lid2];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if (lid == 0)
        dst[0] = convertToDT(localmem[0]);
}

#elif defined FIRST_CHANNEL

__kernel void extractFirstChannel( const __global uchar* img, int img_step, int img_offset,
                                   __global uchar* res, int res_step, int res_offset, int rows, int cols)
{
    int x = get_global_id(0);
    int y = get_global_id(1)*PIX_PER_WI_Y;

    if(x < cols )
    {
        #pragma unroll
        for (int cy=0; cy < PIX_PER_WI_Y && y < rows; ++cy, ++y)
        {
            T1 image = *(__global const T1*)(img + mad24(y, img_step, mad24(x, (int)sizeof(T1)*cn, img_offset)));;
            int res_idx = mad24(y, res_step, mad24(x, (int)sizeof(float), res_offset));
            *(__global float *)(res + res_idx) = image;
        }
    }
}

#elif defined CCORR

#if cn==1 && PIX_PER_WI_X==4

__kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_step, int src_offset,
                                        __global const uchar * templateptr, int template_step, int template_offset, int template_rows, int template_cols,
                                        __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
    int x0 = get_global_id(0)*PIX_PER_WI_X;
    int y = get_global_id(1);

    if (y < dst_rows)
    {
        if (x0 + PIX_PER_WI_X <= dst_cols)
        {
            WT sum = (WT)(0);

            int ind = mad24(y, src_step, mad24(x0, (int)sizeof(T1), src_offset));
            __global const T1 * template = (__global const T1*)(templateptr + template_offset);

            for (int i = 0; i < template_rows; ++i)
            {
                for (int j = 0; j < template_cols; ++j)
                {
                    T temp = (T)(template[j]);
                    T src = *(__global const T*)(srcptr + ind + j*(int)sizeof(T1));

                    sum = mad(convertToWT(src), convertToWT(temp), sum);

                }
            ind += src_step;
            template = (__global const T1 *)((__global const uchar *)template + template_step);
            }

            T temp = (T)(template[0]);
            int dst_idx = mad24(y, dst_step, mad24(x0, (int)sizeof(float), dst_offset));
            *(__global float4 *)(dst + dst_idx) = convert_float4(sum);
        }
        else
        {
            WT1 sum [PIX_PER_WI_X];
            #pragma unroll
            for (int i=0; i < PIX_PER_WI_X; i++) sum[i] = 0;

            __global const T1 * src = (__global const T1 *)(srcptr + mad24(y, src_step, mad24(x0, (int)sizeof(T1), src_offset)));
            __global const T1 * template = (__global const T1 *)(templateptr + template_offset);

            for (int i = 0; i < template_rows; ++i)
            {
                for (int j = 0; j < template_cols; ++j)
                {
                    #pragma unroll<commit_msg>Use vload to read unaligned data instead of dereference operator.<commit_after>//                           License Agreement
//                For Open Source Computer Vision Library
//
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
//   * Redistribution's of source code must retain the above copyright notice,
//     this list of conditions and the following disclaimer.
//
//   * Redistribution's in binary form must reproduce the above copyright notice,
//     this list of conditions and the following disclaimer in the documentation
//     and/or other materials provided with the distribution.
//
//   * The name of the copyright holders may not be used to endorse or promote products
//     derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors as is and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.

#if cn != 3
#define loadpix(addr) *(__global const T *)(addr)
#define TSIZE (int)sizeof(T)
#else
#define loadpix(addr) vload3(0, (__global const T1 *)(addr))
#define TSIZE ((int)sizeof(T1)*3)
#endif

#define SQSUMS_PTR(ox, oy) mad24(y + oy, src_sqsums_step, mad24(x + ox, cn, src_sqsums_offset))
#define SUMS_PTR(ox, oy) mad24(y + oy, src_sums_step, mad24(x + ox, cn, src_sums_offset))
#define SUMS(ox, oy)    mad24(y+oy, src_sums_step, mad24(x+ox, (int)sizeof(T1)*cn, src_sums_offset))
#define SQ_SUMS(ox, oy) mad24(y+oy, src_sqsums_step, mad24(x+ox, (int)sizeof(T1)*cn, src_sqsums_offset))

inline float normAcc(float num, float denum)
{
    if (fabs(num) < denum)
        return num / denum;
    if (fabs(num) < denum * 1.125f)
        return num > 0 ? 1 : -1;
    return 0;
}

inline float normAcc_SQDIFF(float num, float denum)
{
    if (fabs(num) < denum)
        return num / denum;
    if (fabs(num) < denum * 1.125f)
        return num > 0 ? 1 : -1;
    return 1;
}

#define noconvert

#if cn == 1
#define convertToDT(value) (float)(value)
#elif cn == 2
#define convertToDT(value) (float)(value.x + value.y)
#elif cn == 3
#define convertToDT(value) (float)(value.x + value.y + value.z)
#elif cn == 4
#define convertToDT(value) (float)(value.x + value.y + value.z + value.w)
#else
#error "cn should be 1-4"
#endif

#ifdef CALC_SUM

__kernel void calcSum(__global const uchar * srcptr, int src_step, int src_offset,
                      int cols, int total, __global float * dst)
{
    int lid = get_local_id(0), id = get_global_id(0);

    __local WT localmem[WGS2_ALIGNED];
    WT accumulator = (WT)(0), tmp;

    for ( ; id < total; id += WGS)
    {
        int src_index = mad24(id / cols, src_step, mad24(id % cols, TSIZE, src_offset));
        T src = loadpix(srcptr + src_index);

        tmp = convertToWT(src);

        accumulator = mad(tmp, tmp, accumulator);
    }

    if (lid < WGS2_ALIGNED)
        localmem[lid] = accumulator;
    barrier(CLK_LOCAL_MEM_FENCE);

    if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
        localmem[lid - WGS2_ALIGNED] += accumulator;
    barrier(CLK_LOCAL_MEM_FENCE);

    for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
    {
        if (lid < lsize)
        {
            int lid2 = lsize + lid;
            localmem[lid] += localmem[lid2];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if (lid == 0)
        dst[0] = convertToDT(localmem[0]);
}

#elif defined FIRST_CHANNEL

__kernel void extractFirstChannel( const __global uchar* img, int img_step, int img_offset,
                                   __global uchar* res, int res_step, int res_offset, int rows, int cols)
{
    int x = get_global_id(0);
    int y = get_global_id(1)*PIX_PER_WI_Y;

    if(x < cols )
    {
        #pragma unroll
        for (int cy=0; cy < PIX_PER_WI_Y && y < rows; ++cy, ++y)
        {
            T1 image = *(__global const T1*)(img + mad24(y, img_step, mad24(x, (int)sizeof(T1)*cn, img_offset)));;
            int res_idx = mad24(y, res_step, mad24(x, (int)sizeof(float), res_offset));
            *(__global float *)(res + res_idx) = image;
        }
    }
}

#elif defined CCORR

#if cn==1 && PIX_PER_WI_X==4

__kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_step, int src_offset,
                                        __global const uchar * templateptr, int template_step, int template_offset, int template_rows, int template_cols,
                                        __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
    int x0 = get_global_id(0)*PIX_PER_WI_X;
    int y = get_global_id(1);

    if (y < dst_rows)
    {
        if (x0 + PIX_PER_WI_X <= dst_cols)
        {
            WT sum = (WT)(0);

            int ind = mad24(y, src_step, mad24(x0, (int)sizeof(T1), src_offset));
            __global const T1 * template = (__global const T1*)(templateptr + template_offset);

            for (int i = 0; i < template_rows; ++i)
            {
                for (int j = 0; j < template_cols; ++j)
                {
                    T temp = (T)(template[j]);
                    T src = vload4(0, (__global const T1*)(srcptr + ind + j*(int)sizeof(T1)));

                    sum = mad(convertToWT(src), convertToWT(temp), sum);

                }
            ind += src_step;
            template = (__global const T1 *)((__global const uchar *)template + template_step);
            }

            T temp = (T)(template[0]);
            int dst_idx = mad24(y, dst_step, mad24(x0, (int)sizeof(float), dst_offset));
            *(__global float4 *)(dst + dst_idx) = convert_float4(sum);
        }
        else
        {
            WT1 sum [PIX_PER_WI_X];
            #pragma unroll
            for (int i=0; i < PIX_PER_WI_X; i++) sum[i] = 0;

            __global const T1 * src = (__global const T1 *)(srcptr + mad24(y, src_step, mad24(x0, (int)sizeof(T1), src_offset)));
            __global const T1 * template = (__global const T1 *)(templateptr + template_offset);

            for (int i = 0; i < template_rows; ++i)
            {
                for (int j = 0; j < template_cols; ++j)
                {
                    #pragma unroll
{'fraction_non_alphanumeric': 0.11313430729056391, 'fraction_numerical': 0.01983202766603148, 'mean_word_length': 2.5819009100101113, 'pattern_counts': {'":': 0, '<': 37, '<?xml version=': 0, '>': 21, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 6, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '403884', 'n_tokens_mistral': 4761, 'n_tokens_neox': 4554, 'n_words': 1467}