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

   1 // Copyright (c) 2016-2017 Intel Corporation
   2 //
   3 // Licensed under the Apache License, Version 2.0 (the "License");
   4 // you may not use this file except in compliance with the License.
   5 // You may obtain a copy of the License at
   6 //
   7 //      http://www.apache.org/licenses/LICENSE-2.0
   8 //
   9 // Unless required by applicable law or agreed to in writing, software
  10 // distributed under the License is distributed on an "AS IS" BASIS,
  11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  12 // See the License for the specific language governing permissions and
  13 // limitations under the License.
  14
  15
  16 #include "include/include_all.cl"
  17
  18 #define ACTIVATION_LOGISTIC(input)                      (UNIT_VAL_ONE/(UNIT_VAL_ONE + exp(-input)))
  19 #define ACTIVATION_HYPERBOLIC_TAN(input)                (tanh(input))
  20
  21 // tempGEMM = [ batch, 1, direction, 4 * hidden_size ]
  22 // cell     = [ batch, 1, direction, hidden_size ] optional
  23 // output   = [ batch, 1, direction, hidden_size ] output
  24 KERNEL(lstm_elt)(
  25     const __global INPUT0_TYPE* input,
  26     __global OUTPUT_TYPE* output
  27 #if CELL_TERM
  28     ,const __global OUTPUT_TYPE* cell
  29 #endif
  30     )
  31 {
  32     const uint x = get_global_id(0);
  33     const uint b = get_global_id(1);
  34
  35     ACCUMULATOR_TYPE it = input[GET_DATA_INDEX(INPUT0, b, 0, 0, x + GEMM_OFFSET_I)];
  36     ACCUMULATOR_TYPE ot = input[GET_DATA_INDEX(INPUT0, b, 0, 0, x + GEMM_OFFSET_O)]; // pass constant offsets here
  37     ACCUMULATOR_TYPE zt = input[GET_DATA_INDEX(INPUT0, b, 0, 0, x + GEMM_OFFSET_Z)];
  38
  39     ACCUMULATOR_TYPE val = ACTIVATION_LOGISTIC(CLIP(it)) * ACTIVATION_HYPERBOLIC_TAN(CLIP(zt));
  40
  41 #if CELL_TERM || INPUT_FORGET
  42     ACCUMULATOR_TYPE ft = input[GET_DATA_INDEX(INPUT0, b, 0, 0, x + GEMM_OFFSET_F)];
  43 #endif
  44
  45 #if INPUT_FORGET
  46     val *= ((ACCUMULATOR_TYPE)1 - ft);
  47 #endif
  48
  49 #if CELL_TERM
  50     val += cell[GET_DATA_INDEX(CELL, b, 0, CELL_DIRECTION, x)] * ACTIVATION_LOGISTIC(CLIP(ft));
  51 #endif
  52
  53     output[GET_DATA_INDEX(OUTPUT, b, 0, 0, x)] = (OUTPUT_TYPE)(ACTIVATION_HYPERBOLIC_TAN(val) * ACTIVATION_LOGISTIC(ot)); // hidden
  54     output[GET_DATA_INDEX(OUTPUT, b, 1, 0, x)] = (OUTPUT_TYPE)val; // cell
  55 }