<commit_before>/*M///////////////////////////////////////////////////////////////////////////////////////
//
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
//  By downloading, copying, installing or using the software you agree to this license.
//  If you do not agree to this license, do not download, install,
//  copy or use the software.
//
//
//                           License Agreement
//                For Open Source Computer Vision Library
//
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
//    <NAME>, <EMAIL>
//
// 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.
//
//M*/


#define CELL_WIDTH 8
#define CELL_HEIGHT 8
#define CELLS_PER_BLOCK_X 2
#define CELLS_PER_BLOCK_Y 2
#define NTHREADS 256
#define CV_PI_F 3.1415926535897932384626433832795f

//----------------------------------------------------------------------------
// Histogram computation

__kernel void compute_hists_kernel(const int width, const int cblock_stride_x, const int cblock_stride_y,
                                   const int cnbins, const int cblock_hist_size, const int img_block_width,
                                   const int grad_quadstep, const int qangle_step,
                                   __global const float* grad, __global const uchar* qangle,
                                   const float scale, __global float* block_hists, __local float* smem)
{
    const int lidX = get_local_id(0);
    const int lidY = get_local_id(1);
    const int gidX = get_group_id(0);
    const int gidY = get_group_id(1);

    const int cell_x = lidX / 16;
    const int cell_y = lidY;
    const int cell_thread_x = lidX & 0xF;

    __local float* hists = smem;
    __local float* final_hist = smem + cnbins * 48;

    const int offset_x = gidX * cblock_stride_x + (cell_x << 2) + cell_thread_x;
    const int offset_y = gidY * cblock_stride_y + (cell_y << 2);

    __global const float* grad_ptr = grad + offset_y * grad_quadstep + (offset_x << 1);
    __global const uchar* qangle_ptr = qangle + offset_y * qangle_step + (offset_x << 1);

    // 12 means that 12 pixels affect on block's cell (in one row)
    if (cell_thread_x < 12)
    {
        __local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) + cell_thread_x;
        for (int bin_id = 0; bin_id < cnbins; ++bin_id)
            hist[bin_id * 48] = 0.f;

        const int dist_x = -4 + cell_thread_x - 4 * cell_x;

        const int dist_y_begin = -4 - 4 * lidY;
        for (int dist_y = dist_y_begin; dist_y < dist_y_begin + 12; ++dist_y)
        {
            float2 vote = (float2) (grad_ptr[0], grad_ptr[1]);
            uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]);

            grad_ptr += grad_quadstep;
            qangle_ptr += qangle_step;

            int dist_center_y = dist_y - 4 * (1 - 2 * cell_y);
            int dist_center_x = dist_x - 4 * (1 - 2 * cell_x);

            float gaussian = exp(-(dist_center_y * dist_center_y + dist_center_x * dist_center_x) * scale);
            float interp_weight = (8.f - fabs(dist_y + 0.5f)) * (8.f - fabs(dist_x + 0.5f)) / 64.f;

            hist[bin.x * 48] += gaussian * interp_weight * vote.x;
            hist[bin.y * 48] += gaussian * interp_weight * vote.y;
        }

        volatile __local float* hist_ = hist;
        for (int bin_id = 0; bin_id < cnbins; ++bin_id, hist_ += 48)
        {
            if (cell_thread_x < 6) hist_[0] += hist_[6];
            if (cell_thread_x < 3) hist_[0] += hist_[3];
            if (cell_thread_x == 0)
                final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = hist_[0] + hist_[1] + hist_[2];
        }
    }

    barrier(CLK_LOCAL_MEM_FENCE);

    __global float* block_hist = block_hists + (gidY * img_block_width + gidX) * cblock_hist_size;

    int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 16 + cell_thread_x;
    if (tid < cblock_hist_size)
        block_hist[tid] = final_hist[tid];
}

