From 28287afbea9549e8e2904084ae895c04cca88e95 Mon Sep 17 00:00:00 2001 From: Sheri Zhang Date: Tue, 25 Feb 2020 14:13:54 +0000 Subject: COMPMID-2792: Add support for QASYMM8_SIGNED in CLGEMMLowpMatrixMultiplyReshapedKernel Signed-off-by: Sheri Zhang Change-Id: I005e604253394f31173f37ec0296caf76b5e697c Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/c/VisualCompute/ComputeLibrary/+/227008 Tested-by: bsgcomp Reviewed-by: Sang-Hoon Park Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2798 Tested-by: Arm Jenkins Comments-Addressed: Arm Jenkins Reviewed-by: Michele Di Giorgio --- .../CLGEMMLowpMatrixMultiplyReshapedKernel.h | 10 +- src/core/CL/cl_kernels/gemmlowp.cl | 12 +- .../CLGEMMLowpMatrixMultiplyReshapedKernel.cpp | 6 +- .../CL/GEMMLowpMatrixMultiplyReshaped.cpp | 78 ++++++++++-- tests/validation/fixtures/GEMMLowpFixture.h | 139 ++++++++++++++++----- 5 files changed, 184 insertions(+), 61 deletions(-) diff --git a/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyReshapedKernel.h b/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyReshapedKernel.h index 36b9403b60..64a98128ce 100644 --- a/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyReshapedKernel.h +++ b/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyReshapedKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -49,9 +49,9 @@ public: CLGEMMLowpMatrixMultiplyReshapedKernel &operator=(CLGEMMLowpMatrixMultiplyReshapedKernel &&) = default; /** Initialise the kernel's input and output. * - * @param[in] input0 Input tensor containing the LHS reshaped matrix. Data type supported: QASYMM8. The number of dimensions for the LHS matrix must be less or equal than 4. + * @param[in] input0 Input tensor containing the LHS reshaped matrix. Data type supported: QASYMM8/QASYMM8_SIGNED. The number of dimensions for the LHS matrix must be less or equal than 4. * @param[in] input1 Input tensor containing the RHS reshaped matrix. Data type supported: same as @p input0. The number of dimensions for the RHS matrix must be less or equal than 3. - * @param[out] output Output tensor to store the result of matrix multiplication. Data type supported: same as @p input0 + * @param[out] output Output tensor to store the result of matrix multiplication. Data type supported: S32 * @param[in] lhs_info LHS matrix information used for reshaping the input0 tensor. Only the following values are supported: * lhs_info.m0: 2,3,4,5,6,7,8 * lhs_info.k0: 2,3,4,8,16 @@ -67,9 +67,9 @@ public: void configure(const ICLTensor *input0, const ICLTensor *input1, ICLTensor *output, const GEMMLHSMatrixInfo &lhs_info, const GEMMRHSMatrixInfo &rhs_info, const GEMMReshapeInfo &gemm_info); /** Static function to check if given info will lead to a valid configuration of @ref CLGEMMLowpMatrixMultiplyReshapedKernel * - * @param[in] input0 Input tensor info containing the LHS reshaped matrix. Data type supported: QASYMM8. The number of dimensions for the LHS matrix must be less or equal than 4. + * @param[in] input0 Input tensor info containing the LHS reshaped matrix. Data type supported: QASYMM8/QASYMM8_SIGNED. The number of dimensions for the LHS matrix must be less or equal than 4. * @param[in] input1 Input tensor info containing the RHS reshaped matrix. Data type supported: same as @p input0. The number of dimensions for the RHS matrix must be less or equal than 3. - * @param[in] output Output tensor info. Data type supported: same as @p input0 + * @param[in] output Output tensor info. Data type supported: S32 * @param[in] lhs_info LHS matrix information used for reshaping the input0 tensor. Only the following values are supported: * lhs_info.m0: 2,3,4,5,6,7,8 * lhs_info.k0: 2,3,4,8,16 diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index 239e039e10..8140b8f2a5 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -199,7 +199,7 @@ #define VECTOR_TYPE VEC_DATA_TYPE(DATA_TYPE, NUM_ELEMS_PROCESSED_PER_THREAD_X) #define VECTOR_ACC_TYPE VEC_DATA_TYPE(ACC_DATA_TYPE, NUM_ELEMS_PROCESSED_PER_THREAD_X) #define VECTOR_INT VEC_DATA_TYPE(int, NUM_ELEMS_PROCESSED_PER_THREAD_X) -/** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) in case both matrices have not beed reshaped +/** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) in case both matrices have not been reshaped * * @attention The number of matrix A columns needs to be passed at compile time using -DCOLS_A * @@ -481,7 +481,7 @@ __kernel void gemmlowp_mm_midgard(IMAGE_DECLARATION(src0), #endif // defined(NUM_ELEMS_PROCESSED_PER_THREAD_X) && defined(NUM_ELEMS_PROCESSED_PER_THREAD_Y) && defined(COLS_A) #if defined(M0) && defined(N0) && defined(K0) && defined(V0) && defined(H0) && defined(M) && defined(N) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices with QASYMM data type. +/** This OpenCL kernel computes the matrix multiplication between 2 matrices with QASYMM/QASYMM_SIGNED data type. * The LHS matrix must be reshaped with @ref CLGEMMReshapeLHSMatrixKernel and the M0xK0 must be NOT transposed * The RHS matrix must be reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the K0xN0 must be transposed * @@ -507,7 +507,7 @@ __kernel void gemmlowp_mm_midgard(IMAGE_DECLARATION(src0), * -# DEPTH_GEMM3D: The depth of the output in case it has to be reinterpreted as a 3D tensor * (HEIGHT_GEMM3D * DEPTH_GEMM3D) = columns LHS matrix NOT reshaped * - * @param[in] lhs_ptr Pointer to the LHS reshaped matrix. Supported data type: QASYMM8 + * @param[in] lhs_ptr Pointer to the LHS reshaped matrix. Supported data type: QASYMM8/QASYMM_SIGNED * @param[in] lhs_stride_x Stride of the LHS reshaped matrix in X dimension (in bytes) * @param[in] lhs_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] lhs_stride_y Stride of the LHS reshaped matrix in Y dimension (in bytes) @@ -519,7 +519,7 @@ __kernel void gemmlowp_mm_midgard(IMAGE_DECLARATION(src0), * @param[in] rhs_stride_y Stride of the RHS reshaped matrix in Y dimension (in bytes) * @param[in] rhs_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] rhs_offset_first_element_in_bytes The offset of the first element in the RHS reshaped matrix - * @param[out] dst_ptr Pointer to the destination matrix Supported data type: same as @p lhs_ptr + * @param[out] dst_ptr Pointer to the destination matrix Supported data type: S32 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes) * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes) @@ -583,10 +583,10 @@ __kernel void gemmlowp_mm_reshaped_lhs_nt_rhs_t(IMAGE_DECLARATION(lhs), #endif // defined(DUMMY_WORK_ITEMS) // Compute LHS matrix address - __global uchar *lhs_addr = lhs_ptr + lhs_offset_first_element_in_bytes + (y % V0) * (uint)LHS_OFFSET_X + (y / V0) * (uint)lhs_stride_y + (z * lhs_stride_z); + __global DATA_TYPE *lhs_addr = lhs_ptr + lhs_offset_first_element_in_bytes + (y % V0) * (uint)LHS_OFFSET_X + (y / V0) * (uint)lhs_stride_y + (z * lhs_stride_z); // Compute RHS matrix address - __global uchar *rhs_addr = rhs_ptr + rhs_offset_first_element_in_bytes + (x % H0) * (uint)RHS_OFFSET_X + (x / (uint)H0) * rhs_stride_y; + __global DATA_TYPE *rhs_addr = rhs_ptr + rhs_offset_first_element_in_bytes + (x % H0) * (uint)RHS_OFFSET_X + (x / (uint)H0) * rhs_stride_y; #if defined(MATRIX_B_DEPTH) // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 diff --git a/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyReshapedKernel.cpp b/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyReshapedKernel.cpp index 8d3aff6603..8dfeb773fe 100644 --- a/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyReshapedKernel.cpp +++ b/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyReshapedKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -54,7 +54,7 @@ Status validate_arguments(const ITensorInfo *input0, const ITensorInfo *input1, const GEMMReshapeInfo &gemm_info) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input0, input1, output); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input0, 1, DataType::QASYMM8); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input0, 1, DataType::QASYMM8, DataType::QASYMM8_SIGNED); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input0, input1); ARM_COMPUTE_RETURN_ERROR_ON_MSG(input0->num_dimensions() > 4, "The number of dimensions for the LHS matrix must be <= 4"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(input1->num_dimensions() > 3, "The number of dimensions for the RHS matrix must be <= 3"); @@ -309,4 +309,4 @@ void CLGEMMLowpMatrixMultiplyReshapedKernel::run(const Window &window, cl::Comma } while(window.slide_window_slice_3D(slice)); } -} // namespace arm_compute \ No newline at end of file +} // namespace arm_compute diff --git a/tests/validation/CL/GEMMLowpMatrixMultiplyReshaped.cpp b/tests/validation/CL/GEMMLowpMatrixMultiplyReshaped.cpp index e1f2b0c9a9..96e41b3f54 100644 --- a/tests/validation/CL/GEMMLowpMatrixMultiplyReshaped.cpp +++ b/tests/validation/CL/GEMMLowpMatrixMultiplyReshaped.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -81,7 +81,8 @@ const auto k_values = framework::dataset::make("K", 23); const auto b_values = framework::dataset::make("batch_size", 1, 3); /** M0 values to test - Precommit */ -const auto m0_values_precommit = framework::dataset::make("M0", {4, 6}); +const auto m0_values_precommit_1 = framework::dataset::make("M0", { 4 }); +const auto m0_values_precommit_2 = framework::dataset::make("M0", { 6 }); /** N0 values to test - Precommit */ const auto n0_values_precommit = framework::dataset::make("N0", { 4 }); @@ -119,26 +120,31 @@ const auto i_values_rhs = framework::dataset::make("interleave_rhs", { true, fal TEST_SUITE(CL) TEST_SUITE(GEMMLowpMatrixMultiplyReshaped) + +TEST_SUITE(QUANTIZED) + +TEST_SUITE(QASYMM8) FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMLowpMatrixMultiplyReshapedFixture, framework::DatasetMode::ALL, - combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( + combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( m_values, n_values), k_values), b_values), - m0_values_precommit), + m0_values_precommit_1), n0_values_precommit), k0_values_precommit), v0_values_precommit), h0_values_precommit), i_values_lhs), - i_values_rhs)) + i_values_rhs), + framework::dataset::make("DataType", { DataType::QASYMM8 }))) { // Validate output validate(CLAccessor(_target), _reference); } FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMLowpMatrixMultiplyReshapedFixture, framework::DatasetMode::DISABLED, - combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( + combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( m_values, n_values), k_values), @@ -149,33 +155,35 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMLowpMatrixMultiplyReshapedFixture, framew v0_values_nightly), h0_values_nightly), i_values_lhs), - i_values_rhs)) + i_values_rhs), + framework::dataset::make("DataType", { DataType::QASYMM8 }))) { // Validate output validate(CLAccessor(_target), _reference); } FIXTURE_DATA_TEST_CASE(RunSmall3D, CLGEMMLowpMatrixMultiplyReshaped3DFixture, framework::DatasetMode::ALL, - combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( + combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( m_w_values, m_h_values), n_values), k_values), b_values), - m0_values_precommit), + m0_values_precommit_1), n0_values_precommit), k0_values_precommit), v0_values_precommit), h0_values_precommit), i_values_lhs), - i_values_rhs)) + i_values_rhs), + framework::dataset::make("DataType", { DataType::QASYMM8 }))) { // Validate output validate(CLAccessor(_target), _reference); } FIXTURE_DATA_TEST_CASE(RunLarge3D, CLGEMMLowpMatrixMultiplyReshaped3DFixture, framework::DatasetMode::DISABLED, - combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( + combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( m_w_values, m_h_values), n_values), @@ -187,13 +195,57 @@ FIXTURE_DATA_TEST_CASE(RunLarge3D, CLGEMMLowpMatrixMultiplyReshaped3DFixture, fr v0_values_nightly), h0_values_nightly), i_values_lhs), - i_values_rhs)) + i_values_rhs), + framework::dataset::make("DataType", { DataType::QASYMM8 }))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() // QASYMM8 + +TEST_SUITE(QASYMM8_SIGNED) +FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMLowpMatrixMultiplyReshapedFixture, framework::DatasetMode::ALL, + combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( + m_values, + n_values), + k_values), + b_values), + m0_values_precommit_2), + n0_values_precommit), + k0_values_precommit), + v0_values_precommit), + h0_values_precommit), + i_values_lhs), + i_values_rhs), + framework::dataset::make("DataType", { DataType::QASYMM8_SIGNED }))) { // Validate output validate(CLAccessor(_target), _reference); } +FIXTURE_DATA_TEST_CASE(RunSmall3D, CLGEMMLowpMatrixMultiplyReshaped3DFixture, framework::DatasetMode::ALL, + combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(combine( + m_w_values, + m_h_values), + n_values), + k_values), + b_values), + m0_values_precommit_2), + n0_values_precommit), + k0_values_precommit), + v0_values_precommit), + h0_values_precommit), + i_values_lhs), + i_values_rhs), + framework::dataset::make("DataType", { DataType::QASYMM8_SIGNED }))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() // QASYMM8_SIGNED + +TEST_SUITE_END() // QUANTIZED TEST_SUITE_END() // GEMMLowpMatrixMultiplyReshaped TEST_SUITE_END() // CL } // namespace validation } // namespace test -} // namespace arm_compute \ No newline at end of file +} // namespace arm_compute diff --git a/tests/validation/fixtures/GEMMLowpFixture.h b/tests/validation/fixtures/GEMMLowpFixture.h index 1154d6c8de..d13ada9d64 100644 --- a/tests/validation/fixtures/GEMMLowpFixture.h +++ b/tests/validation/fixtures/GEMMLowpFixture.h @@ -656,7 +656,7 @@ class GEMMLowpMatrixMultiplyReshapedValidationFixture : public framework::Fixtur public: template void setup(unsigned int m, unsigned int n, unsigned int k, unsigned int batch_size, unsigned int m0, unsigned int n0, unsigned int k0, unsigned int v0, unsigned int h0, bool interleave_lhs, - bool interleave_rhs) + bool interleave_rhs, DataType data_type) { GEMMLHSMatrixInfo lhs_info; lhs_info.m0 = m0; @@ -676,24 +676,39 @@ public: const TensorShape lhs_shape(k, m, batch_size); const TensorShape rhs_shape(n, k, batch_size); - _target = compute_target(lhs_shape, rhs_shape, lhs_info, rhs_info); - _reference = compute_reference(lhs_shape, rhs_shape); + _target = compute_target(lhs_shape, rhs_shape, lhs_info, rhs_info, data_type); + _reference = compute_reference(lhs_shape, rhs_shape, data_type); } protected: template void fill(U &&tensor, int i) { - // Between 1 and 254 in order to avoid having -128 and 128 for the DOT product path - std::uniform_int_distribution<> distribution(1, 254); - library->fill(tensor, distribution, i); + switch(tensor.data_type()) + { + case DataType::QASYMM8: + { + // Between 1 and 254 in order to avoid having -128 and 128 for the DOT product path + std::uniform_int_distribution<> distribution(1, 254); + library->fill(tensor, distribution, i); + } + break; + case DataType::QASYMM8_SIGNED: + { + std::uniform_int_distribution<> distribution(-127, 126); + library->fill(tensor, distribution, i); + } + break; + default: + ARM_COMPUTE_ERROR("Unsupported data type"); + } } - TensorType compute_target(const TensorShape &lhs_shape, const TensorShape &rhs_shape, const GEMMLHSMatrixInfo &lhs_info, const GEMMRHSMatrixInfo &rhs_info) + TensorType compute_target(const TensorShape &lhs_shape, const TensorShape &rhs_shape, const GEMMLHSMatrixInfo &lhs_info, const GEMMRHSMatrixInfo &rhs_info, DataType data_type) { // Create tensors - TensorType lhs = create_tensor(lhs_shape, DataType::QASYMM8, 1); - TensorType rhs = create_tensor(rhs_shape, DataType::QASYMM8, 1); + TensorType lhs = create_tensor(lhs_shape, data_type, 1); + TensorType rhs = create_tensor(rhs_shape, data_type, 1); TensorType lhs_reshaped; TensorType rhs_reshaped; TensorType dst; @@ -740,21 +755,41 @@ protected: return dst; } - SimpleTensor compute_reference(const TensorShape &lhs_shape, const TensorShape &rhs_shape) + SimpleTensor compute_reference(const TensorShape &lhs_shape, const TensorShape &rhs_shape, DataType data_type) { TensorShape dst_shape = lhs_shape; dst_shape[0] = rhs_shape[0]; dst_shape[1] = lhs_shape[1]; - // Create reference - SimpleTensor lhs{ lhs_shape, DataType::QASYMM8, 1 }; - SimpleTensor rhs{ rhs_shape, DataType::QASYMM8, 1 }; + switch(data_type) + { + case DataType::QASYMM8: + { + // Create reference + SimpleTensor lhs{ lhs_shape, data_type, 1 }; + SimpleTensor rhs{ rhs_shape, data_type, 1 }; - // Fill reference - fill(lhs, 0); - fill(rhs, 1); + // Fill reference + fill(lhs, 0); + fill(rhs, 1); - return reference::gemmlowp_matrix_multiply_core(lhs, rhs, dst_shape, 0, 0); + return reference::gemmlowp_matrix_multiply_core(lhs, rhs, dst_shape, 0, 0); + } + case DataType::QASYMM8_SIGNED: + { + // Create reference + SimpleTensor lhs{ lhs_shape, data_type, 1 }; + SimpleTensor rhs{ rhs_shape, data_type, 1 }; + + // Fill reference + fill(lhs, 0); + fill(rhs, 1); + + return reference::gemmlowp_matrix_multiply_core(lhs, rhs, dst_shape, 0, 0); + } + default: + ARM_COMPUTE_ERROR("Unsupported data type"); + } } TensorType _target{}; @@ -767,7 +802,7 @@ class GEMMLowpMatrixMultiplyReshaped3DValidationFixture : public framework::Fixt public: template void setup(unsigned int m_w, unsigned int m_h, unsigned int n, unsigned int k, unsigned int batch_size, unsigned int m0, unsigned int n0, unsigned int k0, unsigned int v0, unsigned int h0, - bool interleave_lhs, bool interleave_rhs) + bool interleave_lhs, bool interleave_rhs, DataType data_type) { GEMMLHSMatrixInfo lhs_info; lhs_info.m0 = m0; @@ -790,24 +825,40 @@ public: const TensorShape lhs_shape(k, m, batch_size); const TensorShape rhs_shape(n, k, batch_size); - _target = compute_target(lhs_shape, rhs_shape, lhs_info, rhs_info, m_h); - _reference = compute_reference(lhs_shape, rhs_shape, m_h); + _target = compute_target(lhs_shape, rhs_shape, lhs_info, rhs_info, m_h, data_type); + _reference = compute_reference(lhs_shape, rhs_shape, m_h, data_type); } protected: template void fill(U &&tensor, int i) { - // Between 1 and 254 in order to avoid having -128 and 128 for the DOT product path - std::uniform_int_distribution<> distribution(1, 254); - library->fill(tensor, distribution, i); + switch(tensor.data_type()) + { + case DataType::QASYMM8: + { + // Between 1 and 254 in order to avoid having -128 and 128 for the DOT product path + std::uniform_int_distribution<> distribution(1, 254); + library->fill(tensor, distribution, i); + } + break; + case DataType::QASYMM8_SIGNED: + { + std::uniform_int_distribution<> distribution(-127, 126); + library->fill(tensor, distribution, i); + } + break; + default: + ARM_COMPUTE_ERROR("Unsupported data type"); + } } - TensorType compute_target(const TensorShape &lhs_shape, const TensorShape &rhs_shape, const GEMMLHSMatrixInfo &lhs_info, const GEMMRHSMatrixInfo &rhs_info, unsigned int m_h) + TensorType compute_target(const TensorShape &lhs_shape, const TensorShape &rhs_shape, const GEMMLHSMatrixInfo &lhs_info, const GEMMRHSMatrixInfo &rhs_info, unsigned int m_h, + DataType data_type) { // Create tensors - TensorType lhs = create_tensor(lhs_shape, DataType::QASYMM8, 1); - TensorType rhs = create_tensor(rhs_shape, DataType::QASYMM8, 1); + TensorType lhs = create_tensor(lhs_shape, data_type, 1); + TensorType rhs = create_tensor(rhs_shape, data_type, 1); TensorType lhs_reshaped; TensorType rhs_reshaped; TensorType dst; @@ -854,7 +905,7 @@ protected: return dst; } - SimpleTensor compute_reference(const TensorShape &lhs_shape, const TensorShape &rhs_shape, unsigned int m_h) + SimpleTensor compute_reference(const TensorShape &lhs_shape, const TensorShape &rhs_shape, unsigned int m_h, DataType data_type) { TensorShape dst_shape = lhs_shape; dst_shape.set(0, rhs_shape[0]); @@ -862,15 +913,35 @@ protected: dst_shape.set(2, m_h); dst_shape.set(3, lhs_shape[2]); - // Create reference - SimpleTensor lhs{ lhs_shape, DataType::QASYMM8, 1 }; - SimpleTensor rhs{ rhs_shape, DataType::QASYMM8, 1 }; + switch(data_type) + { + case DataType::QASYMM8: + { + // Create reference + SimpleTensor lhs{ lhs_shape, data_type, 1 }; + SimpleTensor rhs{ rhs_shape, data_type, 1 }; - // Fill reference - fill(lhs, 0); - fill(rhs, 1); + // Fill reference + fill(lhs, 0); + fill(rhs, 1); - return reference::gemmlowp_matrix_multiply_core(lhs, rhs, dst_shape, 0, 0); + return reference::gemmlowp_matrix_multiply_core(lhs, rhs, dst_shape, 0, 0); + } + case DataType::QASYMM8_SIGNED: + { + // Create reference + SimpleTensor lhs{ lhs_shape, data_type, 1 }; + SimpleTensor rhs{ rhs_shape, data_type, 1 }; + + // Fill reference + fill(lhs, 0); + fill(rhs, 1); + + return reference::gemmlowp_matrix_multiply_core(lhs, rhs, dst_shape, 0, 0); + } + default: + ARM_COMPUTE_ERROR("Unsupported data type"); + } } TensorType _target{}; -- cgit v1.2.1