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

   1 // Copyright (c) 2019 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 KERNEL(lstm_dynamic_input_ref)(
  19     const __global INPUT0_TYPE* input,
  20     const __global DYN_LENGTH_TYPE* dyn_lengths,
  21     __global OUTPUT_TYPE* output,
  22     const __global WEIGHTS_TYPE* weights
  23 #if BIAS_TERM
  24     , const __global BIAS_TYPE* biases
  25 #endif
  26     )
  27 {
  28     const uint y        = get_global_id(0);
  29     const uint batch    = get_global_id(1) % INPUT0_BATCH_NUM;
  30     const uint dir      = get_global_id(1) / INPUT0_BATCH_NUM;
  31     const uint timestep = get_global_id(2);
  32
  33     if(timestep > (uint)dyn_lengths[batch])
  34         return;
  35
  36     ACCUMULATOR_TYPE dot_prod = 0;
  37     for(uint x = 0; x < INPUT0_SIZE_X; ++x )
  38     {
  39         const uint input_idx   = GET_DATA_INDEX(INPUT0, batch, timestep, dir, x);
  40         const uint weights_idx = GET_FILTER_INDEX(WEIGHTS, 0, dir, y, x);
  41         dot_prod += (ACCUMULATOR_TYPE)(input[input_idx] * weights[weights_idx]);
  42     }
  43
  44 #if BIAS_TERM
  45     dot_prod += (ACCUMULATOR_TYPE)biases[GET_DATA_INDEX(BIAS, 0, 0, dir, y)];
  46 #endif
  47
  48     output[GET_DATA_INDEX(OUTPUT, batch, timestep, dir, y)] = (OUTPUT_TYPE)dot_prod;
  49 }