diff options
| author | Freddie Liardet <frederick.liardet@arm.com> | 2021-05-07 09:03:12 +0100 |
|---|---|---|
| committer | frederick.liardet <frederick.liardet@arm.com> | 2021-05-12 12:48:12 +0000 |
| commit | c2d0e203ac28458b47c25d0129ccff6ca707ca56 (patch) | |
| tree | 1541a35eb88651b49dd0bad9759873de557da53d | |
| parent | 7d2f69fbe2d74d46eca368594760ab98fc6c577f (diff) | |
| download | ComputeLibrary-c2d0e203ac28458b47c25d0129ccff6ca707ca56.tar.gz | |
Fix bug in Select operator when input is 1xN
Fix issue where gpu select kernel would not compile when input was 1xN
size.
Also add relevant tests for 1xN input.
Resolves: COMPMID-4357
Signed-off-by: Freddie Liardet <frederick.liardet@arm.com>
Change-Id: Ib5f339379e9cc7ef05605312889dbad34cbfe5dd
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5620
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/select.cl | 8 | ||||
| -rw-r--r-- | tests/validation/CL/Select.cpp | 35 |
2 files changed, 38 insertions, 5 deletions
diff --git a/src/core/CL/cl_kernels/select.cl b/src/core/CL/cl_kernels/select.cl index 4d22d5bf07..6fd4bd4ce3 100644 --- a/src/core/CL/cl_kernels/select.cl +++ b/src/core/CL/cl_kernels/select.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2020 Arm Limited. + * Copyright (c) 2018-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -86,7 +86,7 @@ __kernel void select_same_rank( // Calculate result VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - res0 = select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0); + res0 = select(in_y, in_x, CONVERT(in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0, SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))); // Boundary-aware store STORE_VECTOR_SELECT(res, DATA_TYPE, (__global DATA_TYPE *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); @@ -152,7 +152,7 @@ __kernel void select_different_rank_2( // Calculate result VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - res0 = select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0); + res0 = select(in_y, in_x, CONVERT(in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0, SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))); // Boundary-aware store STORE_VECTOR_SELECT(res, DATA_TYPE, (__global DATA_TYPE *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); @@ -220,7 +220,7 @@ __kernel void select_different_rank_n( // Calculate result VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - res0 = select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0); + res0 = select(in_y, in_x, CONVERT(in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0, SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))); // Boundary-aware store STORE_VECTOR_SELECT(res, DATA_TYPE, (__global DATA_TYPE *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); diff --git a/tests/validation/CL/Select.cpp b/tests/validation/CL/Select.cpp index 3d7c61aab5..d3540cae48 100644 --- a/tests/validation/CL/Select.cpp +++ b/tests/validation/CL/Select.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2020 Arm Limited. + * Copyright (c) 2018-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -107,6 +107,17 @@ FIXTURE_DATA_TEST_CASE(RunSmall, validate(CLAccessor(_target), _reference); } +FIXTURE_DATA_TEST_CASE(RunOneDim, + CLSelectFixture<half>, + framework::DatasetMode::PRECOMMIT, + combine(combine(framework::dataset::make("Shape", TensorShape(1U, 16U)), + framework::dataset::make("has_same_rank", { false, true })), + framework::dataset::make("DataType", DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} + FIXTURE_DATA_TEST_CASE(RunLarge, CLSelectFixture<half>, framework::DatasetMode::NIGHTLY, @@ -127,6 +138,17 @@ FIXTURE_DATA_TEST_CASE(RunSmall, validate(CLAccessor(_target), _reference); } +FIXTURE_DATA_TEST_CASE(RunOneDim, + CLSelectFixture<float>, + framework::DatasetMode::PRECOMMIT, + combine(combine(framework::dataset::make("Shape", TensorShape(1U, 16U)), + framework::dataset::make("has_same_rank", { false, true })), + framework::dataset::make("DataType", DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} + FIXTURE_DATA_TEST_CASE(RunLarge, CLSelectFixture<float>, framework::DatasetMode::NIGHTLY, @@ -149,6 +171,17 @@ FIXTURE_DATA_TEST_CASE(RunSmall, validate(CLAccessor(_target), _reference); } +FIXTURE_DATA_TEST_CASE(RunOneDim, + CLSelectFixture<uint8_t>, + framework::DatasetMode::PRECOMMIT, + combine(combine(framework::dataset::make("Shape", TensorShape(1U, 16U)), + framework::dataset::make("has_same_rank", { false, true })), + framework::dataset::make("DataType", DataType::QASYMM8))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} + FIXTURE_DATA_TEST_CASE(RunLarge, CLSelectFixture<uint8_t>, framework::DatasetMode::NIGHTLY, |