diff --git a/docs/template_plugin/tests/functional/shared_tests_instances/single_layer_tests/softmax.cpp b/docs/template_plugin/tests/functional/shared_tests_instances/single_layer_tests/softmax.cpp index 2950d6bc02a3b3..0d6efc122a3db7 100644 --- a/docs/template_plugin/tests/functional/shared_tests_instances/single_layer_tests/softmax.cpp +++ b/docs/template_plugin/tests/functional/shared_tests_instances/single_layer_tests/softmax.cpp @@ -29,8 +29,8 @@ const std::vector inputDynamicShape2D = { {{ngraph::Dimension::dynamic(), ngraph::Dimension::dynamic()}, {{1, 10}, {2, 10}, {10, 10}}} }; -const std::vector axis2D = { - 0, 1 +const std::vector axis2D = { + -2, -1, 0, 1 }; const auto params2D_static = testing::Combine( @@ -55,16 +55,16 @@ const auto params2D_dynamic = testing::Combine( INSTANTIATE_TEST_SUITE_P( smoke_SoftMax2D_static, - SoftMaxLayerTest, + SoftMax8LayerTest, params2D_static, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest::getTestCaseName ); INSTANTIATE_TEST_SUITE_P( smoke_SoftMax2D_dynamic, - SoftMaxLayerTest, + SoftMax8LayerTest, params2D_dynamic, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest::getTestCaseName ); const std::vector inputStaticShape4D = { @@ -79,10 +79,14 @@ const std::vector inputDynamicShape4D = { {{1, 100, 1, 1}, {50, 100, 4, 1}, {2, 100, 10, 1}}}, }; -const std::vector axis4D = {0, 1, 2, 3}; +const std::vector netPrecisions4D = { + ov::element::f32, +}; + +const std::vector axis4D = {0, 1, 2, 3, -1, -2, -3, -4}; const auto params4Dstatic = testing::Combine( - testing::ValuesIn(netPrecisions), + testing::ValuesIn(netPrecisions4D), ::testing::Values(ov::element::undefined), ::testing::Values(ov::element::undefined), testing::ValuesIn(ov::test::static_shapes_to_test_representation(inputStaticShape4D)), @@ -92,7 +96,7 @@ const auto params4Dstatic = testing::Combine( ); const auto params4Ddynamic = testing::Combine( - testing::ValuesIn(netPrecisions), + testing::ValuesIn(netPrecisions4D), ::testing::Values(ov::element::undefined), ::testing::Values(ov::element::undefined), testing::ValuesIn(inputDynamicShape4D), @@ -103,16 +107,69 @@ const auto params4Ddynamic = testing::Combine( INSTANTIATE_TEST_SUITE_P( smoke_SoftMax4D_static, - SoftMaxLayerTest, - params2D_static, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest, + params4Dstatic, + SoftMax8LayerTest::getTestCaseName ); INSTANTIATE_TEST_SUITE_P( smoke_SoftMax4D_dynamic, - SoftMaxLayerTest, - params2D_dynamic, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest, + params4Ddynamic, + SoftMax8LayerTest::getTestCaseName +); + + +const std::vector inputStaticShape5D = { + {1, 100, 1, 1, 1}, + {50, 100, 4, 1, 1}, + {2, 100, 10, 1, 1}, +}; + +const std::vector inputDynamicShape5D = { + {{ngraph::Dimension::dynamic(), 100, ngraph::Dimension(1, 10), 1, 1}, {{1, 100, 1, 1, 1}, {100, 100, 5, 1, 1}}}, + {{ngraph::Dimension::dynamic(), ngraph::Dimension::dynamic(), ngraph::Dimension::dynamic(), ngraph::Dimension::dynamic(), ngraph::Dimension::dynamic()}, + {{1, 100, 1, 1, 1}, {50, 100, 4, 1, 1}, {2, 100, 10, 1, 1}}}, +}; + +const std::vector netPrecisions5D = { + ov::element::f32, +}; + +const std::vector axis5D = {0, 1, 2, 3, 4, -1, -2, -3, -4, -5}; + +const auto params5Dstatic = testing::Combine( + testing::ValuesIn(netPrecisions5D), + ::testing::Values(ov::element::undefined), + ::testing::Values(ov::element::undefined), + testing::ValuesIn(ov::test::static_shapes_to_test_representation(inputStaticShape5D)), + testing::ValuesIn(axis5D), + testing::Values(CommonTestUtils::DEVICE_TEMPLATE), + testing::Values(ov::AnyMap()) +); + +const auto params5Ddynamic = testing::Combine( + testing::ValuesIn(netPrecisions5D), + ::testing::Values(ov::element::undefined), + ::testing::Values(ov::element::undefined), + testing::ValuesIn(inputDynamicShape5D), + testing::ValuesIn(axis5D), + testing::Values(CommonTestUtils::DEVICE_TEMPLATE), + testing::Values(ov::AnyMap()) +); + +INSTANTIATE_TEST_SUITE_P( + smoke_SoftMax5D_static, + SoftMax8LayerTest, + params5Dstatic, + SoftMax8LayerTest::getTestCaseName +); + +INSTANTIATE_TEST_SUITE_P( + smoke_SoftMax5D_dynamic, + SoftMax8LayerTest, + params5Ddynamic, + SoftMax8LayerTest::getTestCaseName ); } // namespace diff --git a/src/plugins/intel_gpu/include/intel_gpu/primitives/softmax.hpp b/src/plugins/intel_gpu/include/intel_gpu/primitives/softmax.hpp index 4c1f698fd24b72..733996cb49922a 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/primitives/softmax.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/primitives/softmax.hpp @@ -27,6 +27,7 @@ struct softmax : public primitive_base { /// @brief Enum type to specify softmax's normalization scope (see #dimension). enum dimension_t { + normalize_b, normalize_f, normalize_x, normalize_y, @@ -52,9 +53,10 @@ struct softmax : public primitive_base { /// Specific behaviour is determined by this parameter, as follows: /// - when set to @link softmax::dimension_t softmax::normalize_x @endlink each input row is normalized independently, /// - when set to @link softmax::dimension_t softmax::normalize_y @endlink each input column is normalized independently, + /// - when set to @link softmax::dimension_t softmax::normalize_z @endlink each input z-coordinate is normalized independently, /// - when set to @link softmax::dimension_t softmax::normalize_f @endlink each in-depth vector of input is normalized independently, /// - when set to @link softmax::dimension_t softmax::normalize_fyx @endlink each 3d image within input is normalized independently, - /// - when set to @link softmax::dimension_t softmax::normalize_bfyx @endlink everything is normalized, + /// - when set to @link softmax::dimension_t softmax::normalize_all @endlink everything is normalized, dimension_t dimension; }; /// @} diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp index a31cd6092de3a8..1d253ea0ee421c 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp @@ -47,6 +47,10 @@ struct softmax_impl : typed_primitive_impl_ocl { sm_params.dim = kernel_selector::softmax_dim::FEATURE; break; + case softmax::normalize_b: + sm_params.dim = kernel_selector::softmax_dim::BATCH; + break; + case softmax::normalize_f: sm_params.dim = kernel_selector::softmax_dim::FEATURE; break; diff --git a/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h b/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h index a834b32937919d..84306e2d61416b 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h +++ b/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h @@ -325,6 +325,7 @@ enum class SoftmaxDim { Y, Z, FEATURE, + BATCH, }; //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_items_class_kernel_base.cpp b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_items_class_kernel_base.cpp index 45ceabbd9cb42b..3b02f0530f71c2 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_items_class_kernel_base.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_items_class_kernel_base.cpp @@ -36,6 +36,7 @@ ParamsKey SoftmaxItemsClassKernelBase::GetDefaultSupportedKey() { k.EnableSoftmaxDim(SoftmaxDim::Y); k.EnableSoftmaxDim(SoftmaxDim::Z); k.EnableSoftmaxDim(SoftmaxDim::FEATURE); + k.EnableSoftmaxDim(SoftmaxDim::BATCH); k.EnableDifferentTypes(); k.EnableTensorOffset(); k.EnableTensorPitches(); @@ -53,6 +54,8 @@ std::vector SoftmaxItemsClassKernelBase::GetSoftmaxDimGlobalSizes(Softma return {out.X().v * out.Y().v, out.Feature().v, out.Batch().v}; case SoftmaxDim::FEATURE: return {out.X().v * out.Z().v, out.Y().v, out.Batch().v}; + case SoftmaxDim::BATCH: + return {out.X().v * out.Z().v, out.Y().v, out.Feature().v}; default: return {}; } @@ -70,14 +73,16 @@ JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params& MakeJitConstant("INPUT0_OTHER0_SIZE", "INPUT0_SIZE_Y"), MakeJitConstant("INPUT0_OTHER1_PITCH", "INPUT0_FEATURE_PITCH"), MakeJitConstant("INPUT0_OTHER2_PITCH", "INPUT0_Z_PITCH"), + MakeJitConstant("INPUT0_OTHER3_PITCH", "INPUT0_BATCH_PITCH"), MakeJitConstant("INPUT0_CLASS_PITCH", "INPUT0_X_PITCH"), MakeJitConstant("INPUT0_CLASS_NUM", "INPUT0_SIZE_X"), MakeJitConstant("OUTPUT_OTHER0_PITCH", "OUTPUT_Y_PITCH"), - MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Z_PITCH"), MakeJitConstant("OUTPUT_OTHER1_PITCH", "OUTPUT_FEATURE_PITCH"), + MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Z_PITCH"), + MakeJitConstant("OUTPUT_OTHER3_PITCH", "OUTPUT_BATCH_PITCH"), MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_X_PITCH"), }); - idx_order = {"batch", "other1", ndims == 5 ? "other2" : "0", "other0", "cls"}; + idx_order = {"other3", "other1", ndims == 5 ? "other2" : "0", "other0", "cls"}; break; case SoftmaxDim::Y: jit.AddConstants({ @@ -85,14 +90,16 @@ JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params& MakeJitConstant("INPUT0_OTHER0_SIZE", "INPUT0_SIZE_X"), MakeJitConstant("INPUT0_OTHER1_PITCH", "INPUT0_FEATURE_PITCH"), MakeJitConstant("INPUT0_OTHER2_PITCH", "INPUT0_Z_PITCH"), + MakeJitConstant("INPUT0_OTHER3_PITCH", "INPUT0_BATCH_PITCH"), MakeJitConstant("INPUT0_CLASS_PITCH", "INPUT0_Y_PITCH"), MakeJitConstant("INPUT0_CLASS_NUM", "INPUT0_SIZE_Y"), MakeJitConstant("OUTPUT_OTHER0_PITCH", "OUTPUT_X_PITCH"), MakeJitConstant("OUTPUT_OTHER1_PITCH", "OUTPUT_FEATURE_PITCH"), MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Z_PITCH"), + MakeJitConstant("OUTPUT_OTHER3_PITCH", "OUTPUT_BATCH_PITCH"), MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_Y_PITCH"), }); - idx_order = {"batch", "other1", ndims == 5 ? "other2" : "0", "cls", "other0"}; + idx_order = {"other3", "other1", ndims == 5 ? "other2" : "0", "cls", "other0"}; break; case SoftmaxDim::Z: jit.AddConstants({ @@ -100,14 +107,16 @@ JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params& MakeJitConstant("INPUT0_OTHER0_SIZE", "INPUT0_SIZE_X"), MakeJitConstant("INPUT0_OTHER1_PITCH", "INPUT0_FEATURE_PITCH"), MakeJitConstant("INPUT0_OTHER2_PITCH", "INPUT0_Y_PITCH"), + MakeJitConstant("INPUT0_OTHER3_PITCH", "INPUT0_BATCH_PITCH"), MakeJitConstant("INPUT0_CLASS_PITCH", "INPUT0_Z_PITCH"), MakeJitConstant("INPUT0_CLASS_NUM", "INPUT0_SIZE_Z"), MakeJitConstant("OUTPUT_OTHER0_PITCH", "OUTPUT_X_PITCH"), MakeJitConstant("OUTPUT_OTHER1_PITCH", "OUTPUT_FEATURE_PITCH"), MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Y_PITCH"), + MakeJitConstant("OUTPUT_OTHER3_PITCH", "OUTPUT_BATCH_PITCH"), MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_Z_PITCH"), }); - idx_order = {"batch", "other1", "cls", "other2", "other0"}; + idx_order = {"other3", "other1", "cls", "other2", "other0"}; break; case SoftmaxDim::FEATURE: jit.AddConstants({ @@ -115,14 +124,33 @@ JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params& MakeJitConstant("INPUT0_OTHER0_SIZE", "INPUT0_SIZE_X"), MakeJitConstant("INPUT0_OTHER1_PITCH", "INPUT0_Y_PITCH"), MakeJitConstant("INPUT0_OTHER2_PITCH", "INPUT0_Z_PITCH"), + MakeJitConstant("INPUT0_OTHER3_PITCH", "INPUT0_BATCH_PITCH"), MakeJitConstant("INPUT0_CLASS_PITCH", "INPUT0_FEATURE_PITCH"), MakeJitConstant("INPUT0_CLASS_NUM", "INPUT0_FEATURE_NUM"), MakeJitConstant("OUTPUT_OTHER0_PITCH", "OUTPUT_X_PITCH"), MakeJitConstant("OUTPUT_OTHER1_PITCH", "OUTPUT_Y_PITCH"), MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Z_PITCH"), + MakeJitConstant("OUTPUT_OTHER3_PITCH", "OUTPUT_BATCH_PITCH"), MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_FEATURE_PITCH"), }); - idx_order = {"batch", "cls", ndims == 5 ? "other2" : "0", "other1", "other0"}; + idx_order = {"other3", "cls", ndims == 5 ? "other2" : "0", "other1", "other0"}; + break; + case SoftmaxDim::BATCH: + jit.AddConstants({ + MakeJitConstant("INPUT0_OTHER0_PITCH", "INPUT0_X_PITCH"), + MakeJitConstant("INPUT0_OTHER0_SIZE", "INPUT0_SIZE_X"), + MakeJitConstant("INPUT0_OTHER1_PITCH", "INPUT0_Y_PITCH"), + MakeJitConstant("INPUT0_OTHER2_PITCH", "INPUT0_Z_PITCH"), + MakeJitConstant("INPUT0_OTHER3_PITCH", "INPUT0_FEATURE_PITCH"), + MakeJitConstant("INPUT0_CLASS_PITCH", "INPUT0_BATCH_PITCH"), + MakeJitConstant("INPUT0_CLASS_NUM", "INPUT0_BATCH_NUM"), + MakeJitConstant("OUTPUT_OTHER0_PITCH", "OUTPUT_X_PITCH"), + MakeJitConstant("OUTPUT_OTHER1_PITCH", "OUTPUT_Y_PITCH"), + MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Z_PITCH"), + MakeJitConstant("OUTPUT_OTHER3_PITCH", "OUTPUT_FEATURE_PITCH"), + MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_BATCH_PITCH"), + }); + idx_order = {"cls", "other3", ndims == 5 ? "other2" : "0", "other1", "other0"}; break; default: break; diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_items_class_optimized.cpp b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_items_class_optimized.cpp index e19701de2d6679..940bfb05fae0ea 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_items_class_optimized.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_items_class_optimized.cpp @@ -25,6 +25,9 @@ inline static size_t GetItemClassCount(const DataTensor& input, SoftmaxDim dim) case SoftmaxDim::FEATURE: item_class_count = input.Feature().v; break; + case SoftmaxDim::BATCH: + item_class_count = input.Batch().v; + break; default: break; } diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl b/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl index d0a2e72ad149be..e99ed03d69de6b 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl @@ -24,11 +24,11 @@ KERNEL(softmax_items_class_optimized)( const uint other2 = 0; #endif const uint other1 = get_group_id(1); - const uint batch = get_group_id(2); + const uint other3 = get_group_id(2); const uint simd_lane = get_sub_group_local_id(); - const uint in_depth_offset = batch*INPUT0_BATCH_PITCH + other2*INPUT0_OTHER2_PITCH + other1*INPUT0_OTHER1_PITCH + other0*INPUT0_OTHER0_PITCH + INPUT0_OFFSET; - const uint out_depth_offset = batch*OUTPUT_BATCH_PITCH + other2*OUTPUT_OTHER2_PITCH + other1*OUTPUT_OTHER1_PITCH + other0*OUTPUT_OTHER0_PITCH + OUTPUT_OFFSET; + const uint in_depth_offset = other3*INPUT0_OTHER3_PITCH + other2*INPUT0_OTHER2_PITCH + other1*INPUT0_OTHER1_PITCH + other0*INPUT0_OTHER0_PITCH + INPUT0_OFFSET; + const uint out_depth_offset = other3*OUTPUT_OTHER3_PITCH + other2*OUTPUT_OTHER2_PITCH + other1*OUTPUT_OTHER1_PITCH + other0*OUTPUT_OTHER0_PITCH + OUTPUT_OFFSET; ACCUMULATOR_TYPE max_value = UNIT_VAL_MIN; ACCUMULATOR_TYPE data[DATA_PER_WORKITEM]; diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl index 664c4e94db7c0b..3514c9f0c78a51 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl @@ -21,10 +21,10 @@ KERNEL(softmax)( const uint other2 = 0; #endif const uint other1 = get_global_id(1); - const uint batch = get_global_id(2); + const uint other3 = get_global_id(2); - const uint in_depth_offset = batch*INPUT0_BATCH_PITCH + other2*INPUT0_OTHER2_PITCH + other1*INPUT0_OTHER1_PITCH + other0*INPUT0_OTHER0_PITCH + INPUT0_OFFSET; - const uint out_depth_offset = batch*OUTPUT_BATCH_PITCH + other2*OUTPUT_OTHER2_PITCH + other1*OUTPUT_OTHER1_PITCH + other0*OUTPUT_OTHER0_PITCH + OUTPUT_OFFSET; + const uint in_depth_offset = other3*INPUT0_OTHER3_PITCH + other2*INPUT0_OTHER2_PITCH + other1*INPUT0_OTHER1_PITCH + other0*INPUT0_OTHER0_PITCH + INPUT0_OFFSET; + const uint out_depth_offset = other3*OUTPUT_OTHER3_PITCH + other2*OUTPUT_OTHER2_PITCH + other1*OUTPUT_OTHER1_PITCH + other0*OUTPUT_OTHER0_PITCH + OUTPUT_OFFSET; ACCUMULATOR_TYPE max_value = UNIT_VAL_MIN; ACCUMULATOR_TYPE data[INPUT0_CLASS_NUM]; diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_common.cpp b/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_common.cpp index 52596c9afdf544..44888925475e21 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_common.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_common.cpp @@ -271,7 +271,9 @@ std::string toString(SoftmaxDim d) { switch (d) { case SoftmaxDim::X: return "X"; case SoftmaxDim::Y: return "Y"; + case SoftmaxDim::Z: return "Z"; case SoftmaxDim::FEATURE: return "FEATURE"; + case SoftmaxDim::BATCH: return "BATCH"; default: return ""; } } diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.cpp b/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.cpp index 4badfbb2aa6f81..5d3c108279acd5 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.cpp @@ -270,9 +270,15 @@ void ParamsKey::EnableSoftmaxDim(SoftmaxDim d) { case SoftmaxDim::Y: key.restrict.val.dedicated.softmax.dimY = 1; break; + case SoftmaxDim::Z: + key.restrict.val.dedicated.softmax.dimZ = 1; + break; case SoftmaxDim::FEATURE: key.restrict.val.dedicated.softmax.dimFeature = 1; break; + case SoftmaxDim::BATCH: + key.restrict.val.dedicated.softmax.dimBatch = 1; + break; default: break; } diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.h b/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.h index 0285f13375c9f0..ade3458f65ee05 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.h +++ b/src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.h @@ -141,7 +141,9 @@ class ParamsKey { struct softmax_t { uint32_t dimX : 1; uint32_t dimY : 1; + uint32_t dimZ : 1; uint32_t dimFeature : 1; + uint32_t dimBatch : 1; } softmax; struct region_yolo_t { uint32_t dimX : 1; diff --git a/src/plugins/intel_gpu/src/plugin/ops/softmax.cpp b/src/plugins/intel_gpu/src/plugin/ops/softmax.cpp index f0e5a2921b0ec0..371717105977f5 100644 --- a/src/plugins/intel_gpu/src/plugin/ops/softmax.cpp +++ b/src/plugins/intel_gpu/src/plugin/ops/softmax.cpp @@ -17,8 +17,7 @@ namespace intel_gpu { static cldnn::softmax::dimension_t GetSoftmaxAxis(int64_t axis, size_t rank) { switch (axis) { - // FIXME: it seems that axis=0 should correspond to normalize_b; - case 0: return cldnn::softmax::normalize_all; + case 0: return cldnn::softmax::normalize_b; case 1: return cldnn::softmax::normalize_f; case 2: if (rank > 4) diff --git a/src/plugins/intel_gpu/tests/test_cases/softmax_gpu_test.cpp b/src/plugins/intel_gpu/tests/test_cases/softmax_gpu_test.cpp index 44a438a22edf25..14546bc3c31648 100644 --- a/src/plugins/intel_gpu/tests/test_cases/softmax_gpu_test.cpp +++ b/src/plugins/intel_gpu/tests/test_cases/softmax_gpu_test.cpp @@ -820,6 +820,96 @@ TEST(softmax_gpu_bfzyx_f16, normalize_all) { ASSERT_NEAR(sum, expected_sum, 0.001); } +TEST(softmax_gpu_bfyx_f32, normalize_b) { + // Input : 3x2x2x2 + static const int32_t x_size = 2, y_size = 2, feature_num = 2, + batch_num = 3, buf_size = x_size*y_size * batch_num * feature_num; + auto& engine = get_test_engine(); + + auto input = engine.allocate_memory({ data_types::f32, format::bfyx,{ batch_num, feature_num, x_size , y_size } }); + topology topology; + topology.add(input_layout("input", input->get_layout())); + topology.add(softmax("softmax", "input", softmax::normalize_b)); + + vector input_vec = { + // y0x0 y0x1 y1x0 y1x1 + /*b0f0*/0.1f, -0.1f, 0.9f, 1.5f, + /*b0f1*/3.f, 0.5f, 7.f, 12.f, + + /*b1f0*/0.2f, 0.2f, -10.f, 5.2f, + /*b1f1*/4.f, 0.5f, 8.f, 8.2f, + + /*b2f0*/0.2f, 0.2f, -10.f, 5.2f, + /*b2f1*/0.2f, 0.2f, -10.f, 5.2f + }; + set_values(input, input_vec); + + float expected_max_values[8] = { + 0.344253346f, //f=0, y=0, x=0 + 0.364854551f, //f=0, y=0, x=1 + + 0.999963085f, //f=0, y=1, x=0 + 0.493894592f, //f=0, y=1, x=1 + + 0.719294981f, //f=1, y=0, x=0 + 0.364854551f, //f=1, y=0, x=1 + + 0.73105857f, //f=1, y=1, x=0 + 0.977054322f //f=1, y=1, x=1 + }; + + network network(engine, topology); + + network.set_input_data("input", input); + auto outputs = network.execute(); + + EXPECT_EQ(outputs.size(), size_t(1)); + EXPECT_EQ(outputs.begin()->first, "softmax"); + + auto output = outputs.at("softmax").get_memory(); + cldnn::mem_lock output_ptr(output, get_test_stream()); + float out_buffer[buf_size]; + for (uint32_t i = 0; i < buf_size; i++) + { + out_buffer[i] = output_ptr[i]; + } + + float temp_max = 0; + float expected_sum = 1.0f; + int max_value_buffer_index = 0; + for (uint32_t i = 0; i < feature_num; i++) //this for loops will sum results in a batch per feature, we expect that: sum = 1.0f + { + for (uint32_t j = 0; j < y_size; j++) + { + for (uint32_t k = 0; k < x_size; k++) + { + float sum = 0.0f; + for (uint32_t l = 0; l < batch_num; l++) + { + int index = l * feature_num * x_size * y_size + + i * x_size * y_size + + j * x_size + + k; + + if (out_buffer[index] >= temp_max) + { + temp_max = out_buffer[index]; + } + + sum += out_buffer[index]; + } + EXPECT_EQ(true, are_equal(temp_max, expected_max_values[max_value_buffer_index])); + temp_max = 0; + max_value_buffer_index++; + + EXPECT_EQ(true, are_equal(sum, expected_sum)); + sum = 0.0f; + } + } + } +} + + ////////////////////////////////////////////////////////////////////////////// // // // Exhaustive Negative Matrix tests // diff --git a/src/tests/functional/plugin/cpu/shared_tests_instances/single_layer_tests/softmax.cpp b/src/tests/functional/plugin/cpu/shared_tests_instances/single_layer_tests/softmax.cpp index d2a0c7bb7091ae..e5d99a9ca3a861 100644 --- a/src/tests/functional/plugin/cpu/shared_tests_instances/single_layer_tests/softmax.cpp +++ b/src/tests/functional/plugin/cpu/shared_tests_instances/single_layer_tests/softmax.cpp @@ -28,8 +28,8 @@ const std::vector inputDynamicShape2D = { {{ngraph::Dimension::dynamic(), ngraph::Dimension::dynamic()}, {{1, 10}, {2, 10}, {10, 10}}} }; -const std::vector axis2D = { - 0, 1 +const std::vector axis2D = { + 0, 1, -1 }; const auto params2D_static = testing::Combine( @@ -54,16 +54,16 @@ const auto params2D_dynamic = testing::Combine( INSTANTIATE_TEST_SUITE_P( smoke_SoftMax2D_static, - SoftMaxLayerTest, + SoftMax8LayerTest, params2D_static, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest::getTestCaseName ); INSTANTIATE_TEST_SUITE_P( smoke_SoftMax2D_dynamic, - SoftMaxLayerTest, + SoftMax8LayerTest, params2D_dynamic, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest::getTestCaseName ); const std::vector inputStaticShape4D = { @@ -78,7 +78,7 @@ const std::vector inputDynamicShape4D = { {{1, 100, 1, 1}, {50, 100, 4, 1}, {2, 100, 10, 1}}}, }; -const std::vector axis4D = {0, 1, 2, 3}; +const std::vector axis4D = {0, 1, 2, 3, -1, -2, -3}; const auto params4Dstatic = testing::Combine( testing::ValuesIn(netPrecisions), @@ -102,16 +102,16 @@ const auto params4Ddynamic = testing::Combine( INSTANTIATE_TEST_SUITE_P( smoke_SoftMax4D_static, - SoftMaxLayerTest, + SoftMax8LayerTest, params2D_static, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest::getTestCaseName ); INSTANTIATE_TEST_SUITE_P( smoke_SoftMax4D_dynamic, - SoftMaxLayerTest, + SoftMax8LayerTest, params2D_dynamic, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest::getTestCaseName ); } // namespace diff --git a/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/softmax.cpp b/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/softmax.cpp index f495b99c00ba86..4446bf02dadc43 100644 --- a/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/softmax.cpp +++ b/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/softmax.cpp @@ -12,6 +12,7 @@ using namespace ov::test::subgraph; namespace { const std::vector netPrecisions = { + ov::element::f16, ov::element::f32, }; @@ -19,10 +20,11 @@ const std::vector inputShapes2D = { {1, 100}, {100, 1}, {10, 10}, + {100, 10} }; -const std::vector axis2D = { - 0, 1 +const std::vector axis2D = { + -2, -1, 0, 1 }; const auto params2D = testing::Combine( @@ -37,9 +39,9 @@ const auto params2D = testing::Combine( INSTANTIATE_TEST_SUITE_P( smoke_SoftMax2D, - SoftMaxLayerTest, + SoftMax8LayerTest, params2D, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest::getTestCaseName ); const std::vector inputShapes4D = { @@ -48,7 +50,7 @@ const std::vector inputShapes4D = { {2, 3, 4, 5}, }; -const std::vector axis4D = {0, 1, 2, 3}; +const std::vector axis4D = {-4, -3, -2, -1, 0, 1, 2, 3}; const auto params4D = testing::Combine( testing::ValuesIn(netPrecisions), @@ -62,9 +64,34 @@ const auto params4D = testing::Combine( INSTANTIATE_TEST_SUITE_P( smoke_SoftMax4D, - SoftMaxLayerTest, + SoftMax8LayerTest, params4D, - SoftMaxLayerTest::getTestCaseName + SoftMax8LayerTest::getTestCaseName +); + +const std::vector inputShapes5D = { + {1, 100, 1, 1, 1}, + {1, 3, 4, 3, 4}, + {2, 3, 4, 5, 6}, +}; + +const std::vector axis5D = {-5, -4, -3, -2, -1, 0, 1, 2, 3, 4}; + +const auto params5D = testing::Combine( + testing::ValuesIn(netPrecisions), + ::testing::Values(ov::element::undefined), + ::testing::Values(ov::element::undefined), + testing::ValuesIn(ov::test::static_shapes_to_test_representation(inputShapes5D)), + testing::ValuesIn(axis5D), + testing::Values(CommonTestUtils::DEVICE_GPU), + testing::Values(ov::AnyMap()) +); + +INSTANTIATE_TEST_SUITE_P( + smoke_SoftMax5D, + SoftMax8LayerTest, + params5D, + SoftMax8LayerTest::getTestCaseName ); } // namespace diff --git a/src/tests/functional/plugin/gpu/shared_tests_instances/skip_tests_config.cpp b/src/tests/functional/plugin/gpu/shared_tests_instances/skip_tests_config.cpp index 43a4690fb62290..c912a24048accd 100644 --- a/src/tests/functional/plugin/gpu/shared_tests_instances/skip_tests_config.cpp +++ b/src/tests/functional/plugin/gpu/shared_tests_instances/skip_tests_config.cpp @@ -20,8 +20,6 @@ std::vector disabledTestPatterns() { R"(.*(PreprocessTest).*(SetMeanImagePreProcessSetBlob).*)", R"(.*(PreprocessTest).*(ReverseInputChannelsPreProcessGetBlob).*)", R"(.*(InferRequestPreprocessDynamicallyInSetBlobTest).*)", - // TODO: Issue: 41462 - R"(.*(SoftMaxLayerTest).*axis=0.*)", // TODO: Issue: 46841 R"(.*(QuantGroupConvBackpropData3D).*)", @@ -69,8 +67,6 @@ std::vector disabledTestPatterns() { R"(.*smoke_LSTMSequenceCommonClip.*LSTMSequenceTest.*CompareWithRefs.*)", // Expected behavior. GPU plugin doesn't support i64 for eltwise power operation. R"(.*EltwiseLayerTest.*OpType=Pow.*NetType=i64.*)", - // TODO: Issue: 67486 - R"(.*(SoftMaxLayerTest).*)", // TODO: Issue: 68712 R"(.*.MatMul.*CompareWithRefs.*IS0=\(1.5\)_IS1=\(1.5\).*transpose_a=0.*transpose_b=1.*CONSTANT.*FP16.*UNSPECIFIED.*UNSPECIFIED.*ANY.*)", // TODO: Issue 69187 diff --git a/src/tests/functional/plugin/shared/include/single_layer_tests/softmax.hpp b/src/tests/functional/plugin/shared/include/single_layer_tests/softmax.hpp index ce8d4660b1c3e1..1272453656495b 100644 --- a/src/tests/functional/plugin/shared/include/single_layer_tests/softmax.hpp +++ b/src/tests/functional/plugin/shared/include/single_layer_tests/softmax.hpp @@ -18,6 +18,14 @@ TEST_P(SoftMaxLayerTest, CompareQueryModel) { query_model(); } +TEST_P(SoftMax8LayerTest, CompareWithRefs) { + run(); +} + +TEST_P(SoftMax8LayerTest, CompareQueryModel) { + query_model(); +} + } // namespace subgraph } // namespace test } // namespace ov diff --git a/src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/softmax.hpp b/src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/softmax.hpp index 59e1a72bf1e451..4d9e4ffcfb5dcc 100644 --- a/src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/softmax.hpp +++ b/src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/softmax.hpp @@ -4,30 +4,87 @@ #pragma once +#include "ngraph_functions/builders.hpp" + +#include "common_test_utils/common_utils.hpp" + #include "shared_test_classes/base/ov_subgraph.hpp" namespace ov { namespace test { namespace subgraph { -using SoftMaxTestParams = std::tuple< - ElementType, // netPrecision - ElementType, // inPrecision - ElementType, // outPrecision - InputShape, // Dynamic shape + Target static shapes - size_t, // axis - TargetDevice, // targetDevice - Config // config ->; - -class SoftMaxLayerTest : public testing::WithParamInterface, - virtual public ov::test::SubgraphBaseTest { +namespace aux { + +template +using SoftMaxTestParams = std::tuple; + +template +class SoftMaxLayerTestBase : public testing::WithParamInterface>, + virtual public ov::test::SubgraphBaseTest { public: - static std::string getTestCaseName(const testing::TestParamInfo &obj); + static std::string getTestCaseName(const testing::TestParamInfo> &obj) { + ElementType netType, inType, outType; + InputShape shapes; + AxisType axis; + TargetDevice targetDevice; + Config config; + std::tie(netType, inType, outType, shapes, axis, targetDevice, config) = obj.param; + + std::ostringstream result; + result << "NetType=" << netType << "_"; + result << "InType=" << inType << "_"; + result << "OutType=" << outType << "_"; + result << "IS=" << CommonTestUtils::partialShape2str({shapes.first}) << "_"; + result << "TS="; + for (const auto& item : shapes.second) { + result << CommonTestUtils::vec2str(item) << "_"; + } + result << "Axis=" << axis << "_"; + result << "Device=" << targetDevice; + + return result.str(); + } protected: - void SetUp() override; + void SetUp() override { + InputShape shapes; + ElementType ngPrc; + AxisType axis; + + std::tie(ngPrc, inType, outType, shapes, axis, targetDevice, configuration) = this->GetParam(); + init_input_shapes({shapes}); + + const auto params = ngraph::builder::makeDynamicParams(ngPrc, inputDynamicShapes); + const auto paramOuts = + ngraph::helpers::convert2OutputVector(ngraph::helpers::castOps2Nodes(params)); + + const auto softMax = std::make_shared(paramOuts.at(0), axis); + const ngraph::ResultVector results{std::make_shared(softMax)}; + + // TODO: This workaround is needed as there is no full support for f16 type in the reference implementation + if (ngPrc == element::Type_t::f16) { + abs_threshold = 0.005; + } + + function = std::make_shared(results, params, "softMax"); + } }; + +} // namespace aux + +using SoftMax1LayerTest = aux::SoftMaxLayerTestBase; +using SoftMax8LayerTest = aux::SoftMaxLayerTestBase; + +using SoftMaxLayerTest = SoftMax1LayerTest; + } // namespace subgraph } // namespace test } // namespace ov diff --git a/src/tests/functional/shared_test_classes/src/single_layer/softmax.cpp b/src/tests/functional/shared_test_classes/src/single_layer/softmax.cpp deleted file mode 100644 index d8e1010a03441f..00000000000000 --- a/src/tests/functional/shared_test_classes/src/single_layer/softmax.cpp +++ /dev/null @@ -1,57 +0,0 @@ -// Copyright (C) 2018-2022 Intel Corporation -// SPDX-License-Identifier: Apache-2.0 -// - -#include "ngraph_functions/builders.hpp" - -#include "common_test_utils/common_utils.hpp" - -#include "shared_test_classes/single_layer/softmax.hpp" - -namespace ov { -namespace test { -namespace subgraph { - -std::string SoftMaxLayerTest::getTestCaseName(const testing::TestParamInfo& obj) { - ElementType netType, inType, outType; - InputShape shapes; - size_t axis; - TargetDevice targetDevice; - Config config; - std::tie(netType, inType, outType, shapes, axis, targetDevice, config) = obj.param; - - std::ostringstream result; - result << "NetType=" << netType << "_"; - result << "InType=" << inType << "_"; - result << "OutType=" << outType << "_"; - result << "IS=" << CommonTestUtils::partialShape2str({shapes.first}) << "_"; - result << "TS="; - for (const auto& item : shapes.second) { - result << CommonTestUtils::vec2str(item) << "_"; - } - result << "Axis=" << axis << "_"; - result << "Device=" << targetDevice; - - return result.str(); -} - -void SoftMaxLayerTest::SetUp() { - InputShape shapes; - ElementType ngPrc; - size_t axis; - - std::tie(ngPrc, inType, outType, shapes, axis, targetDevice, configuration) = GetParam(); - init_input_shapes({shapes}); - - const auto params = ngraph::builder::makeDynamicParams(ngPrc, inputDynamicShapes); - const auto paramOuts = - ngraph::helpers::convert2OutputVector(ngraph::helpers::castOps2Nodes(params)); - - const auto softMax = std::make_shared(paramOuts.at(0), axis); - const ngraph::ResultVector results {std::make_shared(softMax)}; - - function = std::make_shared(results, params, "softMax"); -} -} // namespace subgraph -} // namespace test -} // namespace ov \ No newline at end of file