inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/strided_slice_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(strided_slice_ref)(const __global UNIT_TYPE* input, __global UNIT_TYPE* output)
  19 {
  20     const uint batch = get_global_id(0);
  21     const uint feature = get_global_id(1);
  22
  23 #if NEW_AXIS_MODE
  24     // If NEW_AXIS_MODE that just copy input to output
  25     const uint y_input = get_global_id(2) / INPUT0_SIZE_X;
  26     const uint x_input = get_global_id(2) % INPUT0_SIZE_X;
  27     const uint input_index = INPUT0_OFFSET +
  28         batch * INPUT0_BATCH_PITCH +
  29         feature * INPUT0_FEATURE_PITCH +
  30         y_input * INPUT0_Y_PITCH +
  31         x_input * INPUT0_X_PITCH;
  32     output[input_index] = input[input_index];
  33 #else
  34     const uint y = get_global_id(2) / OUTPUT_SIZE_X;
  35     const uint x = get_global_id(2) % OUTPUT_SIZE_X;
  36     const uint input_index = INPUT0_OFFSET +
  37             (SLICE_BEGIN_BATCH + batch * SLICE_STEPS_BATCH) * INPUT0_BATCH_PITCH +
  38             (SLICE_BEGIN_FEATURE + feature * SLICE_STEPS_FEATURE) * INPUT0_FEATURE_PITCH +
  39             (SLICE_BEGIN_Y + y * SLICE_STEPS_Y) * INPUT0_Y_PITCH +
  40             (SLICE_BEGIN_X + x * SLICE_STEPS_X) * INPUT0_X_PITCH;
  41
  42     const uint output_index = OUTPUT_OFFSET +
  43             batch * OUTPUT_BATCH_PITCH +
  44             feature * OUTPUT_FEATURE_PITCH +
  45             y * OUTPUT_Y_PITCH +
  46             x * OUTPUT_X_PITCH;
  47
  48     output[output_index] = input[input_index];
  49 #endif
  50 }