inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/generic_eltwise_ref.cl

   1 /*
   2 // Copyright (c) 2016 Intel Corporation
   3 //
   4 // Licensed under the Apache License, Version 2.0 (the "License");
   5 // you may not use this file except in compliance with the License.
   6 // You may obtain a copy of the License at
   7 //
   8 //      http://www.apache.org/licenses/LICENSE-2.0
   9 //
  10 // Unless required by applicable law or agreed to in writing, software
  11 // distributed under the License is distributed on an "AS IS" BASIS,
  12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  13 // See the License for the specific language governing permissions and
  14 // limitations under the License.
  15 */
  16
  17 #include "include/include_all.cl"
  18
  19 #ifdef INPUT_STRIDED
  20
  21 #define GET_INDEX(prefix, num) \
  22     CAT(CAT(prefix, num), _OFFSET) + \
  23     ((d1 * CAT(CAT(prefix, num), _STRIDE_X)) % CAT(CAT(prefix, num), _SIZE_X))*CAT(CAT(prefix, num), _X_PITCH) +\
  24     ((d2 * CAT(CAT(prefix, num), _STRIDE_Y)) % CAT(CAT(prefix, num), _SIZE_Y))*CAT(CAT(prefix, num), _Y_PITCH) +\
  25     (d3 % CAT(CAT(prefix, num), _FEATURE_NUM))*CAT(CAT(prefix, num), _FEATURE_PITCH) + \
  26     (d4 % CAT(CAT(prefix, num), _BATCH_NUM  ))*CAT(CAT(prefix, num), _BATCH_PITCH)
  27
  28 #else
  29
  30 #if ELTWISE_LAYOUT_BASED || QUANTIZATION_TERM || ELTWISE_BROADCAST
  31
  32 #define GET_INDEX(prefix, num)                                                          \
  33     CAT(CAT(prefix, num), _OFFSET) +                                                    \
  34     (d1 % CAT(CAT(prefix, num), _SIZE_X     ))*CAT(CAT(prefix, num), _X_PITCH) +        \
  35     (d2 % CAT(CAT(prefix, num), _SIZE_Y     ))*CAT(CAT(prefix, num), _Y_PITCH) +        \
  36     (d3 % CAT(CAT(prefix, num), _FEATURE_NUM))*CAT(CAT(prefix, num), _FEATURE_PITCH) +  \
  37     (d4 % CAT(CAT(prefix, num), _BATCH_NUM  ))*CAT(CAT(prefix, num), _BATCH_PITCH)
  38
  39 #elif ELTWISE_NO_PITCH_SAME_DIMS
  40 #define GET_INDEX(prefix, num)                                                      \
  41     CAT(CAT(prefix, num), _OFFSET) + d1
  42
  43 #else
  44
  45 #define GET_INDEX(prefix, num)                                                      \
  46     CAT(CAT(prefix, num), _OFFSET) +                                                \
  47     (d1 % CAT(CAT(prefix, num), _SIZES)[0])*CAT(CAT(prefix, num), _PITCHES)[0] +    \
  48     (d2 % CAT(CAT(prefix, num), _SIZES)[1])*CAT(CAT(prefix, num), _PITCHES)[1] +    \
  49     (d3 % CAT(CAT(prefix, num), _SIZES)[2])*CAT(CAT(prefix, num), _PITCHES)[2] +    \
  50     (d4 % CAT(CAT(prefix, num), _SIZES)[3])*CAT(CAT(prefix, num), _PITCHES)[3]
  51
  52 #endif
  53
  54 #endif
  55
  56
  57 KERNEL(eltwise)(
  58     INPUTS_DECLS
  59     __global UNIT_TYPE* output
  60 #if CALIBRATION_TERM
  61     , const __global float* calibrations
  62 #endif
  63     )
  64 {
  65 #if ELTWISE_LAYOUT_BASED || QUANTIZATION_TERM || ELTWISE_BROADCAST
  66     const uint d1 = get_global_id(GWS_YX) % OUTPUT_SIZE_X;  // X
  67     const uint d2 = get_global_id(GWS_YX) / OUTPUT_SIZE_X;  // Y
  68     const uint d3 = get_global_id(GWS_FEATURE);             // Feature
  69     const uint d4 = get_global_id(GWS_BATCH);               // Batch
  70
  71     uint output_offset = OUTPUT_OFFSET +
  72                          d1*OUTPUT_X_PITCH +
  73                          d2*OUTPUT_Y_PITCH +
  74                          d3*OUTPUT_FEATURE_PITCH +
  75                          d4*OUTPUT_BATCH_PITCH;
  76 #elif ELTWISE_NO_PITCH_SAME_DIMS
  77     const uint d1 = get_global_id(0);
  78     uint output_offset = OUTPUT_OFFSET + d1;
  79 #else
  80     const uint d1 = get_global_id(0);
  81     const uint d2 = get_global_id(1);
  82     const uint d3 = get_global_id(2) % OUTPUT_SIZES[2];
  83     const uint d4 = get_global_id(2) / OUTPUT_SIZES[2];
  84
  85     uint output_offset = OUTPUT_OFFSET +
  86                          d1*OUTPUT_PITCHES[0] +
  87                          d2*OUTPUT_PITCHES[1] +
  88                          d3*OUTPUT_PITCHES[2] +
  89                          d4*OUTPUT_PITCHES[3];
  90 #endif
  91
  92 #if QUANTIZATION_TERM
  93     int res;
  94 #else
  95     UNIT_TYPE res;
  96 #endif
  97
  98     DO_ELTWISE;
  99
 100 #if QUANTIZATION_TERM
 101 #if CALIBRATION_TERM
 102     res = (int)round(((float)res) * calibrations[d3]);
 103 #else  // CALIBRATION_TERM
 104     res = (int)round(((float)res) * O_QF);
 105 #endif // CALIBRATION_TERM
 106 #endif // QUANTIZATION_TERM
 107
 108 #if QUANTIZATION_TERM
 109     output[output_offset] = ACTIVATION(convert_char(res), NL_M, NL_N);
 110 #else
 111     output[output_offset] = ACTIVATION(res, NL_M, NL_N);
 112 #endif
 113 }