//-------------------------------------------------------------
//  Normalization of histograms via L2Hys_norm
//
float reduce_smem(volatile __local float* smem, int size)
{
    unsigned int tid = get_local_id(0);
    float sum = smem[tid];

    if (size >= 512) { if (tid < 256) smem[tid] = sum = sum + smem[tid + 256]; barrier(CLK_LOCAL_MEM_FENCE); }
    if (size >= 256) { if (tid < 128) smem[tid] = sum = sum + smem[tid + 128]; barrier(CLK_LOCAL_MEM_FENCE); }
    if (size >= 128) { if (tid < 64) smem[tid] = sum = sum + smem[tid + 64]; barrier(CLK_LOCAL_MEM_FENCE); }

    if (tid < 32)
    {
        if (size >= 64) smem[tid] = sum = sum + smem[tid + 32];
        if (size >= 32) smem[tid] = sum = sum + smem[tid + 16];
        if (size >= 16) smem[tid] = sum = sum + smem[tid + 8];
        if (size >= 8) smem[tid] = sum = sum + smem[tid + 4];
        if (size >= 4) smem[tid] = sum = sum + smem[tid + 2];
        if (size >= 2) smem[tid] = sum = sum + smem[tid + 1];
    }

    barrier(CLK_LOCAL_MEM_FENCE);
    sum = smem[0];

    return sum;
}

__kernel void normalize_hists_kernel(const int nthreads, const int block_hist_size, const int img_block_width,
                                     __global float* block_hists, const float threshold, __local float *squares)
{
    const int tid = get_local_id(0);
    const int gidX = get_group_id(0);
    const int gidY = get_group_id(1);

    __global float* hist = block_hists + (gidY * img_block_width + gidX) * block_hist_size + tid;

    float elem = 0.f;
    if (tid < block_hist_size)
        elem = hist[0];

    squares[tid] = elem * elem;

    barrier(CLK_LOCAL_MEM_FENCE);
    float sum = reduce_smem(squares, nthreads);

    float scale = 1.0f / (sqrt(sum) + 0.1f * block_hist_size);
    elem = min(elem * scale, threshold);

    barrier(CLK_LOCAL_MEM_FENCE);
    squares[tid] = elem * elem;

    barrier(CLK_LOCAL_MEM_FENCE);
    sum = reduce_smem(squares, nthreads);
    scale = 1.0f / (sqrt(sum) + 1e-3f);

    if (tid < block_hist_size)
        hist[0] = elem * scale;
}

