modules/dnn/src/opencl/eltwise.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 __kernel void op_sum4(__global const Dtype * A,
  48                       __global const Dtype * B,
  49                       unsigned int A_col_size,
  50                       const float coeff1,
  51                       const float coeff2,
  52                       __global Dtype * C)
  53 {
  54     unsigned int row_gid = get_group_id(0);
  55     unsigned int lid = get_local_id(0);
  56     const __global Dtype *src0_read = A + row_gid * 4 * A_col_size;
  57     const __global Dtype *src1_read = B + row_gid * 4 * A_col_size;
  58     __global Dtype *dst0_read = C + row_gid * 4 * A_col_size;
  59
  60     Dtype4 a0, a1, a2, a3;
  61     Dtype4 dot0, dot1, dot2, dot3;
  62     unsigned int i = lid;
  63     while( i < A_col_size / 4)
  64     {
  65         const Dtype4 b0 = vload4(i, src1_read);
  66         const Dtype4 b1 = vload4(i, src1_read + A_col_size);
  67         const Dtype4 b2 = vload4(i, src1_read + 2 * A_col_size);
  68         const Dtype4 b3 = vload4(i, src1_read + 3 * A_col_size);
  69
  70 #if LOOP == 0
  71         a0 = vload4(i, src0_read);
  72         a1 = vload4(i, src0_read + A_col_size);
  73         a2 = vload4(i, src0_read + 2 * A_col_size);
  74         a3 = vload4(i, src0_read + 3 * A_col_size);
  75
  76         dot0 = a0 * coeff1 + b0 * coeff2;
  77         dot1 = a1 * coeff1 + b1 * coeff2;
  78         dot2 = a2 * coeff1 + b2 * coeff2;
  79         dot3 = a3 * coeff1 + b3 * coeff2;
  80 #else
  81         a0 = vload4(i, dst0_read);
  82         a1 = vload4(i, dst0_read + A_col_size);
  83         a2 = vload4(i, dst0_read + 2 * A_col_size);
  84         a3 = vload4(i, dst0_read + 3 * A_col_size);
  85
  86         dot0 = a0 + b0 * coeff2;
  87         dot1 = a1 + b1 * coeff2;
  88         dot2 = a2 + b2 * coeff2;
  89         dot3 = a3 + b3 * coeff2;
  90 #endif
  91         vstore4(dot0, i, dst0_read);
  92         vstore4(dot1, i, dst0_read + A_col_size);
  93         vstore4(dot2, i, dst0_read + 2 * A_col_size);
  94         vstore4(dot3, i, dst0_read + 3 * A_col_size);
  95
  96         i += get_local_size(0);
  97     }
  98 }