diff --git a/src/common/transformations/src/transformations/op_conversions/convert_sequences_to_tensor_iterator.cpp b/src/common/transformations/src/transformations/op_conversions/convert_sequences_to_tensor_iterator.cpp index bf3a7d13a2c5fb..b12cc401206e02 100644 --- a/src/common/transformations/src/transformations/op_conversions/convert_sequences_to_tensor_iterator.cpp +++ b/src/common/transformations/src/transformations/op_conversions/convert_sequences_to_tensor_iterator.cpp @@ -468,7 +468,7 @@ ov::pass::ConvertLSTMSequenceToTensorIterator::ConvertLSTMSequenceToTensorIterat } ov::pass::ConvertSequenceToTensorIterator::ConvertSequenceToTensorIterator() { - //d_matcher(); + //add_matcher(); add_matcher(); add_matcher(); } diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_seq.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_seq.cpp index 01621bd13856a6..d1482274ac8ca2 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_seq.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_seq.cpp @@ -32,8 +32,8 @@ struct lstm_seq_impl : typed_primitive_impl_ocl { for (size_t i = 0; i < instance.outputs_memory_count(); i++) { args.outputs.push_back(instance.output_memory_ptr(i)); } - args.outputs.push_back(instance.second_output_mem()); - args.outputs.push_back(instance.third_output_mem()); + //args.outputs.push_back(instance.second_output_mem()); + //args.outputs.push_back(instance.third_output_mem()); return args; } diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/lstm_seq_gpu_bfyx_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/lstm_seq_gpu_bfyx_ref.cl index 91b731aa3e58db..cad293a3f04a62 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/lstm_seq_gpu_bfyx_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/lstm_seq_gpu_bfyx_ref.cl @@ -1,17 +1,9 @@ -// Copyright (C) 2018-2024 Intel Corporation +// Copyright (C) 2024 Intel Corporation // SPDX-License-Identifier: Apache-2.0 // #include "include/batch_headers/fetch_data.cl" -// initial_hidden_state -// initial_cell_state -// sequence_lengths -// WR -// B -// output0 -//output1 -//output2 KERNEL(lstm_seq)( const __global INPUT0_TYPE* x, const __global INPUT1_TYPE* initial_hidden_state, @@ -25,8 +17,39 @@ KERNEL(lstm_seq)( __global OUTPUT2_TYPE* cell_state ) { + const uint hidden_idx = get_global_id(0); const uint b = get_global_id(1); + for(int i=0;i