//---------------------------------------------------------------------
//  Linear SVM based classification
//
__kernel void classify_hists_kernel(const int cblock_hist_size, const int cdescr_size, const int cdescr_width,
                                    const int img_win_width, const int img_block_width,
                                    const int win_block_stride_x, const int win_block_stride_y,
                                    __global const float * block_hists, __global const float* coefs,
                                    float free_coef, float threshold, __global uchar* labels)
{
    const int tid = get_local_id(0);
    const int gidX = get_group_id(0);
    const int gidY = get_group_id(1);

    __global const float* hist = block_hists + (gidY * win_block_stride_y * img_block_width + gidX * win_block_stride_x) * cblock_hist_size;

    float product = 0.f;
    for (int i = tid; i < cdescr_size; i += NTHREADS)
    {
        int offset_y = i / cdescr_width;
        int offset_x = i - offset_y * cdescr_width;
        product += coefs[i] * hist[offset_y * img_block_width * cblock_hist_size + offset_x];
    }

    __local float products[NTHREADS];

    products[tid] = product;

    barrier(CLK_LOCAL_MEM_FENCE);

    if (tid < 128) products[tid] = product = product + products[tid + 128];
    barrier(CLK_LOCAL_MEM_FENCE);

    if (tid < 64) products[tid] = product = product + products[tid + 64];
    barrier(CLK_LOCAL_MEM_FENCE);

    if (tid < 32)
    {
        volatile __local float* smem = products;
        smem[tid] = product = product + smem[tid + 32];
        smem[tid] = product = product + smem[tid + 16];
        smem[tid] = product = product + smem[tid + 8];
        smem[tid] = product = product + smem[tid + 4];
        smem[tid] = product = product + smem[tid + 2];
        smem[tid] = product = product + smem[tid + 1];
    }<commit_msg>Fix ocl::HOG crash on Intel OCL<commit_after>/*M///////////////////////////////////////////////////////////////////////////////////////
//
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
//  By downloading, copying, installing or using the software you agree to this license.
//  If you do not agree to this license, do not download, install,
//  copy or use the software.
//
//
//                           License Agreement
//                For Open Source Computer Vision Library
//
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
//    <NAME>, <EMAIL>
//
// 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.
//
//M*/


#define CELL_WIDTH 8
#define CELL_HEIGHT 8
#define CELLS_PER_BLOCK_X 2
#define CELLS_PER_BLOCK_Y 2
#define NTHREADS 256
#define CV_PI_F 3.1415926535897932384626433832795f

//----------------------------------------------------------------------------
// Histogram computation

__kernel void compute_hists_kernel(const int width, const int cblock_stride_x, const int cblock_stride_y,
                                   const int cnbins, const int cblock_hist_size, const int img_block_width,
                                   const int grad_quadstep, const int qangle_step,
                                   __global const float* grad, __global const uchar* qangle,
                                   const float scale, __global float* block_hists, __local float* smem)
{
    const int lidX = get_local_id(0);
    const int lidY = get_local_id(1);
    const int gidX = get_group_id(0);
    const int gidY = get_group_id(1);

    const int cell_x = lidX / 16;
    const int cell_y = lidY;
    const int cell_thread_x = lidX & 0xF;

    __local float* hists = smem;
    __local float* final_hist = smem + cnbins * 48;

    const int offset_x = gidX * cblock_stride_x + (cell_x << 2) + cell_thread_x;
    const int offset_y = gidY * cblock_stride_y + (cell_y << 2);

    __global const float* grad_ptr = grad + offset_y * grad_quadstep + (offset_x << 1);
    __global const uchar* qangle_ptr = qangle + offset_y * qangle_step + (offset_x << 1);

    // 12 means that 12 pixels affect on block's cell (in one row)
    if (cell_thread_x < 12)
    {
        __local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) + cell_thread_x;
        for (int bin_id = 0; bin_id < cnbins; ++bin_id)
            hist[bin_id * 48] = 0.f;

        const int dist_x = -4 + cell_thread_x - 4 * cell_x;

        const int dist_y_begin = -4 - 4 * lidY;
        for (int dist_y = dist_y_begin; dist_y < dist_y_begin + 12; ++dist_y)
        {
            float2 vote = (float2) (grad_ptr[0], grad_ptr[1]);
            uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]);

            grad_ptr += grad_quadstep;
            qangle_ptr += qangle_step;

            int dist_center_y = dist_y - 4 * (1 - 2 * cell_y);
            int dist_center_x = dist_x - 4 * (1 - 2 * cell_x);

            float gaussian = exp(-(dist_center_y * dist_center_y + dist_center_x * dist_center_x) * scale);
            float interp_weight = (8.f - fabs(dist_y + 0.5f)) * (8.f - fabs(dist_x + 0.5f)) / 64.f;

            hist[bin.x * 48] += gaussian * interp_weight * vote.x;
            hist[bin.y * 48] += gaussian * interp_weight * vote.y;
        }

        volatile __local float* hist_ = hist;
        for (int bin_id = 0; bin_id < cnbins; ++bin_id, hist_ += 48)
        {
            if (cell_thread_x < 6) hist_[0] += hist_[6];
            if (cell_thread_x < 3) hist_[0] += hist_[3];
            if (cell_thread_x == 0)
                final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = hist_[0] + hist_[1] + hist_[2];
        }
    }

    barrier(CLK_LOCAL_MEM_FENCE);

    __global float* block_hist = block_hists + (gidY * img_block_width + gidX) * cblock_hist_size;

    int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 16 + cell_thread_x;
    if (tid < cblock_hist_size)
        block_hist[tid] = final_hist[tid];
}

