modules/ocl/src/opencl/arithm_log.cl

   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, Institute Of Software Chinese Academy Of Science, 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 //    Wu Zailong, bullet@yeah.net
  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 //////////////////////////////////////////////////////////////////////////////////////////////////////
  55 /////////////////////////////////////////////LOG/////////////////////////////////////////////////////
  56 ///////////////////////////////////////////////////////////////////////////////////////////////////////
  57
  58 __kernel void arithm_log_C1(__global srcT *src, __global srcT *dst,
  59     int cols1, int rows,
  60     int srcOffset1, int dstOffset1,
  61     int srcStep1, int dstStep1)
  62 {
  63     int x = get_global_id(0);
  64     int y = get_global_id(1);
  65
  66     if(x < cols1 && y < rows)
  67     {
  68         int srcIdx = mad24(y, srcStep1, x + srcOffset1);
  69         int dstIdx = mad24(y, dstStep1, x + dstOffset1);
  70
  71         dst[dstIdx] = log(src[srcIdx]);
  72     }
  73 }
  74
  75 __kernel void arithm_log_C2(__global srcT *src, __global srcT *dst,
  76     int cols1, int rows,
  77     int srcOffset1, int dstOffset1,
  78     int srcStep1, int dstStep1)
  79 {
  80     int x1 = get_global_id(0) << 1;
  81     int y = get_global_id(1);
  82
  83     if(x1 < cols1 && y < rows)
  84     {
  85         int srcIdx = mad24(y, srcStep1, x1 + srcOffset1);
  86         int dstIdx = mad24(y, dstStep1, x1 + dstOffset1);
  87
  88         dst[dstIdx] =                      log(src[srcIdx]);
  89         dst[dstIdx + 1] = x1 + 1 < cols1 ? log(src[srcIdx + 1]) : dst[dstIdx + 1];
  90     }
  91 }
  92
  93 __kernel void arithm_log_C4(__global srcT *src, __global srcT *dst,
  94     int cols1, int rows,
  95     int srcOffset1, int dstOffset1,
  96     int srcStep1, int dstStep1)
  97 {
  98     int x1 = get_global_id(0) << 2;
  99     int y = get_global_id(1);
 100
 101     if(x1 < cols1 && y < rows)
 102     {
 103         int srcIdx = mad24(y, srcStep1, x1 + srcOffset1);
 104         int dstIdx = mad24(y, dstStep1, x1 + dstOffset1);
 105
 106         dst[dstIdx] =                      log(src[srcIdx]);
 107         dst[dstIdx + 1] = x1 + 1 < cols1 ? log(src[srcIdx + 1]) : dst[dstIdx + 1];
 108         dst[dstIdx + 2] = x1 + 2 < cols1 ? log(src[srcIdx + 2]) : dst[dstIdx + 2];
 109         dst[dstIdx + 3] = x1 + 3 < cols1 ? log(src[srcIdx + 3]) : dst[dstIdx + 3];
 110     }
 111 }