modules/dnn/src/opencl/mvn.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) 2017, Intel Corporation, all rights reserved.
  14 // Copyright (c) 2016-2017 Fabian David Tschopp, all rights reserved.
  15 // Third party copyrights are property of their respective owners.
  16 //
  17 // Redistribution and use in source and binary forms, with or without modification,
  18 // are permitted provided that the following conditions are met:
  19 //
  20 //   * Redistribution's of source code must retain the above copyright notice,
  21 //     this list of conditions and the following disclaimer.
  22 //
  23 //   * Redistribution's in binary form must reproduce the above copyright notice,
  24 //     this list of conditions and the following disclaimer in the documentation
  25 //     and/or other materials provided with the distribution.
  26 //
  27 //   * The name of the copyright holders may not be used to endorse or promote products
  28 //     derived from this software without specific prior written permission.
  29 //
  30 // This software is provided by the copyright holders and contributors "as is" and
  31 // any express or implied warranties, including, but not limited to, the implied
  32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
  33 // In no event shall the Intel Corporation or contributors be liable for any direct,
  34 // indirect, incidental, special, exemplary, or consequential damages
  35 // (including, but not limited to, procurement of substitute goods or services;
  36 // loss of use, data, or profits; or business interruption) however caused
  37 // and on any theory of liability, whether in contract, strict liability,
  38 // or tort (including negligence or otherwise) arising in any way out of
  39 // the use of this software, even if advised of the possibility of such damage.
  40 //
  41 //M*/
  42
  43 #define Dtype float
  44 #define Dtype4 float4
  45 #define Dtype8 float8
  46
  47 #if NUM == 8
  48     #define load(src, index) vload8(0, src + index)
  49     #define store(vec, dst, index) vstore8(vec, 0, dst + index)
  50     #define vec_type Dtype8
  51     #define CALC_MEAN calc_mean8
  52     #define MVN mvn8
  53 #elif NUM == 4
  54     #define load(src, index) vload4(0, src + index)
  55     #define store(vec, dst, index) vstore4(vec, 0, dst + index)
  56     #define vec_type Dtype4
  57     #define CALC_MEAN calc_mean4
  58     #define MVN mvn4
  59 #elif NUM == 1
  60     #define load(src, index) src[index]
  61     #define store(vec, dst, index) dst[index] = vec
  62     #define vec_type Dtype
  63     #define CALC_MEAN calc_mean1
  64     #define MVN mvn1
  65 #endif
  66
  67 __kernel void CALC_MEAN(__global const Dtype* src,
  68                         const int rows,
  69                         const int cols,
  70                         __global Dtype* mean,
  71                         __global Dtype* dst)
  72 {
  73     int x = get_global_id(0);
  74     int y = get_global_id(1) * NUM;
  75     int index = x * cols + y;
  76
  77     if (x >= rows || y >= cols)
  78         return;
  79
  80     Dtype mean_val = mean[x];
  81     vec_type src_vec = load(src, index);
  82     vec_type dst_vec = native_powr(src_vec - (vec_type)mean_val, 2);
  83     store(dst_vec, dst, index);
  84 }
  85
  86 __kernel void MVN(__global const Dtype* src,
  87                   const int rows,
  88                   const int cols,
  89                   const Dtype eps,
  90                   __global const Dtype* mean,
  91                   __global const Dtype* dev,
  92                   __global Dtype* dst)
  93 {
  94     int x = get_global_id(0);
  95     int y = get_global_id(1) * NUM;
  96     int index = x * cols + y;
  97
  98     if (x >= rows || y >= cols)
  99         return;
 100
 101     Dtype mean_val = mean[x];
 102     Dtype dev_val = sqrt(dev[x]);
 103     Dtype alpha;
 104 #ifdef NORM_VARIANCE
 105     alpha = 1 / (eps + dev_val);
 106 #else
 107     alpha = 1;
 108 #endif
 109     vec_type src_vec = load(src, index) - (vec_type)mean_val;
 110     vec_type dst_vec = src_vec * alpha;
 111     store(dst_vec, dst, index);
 112 }