Home | History | Annotate | Download | only in opencl
      1 /*M///////////////////////////////////////////////////////////////////////////////////////
      2 //
      3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
      4 //
      5 //  By downloading, copying, installing or using the software you agree to this license.
      6 //  If you do not agree to this license, do not download, install,
      7 //  copy or use the software.
      8 //
      9 //
     10 //                           License Agreement
     11 //                For Open Source Computer Vision Library
     12 //
     13 // Copyright (C) 2010-2012, MulticoreWare Inc., all rights reserved.
     14 // Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
     15 // Third party copyrights are property of their respective owners.
     16 //
     17 // @Authors
     18 //    Liu Liujun, liujun (at) multicorewareinc.com
     19 //
     20 // Redistribution and use in source and binary forms, with or without modification,
     21 // are permitted provided that the following conditions are met:
     22 //
     23 //   * Redistribution's of source code must retain the above copyright notice,
     24 //     this list of conditions and the following disclaimer.
     25 //
     26 //   * Redistribution's in binary form must reproduce the above copyright notice,
     27 //     this list of conditions and the following disclaimer in the documentation
     28 //     and/or other materials provided with the distribution.
     29 //
     30 //   * The name of the copyright holders may not be used to endorse or promote products
     31 //     derived from this software without specific prior written permission.
     32 //
     33 // This software is provided by the copyright holders and contributors as is and
     34 // any express or implied warranties, including, but not limited to, the implied
     35 // warranties of merchantability and fitness for a particular purpose are disclaimed.
     36 // In no event shall the Intel Corporation or contributors be liable for any direct,
     37 // indirect, incidental, special, exemplary, or consequential damages
     38 // (including, but not limited to, procurement of substitute goods or services;
     39 // loss of use, data, or profits; or business interruption) however caused
     40 // and on any theory of liability, whether in contract, strict liability,
     41 // or tort (including negligence or otherwise) arising in any way out of
     42 // the use of this software, even if advised of the possibility of such damage.
     43 //
     44 //M*/
     45 
     46 #ifdef DOUBLE_SUPPORT
     47 #ifdef cl_amd_fp64
     48 #pragma OPENCL EXTENSION cl_amd_fp64:enable
     49 #elif defined (cl_khr_fp64)
     50 #pragma OPENCL EXTENSION cl_khr_fp64:enable
     51 #endif
     52 #endif
     53 
     54 #define noconvert
     55 
     56 __kernel void blendLinear(__global const uchar * src1ptr, int src1_step, int src1_offset,
     57                           __global const uchar * src2ptr, int src2_step, int src2_offset,
     58                           __global const uchar * weight1, int weight1_step, int weight1_offset,
     59                           __global const uchar * weight2, int weight2_step, int weight2_offset,
     60                           __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols)
     61 {
     62     int x = get_global_id(0);
     63     int y = get_global_id(1);
     64 
     65     if (x < dst_cols && y < dst_rows)
     66     {
     67         int src1_index = mad24(y, src1_step, src1_offset + x * cn * (int)sizeof(T));
     68         int src2_index = mad24(y, src2_step, src2_offset + x * cn * (int)sizeof(T));
     69         int weight1_index = mad24(y, weight1_step, weight1_offset + x * (int)sizeof(float));
     70         int weight2_index = mad24(y, weight2_step, weight2_offset + x * (int)sizeof(float));
     71         int dst_index = mad24(y, dst_step, dst_offset + x * cn * (int)sizeof(T));
     72 
     73         float w1 = *(__global const float *)(weight1 + weight1_index),
     74               w2 = *(__global const float *)(weight2 + weight2_index);
     75         float den = w1 + w2 + 1e-5f;
     76 
     77         __global const T * src1 = (__global const T *)(src1ptr + src1_index);
     78         __global const T * src2 = (__global const T *)(src2ptr + src2_index);
     79         __global T * dst = (__global T *)(dstptr + dst_index);
     80 
     81         #pragma unroll
     82         for (int i = 0; i < cn; ++i)
     83         {
     84             float num = w1 * convert_float(src1[i]) + w2 * convert_float(src2[i]);
     85             dst[i] = convertToT(num / den);
     86         }
     87     }
     88 }
     89