//-------------------------------------------------------------
//  Normalization of histograms via L2Hys_norm
//
float reduce_smem(volatile __local float* smem, int size)
{
    unsigned int tid = get_local_id(0);
    float sum = smem[tid];

    if (size >= 512) { if (tid < 256) smem[tid] = sum = sum + smem[tid + 256]; barrier(CLK_LOCAL_MEM_FENCE); }
    if (size >= 256) { if (tid < 128) smem[tid] = sum = sum + smem[tid + 128]; barrier(CLK_LOCAL_MEM_FENCE); }
    if (size >= 128) { if (tid < 64) smem[tid] = sum = sum + smem[tid + 64]; barrier(CLK_LOCAL_MEM_FENCE); }

    if (tid < 32)
    {
        if (size >= 64) smem[tid] = sum = sum + smem[tid + 32];
    }
    barrier(CLK_LOCAL_MEM_FENCE);
    if (tid < 16)
    {
        if (size >= 32) smem[tid] = sum = sum + smem[tid + 16];
        if (size >= 16) smem[tid] = sum = sum + smem[tid + 8];
        if (size >= 8) smem[tid] = sum = sum + smem[tid + 4];
        if (size >= 4) smem[tid] = sum = sum + smem[tid + 2];
        if (size >= 2) smem[tid] = sum = sum + smem[tid + 1];
    }

    barrier(CLK_LOCAL_MEM_FENCE);
    sum = smem[0];

    return sum;
}

__kernel void normalize_hists_kernel(const int nthreads, const int block_hist_size, const int img_block_width,
                                     __global float* block_hists, const float threshold, __local float *squares)
{
    const int tid = get_local_id(0);
    const int gidX = get_group_id(0);
    const int gidY = get_group_id(1);

    __global float* hist = block_hists + (gidY * img_block_width + gidX) * block_hist_size + tid;

    float elem = 0.f;
    if (tid < block_hist_size)
        elem = hist[0];

    squares[tid] = elem * elem;

    barrier(CLK_LOCAL_MEM_FENCE);
    float sum = reduce_smem(squares, nthreads);

    float scale = 1.0f / (sqrt(sum) + 0.1f * block_hist_size);
    elem = min(elem * scale, threshold);

    barrier(CLK_LOCAL_MEM_FENCE);
    squares[tid] = elem * elem;

    barrier(CLK_LOCAL_MEM_FENCE);
    sum = reduce_smem(squares, nthreads);
    scale = 1.0f / (sqrt(sum) + 1e-3f);

    if (tid < block_hist_size)
        hist[0] = elem * scale;
}

//---------------------------------------------------------------------
//  Linear SVM based classification
//
__kernel void classify_hists_kernel(const int cblock_hist_size, const int cdescr_size, const int cdescr_width,
                                    const int img_win_width, const int img_block_width,
                                    const int win_block_stride_x, const int win_block_stride_y,
                                    __global const float * block_hists, __global const float* coefs,
                                    float free_coef, float threshold, __global uchar* labels)
{
    const int tid = get_local_id(0);
    const int gidX = get_group_id(0);
    const int gidY = get_group_id(1);

    __global const float* hist = block_hists + (gidY * win_block_stride_y * img_block_width + gidX * win_block_stride_x) * cblock_hist_size;

    float product = 0.f;
    for (int i = tid; i < cdescr_size; i += NTHREADS)
    {
        int offset_y = i / cdescr_width;
        int offset_x = i - offset_y * cdescr_width;
        product += coefs[i] * hist[offset_y * img_block_width * cblock_hist_size + offset_x];
    }

    __local float products[NTHREADS];

    products[tid] = product;

    barrier(CLK_LOCAL_MEM_FENCE);

    if (tid < 128) products[tid] = product = product + products[tid + 128];
    barrier(CLK_LOCAL_MEM_FENCE);

    if (tid < 64) products[tid] = product = product + products[tid + 64];
    barrier(CLK_LOCAL_MEM_FENCE);

    if (tid < 32)
    {
        volatile __local float* smem = products;
        smem[tid] = product = product + smem[tid + 32];
    }
    barrier(CLK_LOCAL_MEM_FENCE);
    if (tid < 16)
    {
        volatile __local float* smem = products;
        smem[tid] = product = product + smem[tid + 16];
        smem[tid] = product = product + smem[tid + 8];
        smem[tid] = product = product + smem[tid + 4];
        smem[tid] = product = product + smem[tid + 2];
        smem[tid] = product = product + smem[tid + 1];
    }
{'fraction_non_alphanumeric': 0.13266000313988174, 'fraction_numerical': 0.022502485739703804, 'mean_word_length': 2.3165567511280805, 'pattern_counts': {'":': 0, '<': 59, '<?xml version=': 0, '>': 25, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 22, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '6527395', 'n_tokens_mistral': 6487, 'n_tokens_neox': 6021, 'n_words': 1888}