Skip to content

Commit

Permalink
[GPU] Support axis 0 for Softmax (openvinotoolkit#10364)
Browse files Browse the repository at this point in the history 
* [GPU] Modify Softmax single layer tests to check Softmax-8 is supported with axes in [-rank, rank) interval

* [GPU] Fix cldnn::softmax::dimension_t documentation

* [GPU] Fix ParamsKey::EnableSoftmaxDim

Support Z dimension.

* [GPU] Add Softmax single layer test that checks 5D case

Since some Softmax kernel code contains ifdef on 5-dimensional case,
a test case is needed that covers this functionality.

* [GPU] Support axis 0 in Softmax

* [GPU] Modify Softmax single layer tests to check axis 0

* [GPU] Modify Softmax items class optimized kernel to handle axis 0 correctly

Modify single layer test accordingly.

* [GPU] Modify Softmax unit-test to check softmax::normalize_b

* Split SoftMaxLayerTest into opset1 and opset8 versions

Use SoftMax8LayerTest in the tests throughout repository.
SoftMaxLayerTest now defaults to SoftMax1LayerTest for compatibility.

* [GPU] Add f16 test-case for Softmax single-layer test

Co-authored-by: tgubanova-lohika <[email protected]>
  • Loading branch information
@tgubanova-lohika
opoluektov-lohika and tgubanova-lohika authored May 26, 2022
1 parent 1cce278 commit ccd001f
Show file tree
Hide file tree
Showing 19 changed files with 347 additions and 122 deletions.
87 changes: 72 additions & 15 deletions docs/template_plugin/tests/functional/shared_tests_instances/single_layer_tests/softmax.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -29,8 +29,8 @@ const std::vector<ov::test::InputShape> inputDynamicShape2D = {
{{ngraph::Dimension::dynamic(), ngraph::Dimension::dynamic()}, {{1, 10}, {2, 10}, {10, 10}}}
};

const std::vector<size_t> axis2D = {
0, 1
const std::vector<int64_t> axis2D = {
-2, -1, 0, 1
};

const auto params2D_static = testing::Combine(
Expand All @@ -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<ov::Shape> inputStaticShape4D = {
Expand All @@ -79,10 +79,14 @@ const std::vector<ov::test::InputShape> inputDynamicShape4D = {
{{1, 100, 1, 1}, {50, 100, 4, 1}, {2, 100, 10, 1}}},
};

const std::vector<size_t> axis4D = {0, 1, 2, 3};
const std::vector<ov::test::ElementType> netPrecisions4D = {
ov::element::f32,
};

const std::vector<int64_t> 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)),
Expand All @@ -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),
Expand All @@ -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<ov::Shape> inputStaticShape5D = {
{1, 100, 1, 1, 1},
{50, 100, 4, 1, 1},
{2, 100, 10, 1, 1},
};

const std::vector<ov::test::InputShape> 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<ov::test::ElementType> netPrecisions5D = {
ov::element::f32,
};

const std::vector<int64_t> 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
4 changes: 3 additions & 1 deletion src/plugins/intel_gpu/include/intel_gpu/primitives/softmax.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,7 @@ struct softmax : public primitive_base<softmax> {

/// @brief Enum type to specify softmax's normalization scope (see #dimension).
enum dimension_t {
normalize_b,
normalize_f,
normalize_x,
normalize_y,
Expand All @@ -52,9 +53,10 @@ struct softmax : public primitive_base<softmax> {
/// 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;
};
/// @}
Expand Down
4 changes: 4 additions & 0 deletions src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,10 @@ struct softmax_impl : typed_primitive_impl_ocl<softmax> {
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;
Expand Down
1 change: 1 addition & 0 deletions src/plugins/intel_gpu/src/kernel_selector/common/common_types.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -325,6 +325,7 @@ enum class SoftmaxDim {
Y,
Z,
FEATURE,
BATCH,
};

////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
Expand Down
38 changes: 33 additions & 5 deletions ...l_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_items_class_kernel_base.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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();
Expand All @@ -53,6 +54,8 @@ std::vector<size_t> 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 {};
}
Expand All @@ -70,59 +73,84 @@ 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({
MakeJitConstant("INPUT0_OTHER0_PITCH", "INPUT0_X_PITCH"),
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({
MakeJitConstant("INPUT0_OTHER0_PITCH", "INPUT0_X_PITCH"),
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({
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_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;
Expand Down
3 changes: 3 additions & 0 deletions .../src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_items_class_optimized.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
Expand Down
6 changes: 3 additions & 3 deletions ...lugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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];
Expand Down
6 changes: 3 additions & 3 deletions src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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];
Expand Down
2 changes: 2 additions & 0 deletions src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_common.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 "";
}
}
Expand Down
6 changes: 6 additions & 0 deletions src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
Expand Down
2 changes: 2 additions & 0 deletions src/plugins/intel_gpu/src/kernel_selector/core/kernel_selector_params.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
3 changes: 1 addition & 2 deletions src/plugins/intel_gpu/src/plugin/ops/softmax.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
Loading

0 comments on commit ccd001f

Please sign in to comment.