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 //    Jin Ma jin (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 #if kercn != 3
     47 #define storepix(val, addr)  *(__global T *)(addr) = val
     48 #define TSIZE (int)sizeof(T)
     49 #define scalar scalar_
     50 #else
     51 #define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr))
     52 #define TSIZE ((int)sizeof(T1)*3)
     53 #define scalar (T)(scalar_.x, scalar_.y, scalar_.z)
     54 #endif
     55 
     56 __kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset, int rows, int cols,
     57                           ST scalar_)
     58 {
     59     int x = get_global_id(0);
     60     int y0 = get_global_id(1) * rowsPerWI;
     61 
     62     if (x < cols)
     63     {
     64         int src_index = mad24(y0, src_step, mad24(x, TSIZE, src_offset));
     65 
     66 #if kercn == cn
     67         #pragma unroll
     68         for (int y = y0, i = 0, y1 = min(rows, y0 + rowsPerWI); i < rowsPerWI; ++y, ++i, src_index += src_step)
     69             if (y < y1)
     70                 storepix(x == y ? scalar : (T)(0), srcptr + src_index);
     71 #elif kercn == 4 && cn == 1
     72         if (y0 < rows)
     73         {
     74             storepix(x == y0 >> 2 ? (T)(scalar, 0, 0, 0) : (T)(0), srcptr + src_index);
     75             if (++y0 < rows)
     76             {
     77                 src_index += src_step;
     78                 storepix(x == y0 >> 2 ? (T)(0, scalar, 0, 0) : (T)(0), srcptr + src_index);
     79 
     80                 if (++y0 < rows)
     81                 {
     82                     src_index += src_step;
     83                     storepix(x == y0 >> 2 ? (T)(0, 0, scalar, 0) : (T)(0), srcptr + src_index);
     84 
     85                     if (++y0 < rows)
     86                     {
     87                         src_index += src_step;
     88                         storepix(x == y0 >> 2 ? (T)(0, 0, 0, scalar) : (T)(0), srcptr + src_index);
     89                     }
     90                 }
     91             }
     92         }
     93 #else
     94 #error "Incorrect combination of cn && kercn"
     95 #endif
     96     }
     97 }
     98