1 // Copyright (c) 2019 Intel Corporation
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
7 // http://www.apache.org/licenses/LICENSE-2.0
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.
16 #include "include/include_all.cl"
18 KERNEL(reverse_sequence_ref)(const __global UNIT_TYPE* input, const __global float* seq_lengths, __global UNIT_TYPE* output)
20 const uint batch = get_global_id(0);
21 const uint feature = get_global_id(1);
22 const uint y = get_global_id(2) / INPUT0_SIZE_X;
23 const uint x = get_global_id(2) % INPUT0_SIZE_X;
24 uint dimensions[] = { batch, feature, y, x };
26 const uint input_index = INPUT0_OFFSET +
27 batch * INPUT0_BATCH_PITCH +
28 feature * INPUT0_FEATURE_PITCH +
32 const uint length = seq_lengths[dimensions[BATCH_AXIS]];
33 if (dimensions[SEQ_AXIS] < length)
34 dimensions[SEQ_AXIS] = length - dimensions[SEQ_AXIS] - 1;
36 const uint output_index = OUTPUT_OFFSET +
37 dimensions[0] * OUTPUT_BATCH_PITCH +
38 dimensions[1] * OUTPUT_FEATURE_PITCH +
39 dimensions[2] * OUTPUT_Y_PITCH +
40 dimensions[3] * OUTPUT_X_PITCH;
42 output[output_index] = input[input_index];