From ebf6b8a00b77ea796d877bc1d0e6850c055318a6 Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Mon, 24 Sep 2018 16:31:08 +0100 Subject: COMPMID-1518: Add support for GEMM3D in CLGEMMLowpMatrixMultiplyCore Change-Id: Ib14ac821ee5d4aff80bd602cd3e76e7018abb5e6 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/150268 Tested-by: bsgcomp Reviewed-by: Isabella Gottardi Reviewed-by: Michele DiGiorgio --- .../CL/kernels/CLGEMMLowpMatrixMultiplyKernel.h | 3 + .../core/CL/kernels/CLGEMMLowpReductionKernel.h | 3 +- src/core/CL/cl_kernels/gemmlowp.cl | 673 +++++++++++++++++++-- .../CL/kernels/CLGEMMLowpMatrixMultiplyKernel.cpp | 164 +++-- .../kernels/CLGEMMLowpOffsetContributionKernel.cpp | 32 +- src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp | 4 +- src/runtime/CL/functions/CLGEMM.cpp | 19 +- .../CL/functions/CLGEMMConvolutionLayer.cpp | 5 +- .../CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp | 43 +- tests/datasets/LargeGEMMLowpDataset.h | 22 +- tests/datasets/SmallGEMMLowpDataset.h | 28 +- tests/validate_examples/cl_gemm.cpp | 2 +- tests/validation/CL/GEMMLowp.cpp | 30 +- .../validation/fixtures/GEMMLowpAssemblyFixture.h | 2 +- tests/validation/fixtures/GEMMLowpFixture.h | 24 +- tests/validation/reference/GEMMLowp.cpp | 71 ++- tests/validation/reference/GEMMLowp.h | 6 +- 17 files changed, 952 insertions(+), 179 deletions(-) diff --git a/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyKernel.h b/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyKernel.h index b96e978b66..82dcd93ce6 100644 --- a/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyKernel.h +++ b/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyKernel.h @@ -81,6 +81,9 @@ private: const ICLTensor *_input0; const ICLTensor *_input1; ICLTensor *_output; + bool _slide_matrix_b; + bool _reinterpret_input_as_3d; + bool _reinterpret_output_as_3d; }; } // namespace arm_compute #endif /*__ARM_COMPUTE_CLGEMMLOWPMATRIXMULTIPLYKERNEL_H__*/ diff --git a/arm_compute/core/CL/kernels/CLGEMMLowpReductionKernel.h b/arm_compute/core/CL/kernels/CLGEMMLowpReductionKernel.h index 12c12ef99a..690afdbf84 100644 --- a/arm_compute/core/CL/kernels/CLGEMMLowpReductionKernel.h +++ b/arm_compute/core/CL/kernels/CLGEMMLowpReductionKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -111,5 +111,4 @@ public: void run(const Window &window, cl::CommandQueue &queue) override; }; } // namespace arm_compute - #endif /* __ARM_COMPUTE_CLGEMMLOWREDUCTIONKERNEL_H__ */ diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index 12ac811cc7..e52f1ea486 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -40,6 +40,12 @@ * @note The transposition width step (mult_transpose1xW_width * 4) must be passed at compile time using -DTRANSPOSE1XW_WIDTH_STEP (i.e. -DTRANSPOSE1XW_WIDTH_STEP=2) * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * + * @note In case the output has to be reinterpreted as a 3D tensor (i.e. output of convolution layer), the following information must be passed at compile time: + * -# REINTERPRET_OUTPUT_AS_3D: To reinterpret the output as 3D + * -# HEIGHT_GEMM3D: The height of the output in case it has to be reinterpreted as a 3D tensor. + * -# DEPTH_GEMM3D: The depth of the output in case it has to be reinterpreted as a 3D tensor + * (HEIGHT_GEMM3D * DEPTH_GEMM3D) = columns matrix A NOT reshaped + * * @param[in] src0_ptr Pointer to the source matrix. Supported data type: QASYMM8 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes) @@ -58,13 +64,26 @@ * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes) * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix + * @param[in] src0_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] src1_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] cross_plane_pad (Optional) Bottom paddings in unit of elements (only if defined REINTERPRET_OUTPUT_AS_3D) */ __kernel void gemmlowp_mm_interleaved_transposed_midgard(IMAGE_DECLARATION(src0), IMAGE_DECLARATION(src1), - IMAGE_DECLARATION(dst)) + IMAGE_DECLARATION(dst), + uint src0_stride_z, + uint src1_stride_z, + uint dst_stride_z +#if defined(REINTERPRET_OUTPUT_AS_3D) + , + uint cross_plane_pad +#endif // REINTERPRET_OUTPUT_AS_3D + ) { - int x = get_global_id(0) / TRANSPOSE1XW_WIDTH_STEP; - int y = get_global_id(1) / MULT_INTERLEAVE4X4_HEIGHT; + const int x = get_global_id(0) / TRANSPOSE1XW_WIDTH_STEP; + const int y = get_global_id(1) / MULT_INTERLEAVE4X4_HEIGHT; + const int z = get_global_id(2); // Offset const int offset_row_a = (get_global_id(1) % MULT_INTERLEAVE4X4_HEIGHT) * 4; @@ -75,6 +94,13 @@ __kernel void gemmlowp_mm_interleaved_transposed_midgard(IMAGE_DECLARATION(src0) __global uchar *src_addr_a = (__global uchar *)(src0_ptr + y * src0_stride_y + src0_offset_first_element_in_bytes); __global uchar *src_addr_b = (__global uchar *)(src1_ptr + x * src1_stride_y + src1_offset_first_element_in_bytes); +#if defined(MATRIX_B_DEPTH) + // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 + src_addr_b += (z % MATRIX_B_DEPTH) * src1_stride_z; +#else // defined(MATRIX_B_DEPTH) + src_addr_b += z * src1_stride_z; +#endif // defined(MATRIX_B_DEPTH) + // Compute end row address for matrix B __global uchar *src_end_addr_b = src_addr_b + COLS_B; @@ -122,11 +148,49 @@ __kernel void gemmlowp_mm_interleaved_transposed_midgard(IMAGE_DECLARATION(src0) // Compute destination address Image dst = CONVERT_TO_IMAGE_STRUCT(dst); +#if defined(REINTERPRET_OUTPUT_AS_3D) + // Since we store a 2D output tile in a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zout) is calculated dividing M (get_global_id(1) * 4) by HEIGHT_GEMM3D + uint4 zout = ((uint4)(0, 1, 2, 3) + (uint4)(get_global_id(1) * 4)) / (uint4)HEIGHT_GEMM3D; + zout = min(DEPTH_GEMM3D - 1, zout); + + // Add offset due to the cross plane paddings + zout *= (cross_plane_pad * dst_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply dst_stride_z by DEPTH_GEMM3D + dst.ptr += z * dst_stride_z * DEPTH_GEMM3D; + // Store 4x4 block - vstore4(c00, 0, (__global int *)(offset(&dst, 0, 0))); - vstore4(c10, 0, (__global int *)(offset(&dst, 0, 1))); - vstore4(c20, 0, (__global int *)(offset(&dst, 0, 2))); - vstore4(c30, 0, (__global int *)(offset(&dst, 0, 3))); + vstore4(c00, 0, (__global int *)(dst.ptr + 0 * dst_stride_y + zout.s0)); + vstore4(c10, 0, (__global int *)(dst.ptr + 1 * dst_stride_y + zout.s1)); + vstore4(c20, 0, (__global int *)(dst.ptr + 2 * dst_stride_y + zout.s2)); + vstore4(c30, 0, (__global int *)(dst.ptr + 3 * dst_stride_y + zout.s3)); + +#else // defined(REINTERPRET_OUTPUT_AS_3D) + // Add offset for batched GEMM + dst.ptr += z * dst_stride_z; + + // Store 4x4 block + vstore4(c00, 0, (__global int *)(dst.ptr + 0 * dst_stride_y)); + vstore4(c10, 0, (__global int *)(dst.ptr + 1 * dst_stride_y)); + vstore4(c20, 0, (__global int *)(dst.ptr + 2 * dst_stride_y)); + vstore4(c30, 0, (__global int *)(dst.ptr + 3 * dst_stride_y)); +#endif // defined(REINTERPRET_OUTPUT_AS_3D) } /** This OpenCL kernel is optimized for Bifrost and computes the matrix multiplication between matrix A (src0) and matrix B (src1) @@ -136,6 +200,12 @@ __kernel void gemmlowp_mm_interleaved_transposed_midgard(IMAGE_DECLARATION(src0) * @note The transposition width step (mult_transpose1xW_width * 4) must be passed at compile time using -DTRANSPOSE1XW_WIDTH_STEP (i.e. -DTRANSPOSE1XW_WIDTH_STEP=2) * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * + * @note In case the output has to be reinterpreted as a 3D tensor (i.e. output of convolution layer), the following information must be passed at compile time: + * -# REINTERPRET_OUTPUT_AS_3D: To reinterpret the output as 3D + * -# HEIGHT_GEMM3D: The height of the output in case it has to be reinterpreted as a 3D tensor. + * -# DEPTH_GEMM3D: The depth of the output in case it has to be reinterpreted as a 3D tensor + * (HEIGHT_GEMM3D * DEPTH_GEMM3D) = columns matrix A NOT reshaped + * * @param[in] src0_ptr Pointer to the source matrix. Supported data type: QASYMM8 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes) @@ -154,13 +224,26 @@ __kernel void gemmlowp_mm_interleaved_transposed_midgard(IMAGE_DECLARATION(src0) * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes) * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix + * @param[in] src0_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] src1_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] cross_plane_pad (Optional) Bottom paddings in unit of elements (only if defined REINTERPRET_OUTPUT_AS_3D) */ __kernel void gemmlowp_mm_interleaved_transposed_bifrost(IMAGE_DECLARATION(src0), IMAGE_DECLARATION(src1), - IMAGE_DECLARATION(dst)) + IMAGE_DECLARATION(dst), + uint src0_stride_z, + uint src1_stride_z, + uint dst_stride_z +#if defined(REINTERPRET_OUTPUT_AS_3D) + , + uint cross_plane_pad +#endif // REINTERPRET_OUTPUT_AS_3D + ) { - int x = get_global_id(0) / TRANSPOSE1XW_WIDTH_STEP; - int y = get_global_id(1) / MULT_INTERLEAVE4X4_HEIGHT; + const int x = get_global_id(0) / TRANSPOSE1XW_WIDTH_STEP; + const int y = get_global_id(1) / MULT_INTERLEAVE4X4_HEIGHT; + const int z = get_global_id(2); // Offset const int offset_row_a = (get_global_id(1) % MULT_INTERLEAVE4X4_HEIGHT) * 4; @@ -171,6 +254,13 @@ __kernel void gemmlowp_mm_interleaved_transposed_bifrost(IMAGE_DECLARATION(src0) __global uchar *src_addr_a = (__global uchar *)(src0_ptr + y * src0_stride_y + src0_offset_first_element_in_bytes); __global uchar *src_addr_b = (__global uchar *)(src1_ptr + x * src1_stride_y + src1_offset_first_element_in_bytes); +#if defined(MATRIX_B_DEPTH) + // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 + src_addr_b += (z % MATRIX_B_DEPTH) * src1_stride_z; +#else // defined(MATRIX_B_DEPTH) + src_addr_b += z * src1_stride_z; +#endif // defined(MATRIX_B_DEPTH) + // Compute end row address for matrix B __global uchar *src_end_addr_b = src_addr_b + COLS_B; @@ -416,11 +506,49 @@ __kernel void gemmlowp_mm_interleaved_transposed_bifrost(IMAGE_DECLARATION(src0) // Compute destination address Image dst = CONVERT_TO_IMAGE_STRUCT(dst); +#if defined(REINTERPRET_OUTPUT_AS_3D) + // Since we store a 2D output tile in a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zout) is calculated dividing M (get_global_id(1) * 4) by HEIGHT_GEMM3D + uint4 zout = ((uint4)(0, 1, 2, 3) + (uint4)(get_global_id(1) * 4)) / (uint4)HEIGHT_GEMM3D; + zout = min(DEPTH_GEMM3D - 1, zout); + + // Add offset due to the cross plane paddings + zout *= (cross_plane_pad * dst_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply dst_stride_z by DEPTH_GEMM3D + dst.ptr += z * dst_stride_z * DEPTH_GEMM3D; + // Store 4x4 block - vstore4((int4)(c00, c01, c02, c03), 0, (__global int *)(offset(&dst, 0, 0))); - vstore4((int4)(c10, c11, c12, c13), 0, (__global int *)(offset(&dst, 0, 1))); - vstore4((int4)(c20, c21, c22, c23), 0, (__global int *)(offset(&dst, 0, 2))); - vstore4((int4)(c30, c31, c32, c33), 0, (__global int *)(offset(&dst, 0, 3))); + vstore4((int4)(c00, c01, c02, c03), 0, (__global int *)(dst.ptr + 0 * dst_stride_y + zout.s0)); + vstore4((int4)(c10, c11, c12, c13), 0, (__global int *)(dst.ptr + 1 * dst_stride_y + zout.s1)); + vstore4((int4)(c20, c21, c22, c23), 0, (__global int *)(dst.ptr + 2 * dst_stride_y + zout.s2)); + vstore4((int4)(c30, c31, c32, c33), 0, (__global int *)(dst.ptr + 3 * dst_stride_y + zout.s3)); + +#else // defined(REINTERPRET_OUTPUT_AS_3D) + // Add offset for batched GEMM + dst.ptr += z * dst_stride_z; + + // Store 4x4 block + vstore4((int4)(c00, c01, c02, c03), 0, (__global int *)(dst.ptr + 0 * dst_stride_y)); + vstore4((int4)(c10, c11, c12, c13), 0, (__global int *)(dst.ptr + 1 * dst_stride_y)); + vstore4((int4)(c20, c21, c22, c23), 0, (__global int *)(dst.ptr + 2 * dst_stride_y)); + vstore4((int4)(c30, c31, c32, c33), 0, (__global int *)(dst.ptr + 3 * dst_stride_y)); +#endif // defined(REINTERPRET_OUTPUT_AS_3D) } #if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) @@ -431,6 +559,12 @@ __kernel void gemmlowp_mm_interleaved_transposed_bifrost(IMAGE_DECLARATION(src0) * @note The transposition width step (mult_transpose1xW_width * 4) must be passed at compile time using -DTRANSPOSE1XW_WIDTH_STEP (i.e. -DTRANSPOSE1XW_WIDTH_STEP=2) * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * + * @note In case the output has to be reinterpreted as a 3D tensor (i.e. output of convolution layer), the following information must be passed at compile time: + * -# REINTERPRET_OUTPUT_AS_3D: To reinterpret the output as 3D + * -# HEIGHT_GEMM3D: The height of the output in case it has to be reinterpreted as a 3D tensor. + * -# DEPTH_GEMM3D: The depth of the output in case it has to be reinterpreted as a 3D tensor + * (HEIGHT_GEMM3D * DEPTH_GEMM3D) = columns matrix A NOT reshaped + * * @param[in] src0_ptr Pointer to the source matrix. Supported data type: QASYMM8 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes) @@ -449,13 +583,26 @@ __kernel void gemmlowp_mm_interleaved_transposed_bifrost(IMAGE_DECLARATION(src0) * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes) * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix + * @param[in] src0_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] src1_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] cross_plane_pad (Optional) Bottom paddings in unit of elements (only if defined REINTERPRET_OUTPUT_AS_3D) */ __kernel void gemmlowp_mm_interleaved_transposed_bifrost_dot8(IMAGE_DECLARATION(src0), IMAGE_DECLARATION(src1), - IMAGE_DECLARATION(dst)) + IMAGE_DECLARATION(dst), + uint src0_stride_z, + uint src1_stride_z, + uint dst_stride_z +#if defined(REINTERPRET_OUTPUT_AS_3D) + , + uint cross_plane_pad +#endif // REINTERPRET_OUTPUT_AS_3D + ) { - int x = get_global_id(0) / TRANSPOSE1XW_WIDTH_STEP; - int y = get_global_id(1) / MULT_INTERLEAVE4X4_HEIGHT; + const int x = get_global_id(0) / TRANSPOSE1XW_WIDTH_STEP; + const int y = get_global_id(1) / MULT_INTERLEAVE4X4_HEIGHT; + const int z = get_global_id(2); // Offset const int offset_row_a = (get_global_id(1) % MULT_INTERLEAVE4X4_HEIGHT) * 4; @@ -466,6 +613,13 @@ __kernel void gemmlowp_mm_interleaved_transposed_bifrost_dot8(IMAGE_DECLARATION( __global uchar *src_addr_a = (__global uchar *)(src0_ptr + y * src0_stride_y + src0_offset_first_element_in_bytes); __global uchar *src_addr_b = (__global uchar *)(src1_ptr + x * src1_stride_y + src1_offset_first_element_in_bytes); +#if defined(MATRIX_B_DEPTH) + // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 + src_addr_b += (z % MATRIX_B_DEPTH) * src1_stride_z; +#else // defined(MATRIX_B_DEPTH) + src_addr_b += z * src1_stride_z; +#endif // defined(MATRIX_B_DEPTH) + // Compute end row address for matrix B __global uchar *src_end_addr_b = src_addr_b + COLS_B; @@ -581,11 +735,49 @@ __kernel void gemmlowp_mm_interleaved_transposed_bifrost_dot8(IMAGE_DECLARATION( // Compute destination address Image dst = CONVERT_TO_IMAGE_STRUCT(dst); +#if defined(REINTERPRET_OUTPUT_AS_3D) + // Since we store a 2D output tile in a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zout) is calculated dividing M (get_global_id(1) * 4) by HEIGHT_GEMM3D + uint4 zout = ((uint4)(0, 1, 2, 3) + (uint4)(get_global_id(1) * 4)) / (uint4)HEIGHT_GEMM3D; + zout = min(DEPTH_GEMM3D - 1, zout); + + // Add offset due to the cross plane paddings + zout *= (cross_plane_pad * dst_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply dst_stride_z by DEPTH_GEMM3D + dst.ptr += z * dst_stride_z * DEPTH_GEMM3D; + + // Store 4x4 block + vstore4((int4)(c00, c01, c02, c03), 0, (__global int *)(dst.ptr + 0 * dst_stride_y + zout.s0)); + vstore4((int4)(c10, c11, c12, c13), 0, (__global int *)(dst.ptr + 1 * dst_stride_y + zout.s1)); + vstore4((int4)(c20, c21, c22, c23), 0, (__global int *)(dst.ptr + 2 * dst_stride_y + zout.s2)); + vstore4((int4)(c30, c31, c32, c33), 0, (__global int *)(dst.ptr + 3 * dst_stride_y + zout.s3)); + +#else // defined(REINTERPRET_OUTPUT_AS_3D) + // Add offset for batched GEMM + dst.ptr += z * dst_stride_z; + // Store 4x4 block - vstore4((int4)(c00, c01, c02, c03), 0, (__global int *)(offset(&dst, 0, 0))); - vstore4((int4)(c10, c11, c12, c13), 0, (__global int *)(offset(&dst, 0, 1))); - vstore4((int4)(c20, c21, c22, c23), 0, (__global int *)(offset(&dst, 0, 2))); - vstore4((int4)(c30, c31, c32, c33), 0, (__global int *)(offset(&dst, 0, 3))); + vstore4((int4)(c00, c01, c02, c03), 0, (__global int *)(dst.ptr + 0 * dst_stride_y)); + vstore4((int4)(c10, c11, c12, c13), 0, (__global int *)(dst.ptr + 1 * dst_stride_y)); + vstore4((int4)(c20, c21, c22, c23), 0, (__global int *)(dst.ptr + 2 * dst_stride_y)); + vstore4((int4)(c30, c31, c32, c33), 0, (__global int *)(dst.ptr + 3 * dst_stride_y)); +#endif // defined(REINTERPRET_OUTPUT_AS_3D) } #endif // defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) @@ -599,6 +791,13 @@ __kernel void gemmlowp_mm_interleaved_transposed_bifrost_dot8(IMAGE_DECLARATION( * * @attention The number of matrix A columns needs to be passed at compile time using -DCOLS_A * + * @note In case the input or output have to be reinterpreted as a 3D tensor, the following information must be passed at compile time: + * -# REINTERPRET_INPUT_AS_3D: To reinterpret the input as 3D + * -# REINTERPRET_OUTPUT_AS_3D: To reinterpret the output as 3D + * -# HEIGHT_GEMM3D: The height of the output in case it has to be reinterpreted as a 3D tensor. + * -# DEPTH_GEMM3D: The depth of the output in case it has to be reinterpreted as a 3D tensor + * (HEIGHT_GEMM3D * DEPTH_GEMM3D) = columns matrix A NOT reshaped + * * @param[in] src0_ptr Pointer to the source matrix. Supported data type: QASYMM8 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes) @@ -617,10 +816,27 @@ __kernel void gemmlowp_mm_interleaved_transposed_bifrost_dot8(IMAGE_DECLARATION( * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes) * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix + * @param[in] src0_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] src1_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] src_cross_plane_pad (Optional) Bottom paddings in unit of elements for the input tensor (only if defined REINTERPRET_INPUT_AS_3D) + * @param[in] dst_cross_plane_pad (Optional) Bottom paddings in unit of elements for the output tensor (only if defined REINTERPRET_OUTPUT_AS_3D) */ __kernel void gemmlowp_mm_midgard(IMAGE_DECLARATION(src0), IMAGE_DECLARATION(src1), - IMAGE_DECLARATION(dst)) + IMAGE_DECLARATION(dst), + uint src0_stride_z, + uint src1_stride_z, + uint dst_stride_z +#if defined(REINTERPRET_INPUT_AS_3D) + , + uint src_cross_plane_pad +#endif // REINTERPRET_INPUT_AS_3D +#if defined(REINTERPRET_OUTPUT_AS_3D) + , + uint dst_cross_plane_pad +#endif // REINTERPRET_OUTPUT_AS_3D + ) { int idx = get_global_id(0) * NUM_ELEMS_PROCESSED_PER_THREAD_X; @@ -633,6 +849,47 @@ __kernel void gemmlowp_mm_midgard(IMAGE_DECLARATION(src0), // Update address for the matrix B src_addr.s1 += idx; +#if defined(REINTERPRET_INPUT_AS_3D) + // Since we load a 2D input tile from a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zin) is calculated dividing M (get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y) by HEIGHT_GEMM3D + uint4 zin = ((uint4)(0, 1, 2, 3) + (uint4)(get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y)) / (uint4)HEIGHT_GEMM3D; + zin = min(DEPTH_GEMM3D - 1, zin); + + // Add offset due to the cross plane paddings + zin *= (src_cross_plane_pad * src0_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply src0_stride_z by DEPTH_GEMM3D + src_addr.s0 += get_global_id(2) * src0_stride_z * DEPTH_GEMM3D; + +#else // defined(REINTERPRET_INPUT_AS_3D) + + // Add offset for batched GEMM + src_addr.s0 += get_global_id(2) * src0_stride_z; + +#endif // defined(REINTERPRET_INPUT_AS_3D) + +#if defined(MATRIX_B_DEPTH) + // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 + src_addr.s1 += (get_global_id(2) % MATRIX_B_DEPTH) * src1_stride_z; +#else // defined(MATRIX_B_DEPTH) + src_addr.s1 += get_global_id(2) * src1_stride_z; +#endif // defined(MATRIX_B_DEPTH) + int end_row_vec_a = src_addr.s0 + COLS_A; VECTOR_UINT acc0 = 0; @@ -725,34 +982,95 @@ __kernel void gemmlowp_mm_midgard(IMAGE_DECLARATION(src0), #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 } + const int z = get_global_id(2); + // Compute destination address Image dst = CONVERT_TO_IMAGE_STRUCT(dst); +#if defined(REINTERPRET_OUTPUT_AS_3D) + // Since we store a 2D output tile in a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zout) is calculated dividing M (get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y) by HEIGHT_GEMM3D + uint8 zout = ((uint8)(0, 1, 2, 3, 4, 5, 6, 7) + (uint8)(get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y)) / (uint8)HEIGHT_GEMM3D; + zout = min(DEPTH_GEMM3D - 1, zout); + + // Add offset due to the cross plane paddings + zout *= (dst_cross_plane_pad * dst_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply dst_stride_z by DEPTH_GEMM3D + dst.ptr += z * dst_stride_z * DEPTH_GEMM3D; + + // Store the result + VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) + (CONVERT(acc0, VECTOR_INT), 0, (__global int *)(dst.ptr + 0 * dst_stride_y + zout.s0)); +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 + VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) + (CONVERT(acc1, VECTOR_INT), 0, (__global int *)(dst.ptr + 1 * dst_stride_y + zout.s1)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 + VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) + (CONVERT(acc2, VECTOR_INT), 0, (__global int *)(dst.ptr + 2 * dst_stride_y + zout.s2)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 + VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) + (CONVERT(acc3, VECTOR_INT), 0, (__global int *)(dst.ptr + 3 * dst_stride_y + zout.s3)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 + VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) + (CONVERT(acc4, VECTOR_INT), 0, (__global int *)(dst.ptr + 4 * dst_stride_y + zout.s4)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 + +#else // defined(REINTERPRET_OUTPUT_AS_3D) + // Add offset for batched GEMM + dst.ptr += z * dst_stride_z; + // Store the result VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) - (CONVERT(acc0, VECTOR_INT), 0, (__global int *)(offset(&dst, 0, 0))); + (CONVERT(acc0, VECTOR_INT), 0, (__global int *)(dst.ptr + 0 * dst_stride_y)); #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) - (CONVERT(acc1, VECTOR_INT), 0, (__global int *)(offset(&dst, 0, 1))); + (CONVERT(acc1, VECTOR_INT), 0, (__global int *)(dst.ptr + 1 * dst_stride_y)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) - (CONVERT(acc2, VECTOR_INT), 0, (__global int *)(offset(&dst, 0, 2))); + (CONVERT(acc2, VECTOR_INT), 0, (__global int *)(dst.ptr + 2 * dst_stride_y)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) - (CONVERT(acc3, VECTOR_INT), 0, (__global int *)(offset(&dst, 0, 3))); + (CONVERT(acc3, VECTOR_INT), 0, (__global int *)(dst.ptr + 3 * dst_stride_y)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 VSTORE(NUM_ELEMS_PROCESSED_PER_THREAD_X) - (CONVERT(acc4, VECTOR_INT), 0, (__global int *)(offset(&dst, 0, 4))); + (CONVERT(acc4, VECTOR_INT), 0, (__global int *)(dst.ptr + 4 * dst_stride_y)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 +#endif // defined(REINTERPRET_OUTPUT_AS_3D) } /** OpenCL kernel optimized for Bifrost architectures that computes the matrix multiplication between matrix A (src0) and matrix B (src1) in case both matrices have not beed reshaped * * @attention The number of matrix A columns needs to be passed at compile time using -DCOLS_A * + * @note In case the input or output have to be reinterpreted as a 3D tensor, the following information must be passed at compile time: + * -# REINTERPRET_INPUT_AS_3D: To reinterpret the input as 3D + * -# REINTERPRET_OUTPUT_AS_3D: To reinterpret the output as 3D + * -# HEIGHT_GEMM3D: The height of the output in case it has to be reinterpreted as a 3D tensor. + * -# DEPTH_GEMM3D: The depth of the output in case it has to be reinterpreted as a 3D tensor + * (HEIGHT_GEMM3D * DEPTH_GEMM3D) = columns matrix A NOT reshaped + * * @param[in] src0_ptr Pointer to the source matrix. Supported data type: QASYMM8 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes) @@ -771,10 +1089,27 @@ __kernel void gemmlowp_mm_midgard(IMAGE_DECLARATION(src0), * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes) * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix + * @param[in] src0_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] src1_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] src_cross_plane_pad (Optional) Bottom paddings in unit of elements for the input tensor (only if defined REINTERPRET_INPUT_AS_3D) + * @param[in] dst_cross_plane_pad (Optional) Bottom paddings in unit of elements for the output tensor (only if defined REINTERPRET_OUTPUT_AS_3D) */ __kernel void gemmlowp_mm_bifrost(IMAGE_DECLARATION(src0), IMAGE_DECLARATION(src1), - IMAGE_DECLARATION(dst)) + IMAGE_DECLARATION(dst), + uint src0_stride_z, + uint src1_stride_z, + uint dst_stride_z +#if defined(REINTERPRET_INPUT_AS_3D) + , + uint src_cross_plane_pad +#endif // REINTERPRET_INPUT_AS_3D +#if defined(REINTERPRET_OUTPUT_AS_3D) + , + uint dst_cross_plane_pad +#endif // REINTERPRET_OUTPUT_AS_3D + ) { int idx = get_global_id(0) * NUM_ELEMS_PROCESSED_PER_THREAD_X; @@ -787,6 +1122,47 @@ __kernel void gemmlowp_mm_bifrost(IMAGE_DECLARATION(src0), // Update address for the matrix B src_addr.s1 += idx; +#if defined(REINTERPRET_INPUT_AS_3D) + // Since we load a 2D input tile from a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zin) is calculated dividing M (get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y) by HEIGHT_GEMM3D + uint4 zin = ((uint4)(0, 1, 2, 3) + (uint4)(get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y)) / (uint4)HEIGHT_GEMM3D; + zin = min(DEPTH_GEMM3D - 1, zin); + + // Add offset due to the cross plane paddings + zin *= (src_cross_plane_pad * src0_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply src0_stride_z by DEPTH_GEMM3D + src_addr.s0 += get_global_id(2) * src0_stride_z * DEPTH_GEMM3D; + +#else // defined(REINTERPRET_INPUT_AS_3D) + + // Add offset for batched GEMM + src_addr.s0 += get_global_id(2) * src0_stride_z; + +#endif // defined(REINTERPRET_INPUT_AS_3D) + +#if defined(MATRIX_B_DEPTH) + // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 + src_addr.s1 += (get_global_id(2) % MATRIX_B_DEPTH) * src1_stride_z; +#else // defined(MATRIX_B_DEPTH) + src_addr.s1 += get_global_id(2) * src1_stride_z; +#endif // defined(MATRIX_B_DEPTH) + int end_row_vec_a = src_addr.s0 + COLS_A; uint acc00 = 0; @@ -1075,23 +1451,72 @@ __kernel void gemmlowp_mm_bifrost(IMAGE_DECLARATION(src0), #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 } + const int z = get_global_id(2); + // Compute destination address Image dst = CONVERT_TO_IMAGE_STRUCT(dst); +#if defined(REINTERPRET_OUTPUT_AS_3D) + // Since we store a 2D output tile in a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zout) is calculated dividing M (get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y) by HEIGHT_GEMM3D + uint8 zout = ((uint8)(0, 1, 2, 3, 4, 5, 6, 7) + (uint8)(get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y)) / (uint8)HEIGHT_GEMM3D; + zout = min(DEPTH_GEMM3D - 1, zout); + + // Add offset due to the cross plane paddings + zout *= (dst_cross_plane_pad * dst_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply dst_stride_z by DEPTH_GEMM3D + dst.ptr += z * dst_stride_z * DEPTH_GEMM3D; + + // Store the result + vstore4((int4)(acc00, acc01, acc02, acc03), 0, (__global int *)(dst.ptr + 0 * dst_stride_y + zout.s0)); +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 + vstore4((int4)(acc10, acc11, acc12, acc13), 0, (__global int *)(dst.ptr + 1 * dst_stride_y + zout.s1)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 + vstore4((int4)(acc20, acc21, acc22, acc23), 0, (__global int *)(dst.ptr + 2 * dst_stride_y + zout.s2)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 + vstore4((int4)(acc30, acc31, acc32, acc33), 0, (__global int *)(dst.ptr + 3 * dst_stride_y + zout.s3)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 + vstore4((int4)(acc40, acc41, acc42, acc43), 0, (__global int *)(dst.ptr + 4 * dst_stride_y + zout.s4)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 + +#else // defined(REINTERPRET_OUTPUT_AS_3D) + // Add offset for batched GEMM + dst.ptr += z * dst_stride_z; + // Store the result - vstore4((int4)(acc00, acc01, acc02, acc03), 0, (__global int *)(offset(&dst, 0, 0))); + vstore4((int4)(acc00, acc01, acc02, acc03), 0, (__global int *)(dst.ptr + 0 * dst_stride_y)); #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 - vstore4((int4)(acc10, acc11, acc12, acc13), 0, (__global int *)(offset(&dst, 0, 1))); + vstore4((int4)(acc10, acc11, acc12, acc13), 0, (__global int *)(dst.ptr + 1 * dst_stride_y)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 - vstore4((int4)(acc20, acc21, acc22, acc23), 0, (__global int *)(offset(&dst, 0, 2))); + vstore4((int4)(acc20, acc21, acc22, acc23), 0, (__global int *)(dst.ptr + 2 * dst_stride_y)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 - vstore4((int4)(acc30, acc31, acc32, acc33), 0, (__global int *)(offset(&dst, 0, 3))); + vstore4((int4)(acc30, acc31, acc32, acc33), 0, (__global int *)(dst.ptr + 3 * dst_stride_y)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 - vstore4((int4)(acc40, acc41, acc42, acc43), 0, (__global int *)(offset(&dst, 0, 4))); + vstore4((int4)(acc40, acc41, acc42, acc43), 0, (__global int *)(dst.ptr + 4 * dst_stride_y)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 +#endif // defined(REINTERPRET_OUTPUT_AS_3D) } #if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) @@ -1099,6 +1524,13 @@ __kernel void gemmlowp_mm_bifrost(IMAGE_DECLARATION(src0), * * @attention The number of matrix A columns needs to be passed at compile time using -DCOLS_A * + * @note In case the input or output have to be reinterpreted as a 3D tensor, the following information must be passed at compile time: + * -# REINTERPRET_INPUT_AS_3D: To reinterpret the input as 3D + * -# REINTERPRET_OUTPUT_AS_3D: To reinterpret the output as 3D + * -# HEIGHT_GEMM3D: The height of the output in case it has to be reinterpreted as a 3D tensor. + * -# DEPTH_GEMM3D: The depth of the output in case it has to be reinterpreted as a 3D tensor + * (HEIGHT_GEMM3D * DEPTH_GEMM3D) = columns matrix A NOT reshaped + * * @param[in] src0_ptr Pointer to the source matrix. Supported data type: QASYMM8 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes) @@ -1117,10 +1549,27 @@ __kernel void gemmlowp_mm_bifrost(IMAGE_DECLARATION(src0), * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes) * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix + * @param[in] src0_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] src1_stride_z Stride of the source matrix in Z dimension (in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] src_cross_plane_pad (Optional) Bottom paddings in unit of elements for the input tensor (only if defined REINTERPRET_INPUT_AS_3D) + * @param[in] dst_cross_plane_pad (Optional) Bottom paddings in unit of elements for the output tensor (only if defined REINTERPRET_OUTPUT_AS_3D) */ __kernel void gemmlowp_mm_bifrost_dot8(IMAGE_DECLARATION(src0), IMAGE_DECLARATION(src1), - IMAGE_DECLARATION(dst)) + IMAGE_DECLARATION(dst), + uint src0_stride_z, + uint src1_stride_z, + uint dst_stride_z +#if defined(REINTERPRET_INPUT_AS_3D) + , + uint src_cross_plane_pad +#endif // REINTERPRET_INPUT_AS_3D +#if defined(REINTERPRET_OUTPUT_AS_3D) + , + uint dst_cross_plane_pad +#endif // REINTERPRET_OUTPUT_AS_3D) + ) { int idx = get_global_id(0) * NUM_ELEMS_PROCESSED_PER_THREAD_X; @@ -1133,6 +1582,47 @@ __kernel void gemmlowp_mm_bifrost_dot8(IMAGE_DECLARATION(src0), // Update address for the matrix B src_addr.s1 += idx; +#if defined(REINTERPRET_INPUT_AS_3D) + // Since we load a 2D input tile from a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zin) is calculated dividing M (get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y) by HEIGHT_GEMM3D + uint4 zin = ((uint4)(0, 1, 2, 3) + (uint4)(get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y)) / (uint4)HEIGHT_GEMM3D; + zin = min(DEPTH_GEMM3D - 1, zin); + + // Add offset due to the cross plane paddings + zin *= (src_cross_plane_pad * src0_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply src0_stride_z by DEPTH_GEMM3D + src_addr.s0 += get_global_id(2) * src0_stride_z * DEPTH_GEMM3D; + +#else // defined(REINTERPRET_INPUT_AS_3D) + + // Add offset for batched GEMM + src_addr.s0 += get_global_id(2) * src0_stride_z; + +#endif // defined(REINTERPRET_INPUT_AS_3D) + +#if defined(MATRIX_B_DEPTH) + // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 + src_addr.s1 += (get_global_id(2) % MATRIX_B_DEPTH) * src1_stride_z; +#else // defined(MATRIX_B_DEPTH) + src_addr.s1 += get_global_id(2) * src1_stride_z; +#endif // defined(MATRIX_B_DEPTH) + int end_row_vec_a = src_addr.s0 + COLS_A; uint acc00 = 0; @@ -1321,23 +1811,72 @@ __kernel void gemmlowp_mm_bifrost_dot8(IMAGE_DECLARATION(src0), #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 } + const int z = get_global_id(2); + // Compute destination address Image dst = CONVERT_TO_IMAGE_STRUCT(dst); +#if defined(REINTERPRET_OUTPUT_AS_3D) + // Since we store a 2D output tile in a 3D tensor, we need to check when the plane changes across the z dimension + // in order to take into account the presence of possible cross plane paddings + // + // | | + // | plane0 | + // | | + // |__________________| + // |******************| + // | cross_plane_pad | + // |******************| + // | | + // | plane1 | + // | | + // |__________________| + + // The plane (zout) is calculated dividing M (get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y) by HEIGHT_GEMM3D + uint8 zout = ((uint8)(0, 1, 2, 3, 4, 5, 6, 7) + (uint8)(get_global_id(1) * NUM_ELEMS_PROCESSED_PER_THREAD_Y)) / (uint8)HEIGHT_GEMM3D; + zout = min(DEPTH_GEMM3D - 1, zout); + + // Add offset due to the cross plane paddings + zout *= (dst_cross_plane_pad * dst_stride_y); + + // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we + // multiply dst_stride_z by DEPTH_GEMM3D + dst.ptr += z * dst_stride_z * DEPTH_GEMM3D; + // Store the result - vstore4((int4)(acc00, acc01, acc02, acc03), 0, (__global int *)(offset(&dst, 0, 0))); + vstore4((int4)(acc00, acc01, acc02, acc03), 0, (__global int *)(dst.ptr + 0 * dst_stride_y + zout.s0)); #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 - vstore4((int4)(acc10, acc11, acc12, acc13), 0, (__global int *)(offset(&dst, 0, 1))); + vstore4((int4)(acc10, acc11, acc12, acc13), 0, (__global int *)(dst.ptr + 1 * dst_stride_y + zout.s1)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 - vstore4((int4)(acc20, acc21, acc22, acc23), 0, (__global int *)(offset(&dst, 0, 2))); + vstore4((int4)(acc20, acc21, acc22, acc23), 0, (__global int *)(dst.ptr + 2 * dst_stride_y + zout.s2)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 - vstore4((int4)(acc30, acc31, acc32, acc33), 0, (__global int *)(offset(&dst, 0, 3))); + vstore4((int4)(acc30, acc31, acc32, acc33), 0, (__global int *)(dst.ptr + 3 * dst_stride_y + zout.s3)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 #if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 - vstore4((int4)(acc40, acc41, acc42, acc43), 0, (__global int *)(offset(&dst, 0, 4))); + vstore4((int4)(acc40, acc41, acc42, acc43), 0, (__global int *)(dst.ptr + 4 * dst_stride_y + zout.s4)); #endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 + +#else // defined(REINTERPRET_OUTPUT_AS_3D) + // Add offset for batched GEMM + dst.ptr += z * dst_stride_z; + + // Store the result + vstore4((int4)(acc00, acc01, acc02, acc03), 0, (__global int *)(dst.ptr + 0 * dst_stride_y)); +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 + vstore4((int4)(acc10, acc11, acc12, acc13), 0, (__global int *)(dst.ptr + 1 * dst_stride_y)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 1 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 + vstore4((int4)(acc20, acc21, acc22, acc23), 0, (__global int *)(dst.ptr + 2 * dst_stride_y)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 2 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 + vstore4((int4)(acc30, acc31, acc32, acc33), 0, (__global int *)(dst.ptr + 3 * dst_stride_y)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 3 +#if NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 + vstore4((int4)(acc40, acc41, acc42, acc43), 0, (__global int *)(dst.ptr + 4 * dst_stride_y)); +#endif // NUM_ELEMS_PROCESSED_PER_THREAD_Y > 4 +#endif // defined(REINTERPRET_OUTPUT_AS_3D) } #endif // defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) @@ -1480,26 +2019,26 @@ __kernel void gemmlowp_matrix_b_reduction(TENSOR3D_DECLARATION(src), * (sum_row[i] * B_OFFSET) + * (K_OFFSET) * - * @param[in] mm_result_ptr Pointer to the source tensor. Supported data type: S32 - * @param[in] mm_result_stride_x Stride of the source tensor in X dimension (in bytes) - * @param[in] mm_result_step_x mm_result_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] mm_result_stride_y Stride of the source tensor in Y dimension (in bytes) - * @param[in] mm_result_step_y mm_result_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] mm_result_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] mm_result_step_z mm_result_stride_z * number of elements along Z processed per workitem(in bytes) - * @param[in] mm_result_offset_first_element_in_bytes The offset of the first element in the source tensor - * @param[in] sum_col_result_ptr Pointer to the source tensor. Supported data type: same as @p mm_result_ptr - * @param[in] sum_col_result_stride_x Stride of the source tensor in X dimension (in bytes) - * @param[in] sum_col_result_step_x sum_col_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] sum_col_result_stride_y Stride of the source tensor in Y dimension (in bytes) - * @param[in] sum_col_result_step_y sum_col_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] sum_col_result_offset_first_element_in_bytes The offset of the first element in the source tensor - * @param[in] sum_row_result_ptr Pointer to the source tensor. Supported data type: same as @p mm_result_ptr - * @param[in] sum_row_result_stride_x Stride of the source tensor in X dimension (in bytes) - * @param[in] sum_row_result_step_x sum_row_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] sum_row_result_stride_y Stride of the source tensor in Y dimension (in bytes) - * @param[in] sum_row_result_step_y sum_row_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] sum_row_result_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] mm_result_ptr Pointer to the source tensor. Supported data type: S32 + * @param[in] mm_result_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] mm_result_step_x mm_result_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] mm_result_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] mm_result_step_y mm_result_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] mm_result_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] mm_result_step_z mm_result_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] mm_result_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] sum_col_ptr Pointer to the source tensor. Supported data type: same as @p mm_result_ptr + * @param[in] sum_col_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] sum_col_step_x sum_col_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] sum_col_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] sum_col_step_y sum_col_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] sum_col_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] sum_row_ptr Pointer to the source tensor. Supported data type: same as @p mm_result_ptr + * @param[in] sum_row_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] sum_row_step_x sum_row_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] sum_row_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] sum_row_step_y sum_row_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] sum_row_offset_first_element_in_bytes The offset of the first element in the source tensor */ __kernel void gemmlowp_offset_contribution(TENSOR3D_DECLARATION(mm_result) #if defined(A_OFFSET) @@ -1514,15 +2053,23 @@ __kernel void gemmlowp_offset_contribution(TENSOR3D_DECLARATION(mm_result) { Tensor3D mm_result = CONVERT_TO_TENSOR3D_STRUCT(mm_result); + const int y = get_global_id(1); + const int z = get_global_id(2); + int4 a_offset_s32 = (int4)0; int4 b_offset_s32 = (int4)0; + int batch_id = z; +#if defined(DEPTH_INPUT3D) + batch_id /= (int)DEPTH_INPUT3D; +#endif // defined(DEPTH_INPUT3D) + #if defined(A_OFFSET) Image sum_col = CONVERT_TO_IMAGE_STRUCT(sum_col); // Compute the offset contribution due to A_OFFSET #if defined(SUM_COL_HAS_BATCHES) - a_offset_s32 = vload4(0, (__global int *)(sum_col.ptr + get_global_id(2) * sum_col_stride_y)); + a_offset_s32 = vload4(0, (__global int *)(sum_col.ptr + batch_id * sum_col_stride_y)); #else // defined(MATRIX_B_HAS_BATCHES) a_offset_s32 = vload4(0, (__global int *)(sum_col.ptr)); #endif // defined(MATRIX_B_HAS_BATCHES) @@ -1534,7 +2081,11 @@ __kernel void gemmlowp_offset_contribution(TENSOR3D_DECLARATION(mm_result) Image sum_row = CONVERT_TO_IMAGE_STRUCT(sum_row); // Compute the offset contribution due to B_OFFSET - b_offset_s32 = (int4) * (((__global int *)(sum_row.ptr + get_global_id(2) * sum_row_stride_y)) + get_global_id(1)); +#if defined(HEIGHT_INPUT3D) && defined(DEPTH_INPUT3D) + b_offset_s32 = (int4) * (((__global int *)(sum_row.ptr + batch_id * sum_row_stride_y)) + (z % (int)DEPTH_INPUT3D) * (int)HEIGHT_INPUT3D + y); +#else // defined(HEIGHT_INPUT3D) && defined(DEPTH_INPUT3D) + b_offset_s32 = (int4) * (((__global int *)(sum_row.ptr + batch_id * sum_row_stride_y)) + y); +#endif // defined(HEIGHT_INPUT3D) && defined(DEPTH_INPUT3D) b_offset_s32 *= (int4)B_OFFSET; #endif // defined(B_OFFSET) diff --git a/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyKernel.cpp b/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyKernel.cpp index 9adf95fa33..cf66ebd5fe 100644 --- a/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyKernel.cpp +++ b/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyKernel.cpp @@ -59,17 +59,12 @@ Status validate_arguments(const ITensorInfo *input0, const ITensorInfo *input1, { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input0, 1, DataType::QASYMM8); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input0, input1); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(is_interleaved_transposed && reshape_info.reinterpret_input_as_3d(), "The input tensor cannot be reinterpreted as 3D if is_interleaved_transposed is true"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(input1->num_dimensions() > 2 && reshape_info.reinterpret_input_as_3d(), "The input1 tensor cannot have more than 2 dimensions if input0 has to be reinterpreted as 3D"); if(!is_interleaved_transposed) { ARM_COMPUTE_RETURN_ERROR_ON(input0->dimension(0) != input1->dimension(1)); - - if(output->total_size() != 0) - { - ARM_COMPUTE_RETURN_ERROR_ON(input1->dimension(0) != output->dimension(0)); - ARM_COMPUTE_RETURN_ERROR_ON(input0->dimension(1) != output->dimension(1)); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::S32); - } } else { @@ -95,43 +90,82 @@ Status validate_arguments(const ITensorInfo *input0, const ITensorInfo *input1, ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(input0, &tensor_info_reshaped0); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(input1, &tensor_info_reshaped1); + } - if(output->total_size() != 0) - { - ARM_COMPUTE_RETURN_ERROR_ON(output->dimension(0) != static_cast(n)); - ARM_COMPUTE_RETURN_ERROR_ON(output->dimension(1) != static_cast(m)); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::S32); - } + if(output->total_size() != 0) + { + const TensorInfo tensor_info_output = output->clone()->set_tensor_shape(compute_mm_shape(*input0, *input1, is_interleaved_transposed, reshape_info)); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(output, &tensor_info_output); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::S32); } return Status{}; } std::pair validate_and_configure_window(ITensorInfo *input0, ITensorInfo *input1, ITensorInfo *output, bool is_interleaved_transposed, - ElementsProcessed &num_elements_processed) + const GEMMReshapeInfo &reshape_info, ElementsProcessed &num_elements_processed) { unsigned int &num_elems_processed_per_iteration_x = num_elements_processed[0]; unsigned int &num_elems_processed_per_iteration_y = num_elements_processed[1]; + bool reinterpret_input_as_3d = reshape_info.reinterpret_input_as_3d(); + bool reinterpret_output_as_3d = (reshape_info.depth_output_gemm3d() != 1); Window win{}; + Window win_out{}; bool window_changed = false; + // In case both input and output have to be reinterpreted as 3D tensors, + // force reinterpret_input_as_3d and reinterpret_output_as_3d to be false. + if(reinterpret_input_as_3d == reinterpret_output_as_3d) + { + reinterpret_input_as_3d = false; + reinterpret_output_as_3d = false; + } + + // Output tensor auto inizialitation if not yet initialized + auto_init_if_empty(*output, input0->clone()->set_tensor_shape(compute_mm_shape(*input0, *input1, is_interleaved_transposed, reshape_info))); + + TensorInfo tmp_info(*output); + + if(reinterpret_output_as_3d) + { + // Since the output tensor has to be reinterpreted as 3D and the execute window is based on a 2D GEMM, + // the window needs to be constructed on the 2D collapsed version of the tensor + TensorShape tmp_shape(output->tensor_shape()); + tmp_shape.collapse(2U, 1U); + tmp_info.set_tensor_shape(tmp_shape); + } + // Check if the output tensor is a vector. If so,the kernel runs the vector-matrix multiplication if(is_interleaved_transposed) { + // reinterpret_input_as_3d is not supported if is_interleaved_transposed is set + ARM_COMPUTE_ERROR_ON(reshape_info.reinterpret_input_as_3d()); + // Configure kernel window num_elems_processed_per_iteration_x = 4; num_elems_processed_per_iteration_y = 4; - win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x, num_elems_processed_per_iteration_y)); + // Note: bottom paddings are calculated manually as the output can be reinterpreted as 3D tensor + // The only way to set properly the paddings, it is to set those explicitly through the AccessWindowStatic + const int m = reshape_info.m(); + const int bottom_pad = (num_elems_processed_per_iteration_y - (m % num_elems_processed_per_iteration_y)) % num_elems_processed_per_iteration_y; + + win = calculate_max_window(tmp_info, Steps(num_elems_processed_per_iteration_x, num_elems_processed_per_iteration_y)); + win_out = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x, num_elems_processed_per_iteration_y)); AccessWindowRectangle input0_access(input0, 0, 0, num_elems_processed_per_iteration_y, 1, 1.f, 0.25f); - AccessWindowTranspose input1_access(input1, 0, 0, num_elems_processed_per_iteration_x, 1, 0.f, 0.25f); - AccessWindowRectangle output_access(output, 0, 0, num_elems_processed_per_iteration_x, num_elems_processed_per_iteration_y); + AccessWindowStatic input1_access(input1, 0, 0, + ceil_to_multiple(input1->dimension(0), num_elems_processed_per_iteration_x), + ceil_to_multiple(input1->dimension(1), num_elems_processed_per_iteration_y)); + AccessWindowStatic output_access(output, 0, 0, + ceil_to_multiple(output->dimension(0), num_elems_processed_per_iteration_x), + output->dimension(1) + bottom_pad); - window_changed = update_window_and_padding(win, input0_access, input1_access, output_access); + window_changed = update_window_and_padding(win, input0_access, input1_access) || // window used by the execute_window_loop + update_window_and_padding(win_out, output_access); // window used to update the padding requirements of output tensor - output_access.set_valid_region(win, ValidRegion(Coordinates(0, 0), output->tensor_shape())); + output_access.set_valid_region(win_out, ValidRegion(Coordinates(0, 0), output->tensor_shape())); } else { @@ -139,27 +173,42 @@ std::pair validate_and_configure_window(ITensorInfo *input0, ITe num_elems_processed_per_iteration_x = 4; num_elems_processed_per_iteration_y = std::min(static_cast(output->dimension(1)), 5); + // Note: bottom paddings are calculated manually as the output can be reinterpreted as 3D tensor + // The only way to set properly the paddings, it is to set those explicitly through the AccessWindowStatic + const int m = reinterpret_input_as_3d ? input0->tensor_shape()[1] * input0->tensor_shape()[2] : input0->tensor_shape()[1]; + const int bottom_pad = (num_elems_processed_per_iteration_y - (m % num_elems_processed_per_iteration_y)) % num_elems_processed_per_iteration_y; + // Configure window - win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x, num_elems_processed_per_iteration_y)); + win = calculate_max_window(tmp_info, Steps(num_elems_processed_per_iteration_x, num_elems_processed_per_iteration_y)); + win_out = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x, num_elems_processed_per_iteration_y)); - AccessWindowStatic input0_access(input0, 0, 0, input0->dimension(0), ceil_to_multiple(input0->dimension(1), num_elems_processed_per_iteration_y)); - AccessWindowStatic input1_access(input1, 0, 0, ceil_to_multiple(input1->dimension(0), num_elems_processed_per_iteration_x), input1->dimension(1)); - AccessWindowRectangle output_access(output, 0, 0, num_elems_processed_per_iteration_x, num_elems_processed_per_iteration_y); + AccessWindowStatic input0_access(input0, 0, 0, input0->dimension(0), input0->dimension(1) + bottom_pad); + AccessWindowStatic input1_access(input1, 0, 0, ceil_to_multiple(input1->dimension(0), num_elems_processed_per_iteration_x), input1->dimension(1)); + AccessWindowStatic output_access(output, 0, 0, + ceil_to_multiple(output->dimension(0), num_elems_processed_per_iteration_x), + output->dimension(1) + bottom_pad); - window_changed = update_window_and_padding(win, input0_access, input1_access, output_access); + window_changed = update_window_and_padding(win, input0_access, input1_access) || // window used by the execute_window_loop + update_window_and_padding(win_out, output_access); // window used to update the padding requirements of output tensor Coordinates coord; coord.set_num_dimensions(output->num_dimensions()); output_access.set_valid_region(win, ValidRegion(coord, output->tensor_shape())); } + // Collapse along the Z direction + // This collapse needs to be here in order to tune the Z dimension of LWS + Window collapsed = win; + const unsigned int dimension_to_collapse = std::min(static_cast(output->num_dimensions()), 2u); + collapsed = win.collapse(win, dimension_to_collapse); + Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; - return std::make_pair(err, win); + return std::make_pair(err, collapsed); } } // namespace CLGEMMLowpMatrixMultiplyKernel::CLGEMMLowpMatrixMultiplyKernel() - : _input0(nullptr), _input1(nullptr), _output(nullptr) + : _input0(nullptr), _input1(nullptr), _output(nullptr), _slide_matrix_b(true), _reinterpret_input_as_3d(false), _reinterpret_output_as_3d(false) { } @@ -176,9 +225,23 @@ void CLGEMMLowpMatrixMultiplyKernel::configure(const ICLTensor *input0, const IC ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input0->info(), input1->info(), output->info(), is_interleaved_transposed, reshape_info)); - _input0 = input0; - _input1 = input1; - _output = output; + _input0 = input0; + _input1 = input1; + _output = output; + _reinterpret_input_as_3d = reshape_info.reinterpret_input_as_3d(); + _reinterpret_output_as_3d = (reshape_info.depth_output_gemm3d() != 1); + + // In case both input and output have to be reinterpreted as 3D tensors, + // force reinterpret_input_as_3d and reinterpret_output_as_3d to be false. + if(_reinterpret_input_as_3d == _reinterpret_output_as_3d) + { + _reinterpret_input_as_3d = false; + _reinterpret_output_as_3d = false; + } + + // Check if we need to slide the matrix B + const unsigned int num_dimensions_input0 = _reinterpret_input_as_3d ? _input0->info()->num_dimensions() - 1 : _input0->info()->num_dimensions(); + _slide_matrix_b = (_input1->info()->num_dimensions() >= num_dimensions_input0); ElementsProcessed num_elements_processed{}; @@ -186,15 +249,21 @@ void CLGEMMLowpMatrixMultiplyKernel::configure(const ICLTensor *input0, const IC GPUTarget arch_target = get_arch_from_target(get_target()); // Configure kernel window - auto win_config = validate_and_configure_window(input0->info(), input1->info(), output->info(), is_interleaved_transposed, num_elements_processed); + auto win_config = validate_and_configure_window(input0->info(), input1->info(), output->info(), is_interleaved_transposed, reshape_info, num_elements_processed); ARM_COMPUTE_ERROR_THROW_ON(win_config.first); ICLKernel::configure_internal(win_config.second); const bool is_dot8_supported = dot8_supported(CLKernelLibrary::get().get_device()); // Create build options - CLBuildOptions build_opts; std::string kernel_name(" "); + CLBuildOptions build_opts; + build_opts.add_option_if(_reinterpret_input_as_3d, "-DREINTERPRET_INPUT_AS_3D"); + build_opts.add_option_if(_reinterpret_output_as_3d, "-DREINTERPRET_OUTPUT_AS_3D"); + build_opts.add_option_if(_reinterpret_input_as_3d || _reinterpret_output_as_3d, "-DHEIGHT_GEMM3D=" + support::cpp11::to_string(output->info()->dimension(1))); + build_opts.add_option_if(_reinterpret_input_as_3d || _reinterpret_output_as_3d, "-DDEPTH_GEMM3D=" + support::cpp11::to_string(output->info()->dimension(2))); + build_opts.add_option_if(!_slide_matrix_b, "-DMATRIX_B_DEPTH=" + support::cpp11::to_string(input1->info()->dimension(2))); + if(is_interleaved_transposed) { const int mult_transpose1xW_width = reshape_info.mult_transpose1xW_width(); @@ -225,6 +294,8 @@ void CLGEMMLowpMatrixMultiplyKernel::configure(const ICLTensor *input0, const IC // Set config_id for enabling LWS tuning _config_id = "gemmlowp_"; _config_id += (is_interleaved_transposed ? "reshaped_" : ""); + _config_id += (_reinterpret_input_as_3d ? "3di_" : ""); + _config_id += (_reinterpret_output_as_3d ? "3do_" : ""); _config_id += lower_string(string_from_data_type(input0->info()->data_type())); _config_id += "_"; _config_id += support::cpp11::to_string(output->info()->dimension(1)); @@ -242,6 +313,7 @@ Status CLGEMMLowpMatrixMultiplyKernel::validate(const ITensorInfo *input0, const input1->clone().get(), output->clone().get(), is_interleaved_transposed, + reshape_info, num_elements_processed) .first); @@ -253,18 +325,33 @@ void CLGEMMLowpMatrixMultiplyKernel::run(const Window &window, cl::CommandQueue ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); - Window slice = window.first_slice_window_2D(); + Window slice = window.first_slice_window_3D(); Window slice_matrix_b = slice; - slice_matrix_b.set(Window::DimX, Window::Dimension(0, _input1->info()->dimension(0), 1)); - slice_matrix_b.set(Window::DimY, Window::Dimension(0, _input1->info()->dimension(1), 1)); - slice_matrix_b.set(Window::DimZ, Window::Dimension(0, 1, 1)); + slice_matrix_b.set(Window::DimX, Window::Dimension(0, 1, 1)); + slice_matrix_b.set(Window::DimY, Window::Dimension(0, 1, 1)); + + if(_reinterpret_input_as_3d) + { + // Pass bottom paddings to the kernel if the input has to be reinterpreted as 3D tensor + const unsigned int idx0 = 3 * num_arguments_per_2D_tensor() + 3; + const unsigned int total_cross_plane_pad = _input0->info()->padding().top + _input0->info()->padding().bottom; + _kernel.setArg(idx0, static_cast(total_cross_plane_pad)); + } + + if(_reinterpret_output_as_3d) + { + // Pass bottom paddings to the kernel if the output has to be reinterpreted as 3D tensor + const unsigned int idx0 = 3 * num_arguments_per_2D_tensor() + 3 + (_reinterpret_input_as_3d ? 1 : 0); + const unsigned int total_cross_plane_pad = _output->info()->padding().top + _output->info()->padding().bottom; + _kernel.setArg(idx0, static_cast(total_cross_plane_pad)); + } do { Window slice_b = slice; // Don't slice matrix B along the z dimension if matrix B has just 2 dimensions and matrix A more than 2 // This scenario can happen when the the matrix multiplication is used to perform a convolution operation - if(_input1->info()->num_dimensions() < 3) + if(_slide_matrix_b) { slice_b = slice_matrix_b; } @@ -273,7 +360,10 @@ void CLGEMMLowpMatrixMultiplyKernel::run(const Window &window, cl::CommandQueue add_2D_tensor_argument(idx, _input0, slice); add_2D_tensor_argument(idx, _input1, slice_b); add_2D_tensor_argument(idx, _output, slice); + _kernel.setArg(idx++, static_cast(_input0->info()->strides_in_bytes()[2])); + _kernel.setArg(idx++, static_cast(_input1->info()->strides_in_bytes()[2])); + _kernel.setArg(idx++, static_cast(_output->info()->strides_in_bytes()[2])); enqueue(queue, *this, slice, lws_hint()); } - while(window.slide_window_slice_2D(slice)); + while(window.slide_window_slice_3D(slice)); } diff --git a/src/core/CL/kernels/CLGEMMLowpOffsetContributionKernel.cpp b/src/core/CL/kernels/CLGEMMLowpOffsetContributionKernel.cpp index aa954abde1..3888353ee7 100644 --- a/src/core/CL/kernels/CLGEMMLowpOffsetContributionKernel.cpp +++ b/src/core/CL/kernels/CLGEMMLowpOffsetContributionKernel.cpp @@ -62,16 +62,24 @@ Status validate_arguments(const ITensorInfo *mm_result, const ITensorInfo *vecto if(b_offset != 0) { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(vector_sum_row, 1, DataType::S32); - ARM_COMPUTE_RETURN_ERROR_ON(vector_sum_row->dimension(0) != mm_result->dimension(1)); + + // Check if input is a 3D reinterpretation + const bool reinterpret_as_3d = vector_sum_row != nullptr && mm_result->num_dimensions() > 1 && mm_result->tensor_shape().y() != vector_sum_row->tensor_shape().x(); + + // Validate input + ARM_COMPUTE_RETURN_ERROR_ON(reinterpret_as_3d && vector_sum_row->dimension(0) != (mm_result->dimension(1) * mm_result->dimension(2))); + ARM_COMPUTE_RETURN_ERROR_ON(!reinterpret_as_3d && vector_sum_row != nullptr && vector_sum_row->dimension(0) != mm_result->dimension(1)); TensorShape output_shape = mm_result->tensor_shape(); if(output_shape.num_dimensions() > 1) { + const unsigned int output_batch_idx = reinterpret_as_3d ? 3 : 2; + TensorShape vector_sum_row_shape = vector_sum_row->tensor_shape(); vector_sum_row_shape.collapse_from(1); - output_shape.collapse_from(2); + output_shape.collapse_from(output_batch_idx); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(vector_sum_row_shape[1] != output_shape[2], + ARM_COMPUTE_RETURN_ERROR_ON_MSG(vector_sum_row_shape[1] != output_shape[output_batch_idx], "mm_result tensor must have the same number of batches of output tensor"); if(a_offset != 0) @@ -98,20 +106,17 @@ std::pair validate_and_configure_window(ITensorInfo *mm_result, Window win = calculate_max_window(*mm_result, Steps(num_elems_processed_per_iteration)); AccessWindowHorizontal mm_result_access(mm_result, 0, num_elems_processed_per_iteration); - window_changed = window_changed || update_window_and_padding(win, - mm_result_access); + window_changed = window_changed || update_window_and_padding(win, mm_result_access); if(a_offset != 0) { AccessWindowHorizontal vector_sum_col_access(vector_sum_col, 0, num_elems_processed_per_iteration); - window_changed = window_changed || update_window_and_padding(win, - vector_sum_col_access); + window_changed = window_changed || update_window_and_padding(win, vector_sum_col_access); } if(b_offset != 0) { AccessWindowStatic vector_sum_row_access(vector_sum_row, 0, 0, vector_sum_row->dimension(0), 0); // NOLINT - window_changed = window_changed || update_window_and_padding(win, - vector_sum_row_access); + window_changed = window_changed || update_window_and_padding(win, vector_sum_row_access); } Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; @@ -137,6 +142,11 @@ void CLGEMMLowpOffsetContributionKernel::configure(ICLTensor *mm_result, const I _vector_sum_row = vector_sum_row; _mm_result = mm_result; + // Check if input is a 3D reinterpretation + const bool reinterpret_as_3d = vector_sum_row != nullptr + && mm_result->info()->num_dimensions() > 1 + && mm_result->info()->tensor_shape().y() != vector_sum_row->info()->tensor_shape().x(); + // Set the arguments to pass at compile time CLBuildOptions build_opts; @@ -149,6 +159,8 @@ void CLGEMMLowpOffsetContributionKernel::configure(ICLTensor *mm_result, const I // If b_offset == 0, vector_sum_row can be a nullptr build_opts.add_option_if(b_offset != 0, "-DB_OFFSET=" + support::cpp11::to_string(b_offset)); build_opts.add_option("-DK_OFFSET=" + support::cpp11::to_string(a_offset * b_offset * k)); + build_opts.add_option_if(reinterpret_as_3d, "-DHEIGHT_INPUT3D=" + support::cpp11::to_string(mm_result->info()->dimension(1))); + build_opts.add_option_if(reinterpret_as_3d, "-DDEPTH_INPUT3D=" + support::cpp11::to_string(mm_result->info()->dimension(2))); // Create kernel _kernel = static_cast(CLKernelLibrary::get().create_kernel("gemmlowp_offset_contribution", build_opts.options())); @@ -194,11 +206,13 @@ void CLGEMMLowpOffsetContributionKernel::run(const Window &window, cl::CommandQu // Set window for vector_sum_col Window win_vector_sum_col = slice; win_vector_sum_col.set(Window::DimY, Window::Dimension(0, 0, 0)); + win_vector_sum_col.set(Window::DimZ, Window::Dimension(0, 0, 0)); // Set window for vector_sum_row Window win_vector_sum_row = slice; win_vector_sum_row.set(Window::DimX, Window::Dimension(0, 0, 0)); win_vector_sum_row.set(Window::DimY, Window::Dimension(0, 0, 0)); + win_vector_sum_col.set(Window::DimZ, Window::Dimension(0, 0, 0)); do { diff --git a/src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp b/src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp index cd26cd1597..9cf5d1fb6a 100644 --- a/src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp +++ b/src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp @@ -196,8 +196,8 @@ void CLGEMMLowpMatrixBReductionKernel::run(const Window &window, cl::CommandQueu Window slice_out = collapsed.first_slice_window_2D(); Window slice_in = slice_out; - slice_in.set(Window::DimY, Window::Dimension(0, 1, 1)); - slice_in.set(Window::DimZ, Window::Dimension(0, 1, 1)); + slice_in.set(Window::DimY, Window::Dimension(0, 0, 0)); + slice_in.set(Window::DimZ, Window::Dimension(0, 0, 0)); do { diff --git a/src/runtime/CL/functions/CLGEMM.cpp b/src/runtime/CL/functions/CLGEMM.cpp index 9dbfd3e153..821464e3b3 100644 --- a/src/runtime/CL/functions/CLGEMM.cpp +++ b/src/runtime/CL/functions/CLGEMM.cpp @@ -72,8 +72,18 @@ inline bool is_interleaved_transposed(int m, int n, int k, DataType data_type, b } // namespace CLGEMM::CLGEMM(std::shared_ptr memory_manager) - : _memory_group(std::move(memory_manager)), _interleave_kernel(), _transpose_kernel(), _mm_kernel(), _ma_kernel(), _tmp_a(), _tmp_b(), _original_b(nullptr), _is_interleaved_transposed(false), - _run_addition(false), _reshape_b_only_on_first_run(false), _is_prepared(false) + : _memory_group(std::move(memory_manager)), + _interleave_kernel(), + _transpose_kernel(), + _mm_kernel(), + _ma_kernel(), + _tmp_a(), + _tmp_b(), + _original_b(nullptr), + _is_interleaved_transposed(false), + _run_addition(false), + _reshape_b_only_on_first_run(false), + _is_prepared(false) { } @@ -146,8 +156,9 @@ void CLGEMM::configure(const ICLTensor *a, const ICLTensor *b, const ICLTensor * } // Configure and tune matrix multiply kernel - _mm_kernel.configure(matrix_a, matrix_b, output, alpha, _is_interleaved_transposed, GEMMReshapeInfo(m, n, k, mult_transpose1xW_width, mult_interleave4x4_height, depth_output_gemm3d, - reinterpret_input_as_3d)); + _mm_kernel.configure(matrix_a, matrix_b, output, alpha, _is_interleaved_transposed, GEMMReshapeInfo(m, n, k, + mult_transpose1xW_width, mult_interleave4x4_height, + depth_output_gemm3d, reinterpret_input_as_3d)); CLScheduler::get().tune_kernel_static(_mm_kernel); if(_is_interleaved_transposed) diff --git a/src/runtime/CL/functions/CLGEMMConvolutionLayer.cpp b/src/runtime/CL/functions/CLGEMMConvolutionLayer.cpp index c9daea4169..bd5e969921 100644 --- a/src/runtime/CL/functions/CLGEMMConvolutionLayer.cpp +++ b/src/runtime/CL/functions/CLGEMMConvolutionLayer.cpp @@ -130,9 +130,10 @@ Status CLGEMMConvolutionLayer::validate_mm(const ITensorInfo *input, const ITens { const bool is_quantized = is_data_type_quantized_asymmetric(input->data_type()); - const GEMMInfo &gemm_info = GEMMInfo(false, false, true /* Reshape weights only for the first run */, gemm_3d_depth, skip_im2col /* Reinterpret the input as 3D if im2col is skipped */); if(is_quantized) { + const GEMMInfo &gemm_info = GEMMInfo(false, false, true /* Reshape weights only for the first run */); + // Since we need negative offsets for computing convolution, we need to change QuantizationInfo() // Extract and negate input and weights offset const QuantizationInfo input_quantization_info = input->quantization_info(); @@ -148,6 +149,8 @@ Status CLGEMMConvolutionLayer::validate_mm(const ITensorInfo *input, const ITens } else { + const GEMMInfo &gemm_info = GEMMInfo(false, false, true /* Reshape weights only for the first run */, gemm_3d_depth, skip_im2col /* Reinterpret the input as 3D if im2col is skipped */); + // Perform validation step on Matrix multiply function return CLGEMM::validate(input, weights, nullptr, output, 1.0f, 0.0f, gemm_info); } diff --git a/src/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp b/src/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp index 763ebced83..1d6f343cb2 100644 --- a/src/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp +++ b/src/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp @@ -107,15 +107,23 @@ void CLGEMMLowpMatrixMultiplyCore::configure(const ICLTensor *a, const ICLTensor // Arguments used by GEMMReshapeInfo // If we pass the matrix A and matrix B reshaped to CLGEMMMatrixMultiplyKernel, we need to pass m, n, k, mult_transpose1xW_width and mult_interleave4x4_height to CLGEMMReshapeInfo // in order to know how the matrices have been reshaped + bool reinterpret_input_as_3d = gemm_info.reinterpret_input_as_3d(); const int m = a->info()->dimension(1); const int n = b->info()->dimension(0); const int k = a->info()->dimension(0); + const int depth_output_gemm3d = gemm_info.depth_output_gemm3d(); constexpr int mult_transpose1xW_width = 1; constexpr int mult_interleave4x4_height = 1; // Check if we need to reshape the matrix A and matrix B _is_interleaved_transposed = is_interleaved_transposed(m, n, k, _reshape_b_only_on_first_run, gpu_target); + // if _is_interleaved_transposed is set, force reinterpret_input_as_3d to be false as the output of CLGEMMInterleaveKernel will be 2D + if(_is_interleaved_transposed) + { + reinterpret_input_as_3d = false; + } + if(_is_interleaved_transposed) { matrix_a = &_tmp_a; @@ -128,14 +136,15 @@ void CLGEMMLowpMatrixMultiplyCore::configure(const ICLTensor *a, const ICLTensor } // Configure interleave kernel - _mtx_a_reshape_kernel.configure(a, &_tmp_a, mult_interleave4x4_height); + _mtx_a_reshape_kernel.configure(a, &_tmp_a, mult_interleave4x4_height, gemm_info.reinterpret_input_as_3d()); // Configure transpose kernel _mtx_b_reshape_kernel.configure(b, &_tmp_b, mult_transpose1xW_width); } - // Configure matrix multiply kernel - _mm_kernel.configure(matrix_a, matrix_b, output, _is_interleaved_transposed, GEMMReshapeInfo(m, n, k, mult_transpose1xW_width, mult_interleave4x4_height)); + _mm_kernel.configure(matrix_a, matrix_b, output, _is_interleaved_transposed, GEMMReshapeInfo(m, n, k, + mult_transpose1xW_width, mult_interleave4x4_height, + depth_output_gemm3d, reinterpret_input_as_3d)); // Initialize matrix B reduction kernel only if _a_offset is not equal to 0 if(_a_offset != 0) @@ -191,28 +200,30 @@ Status CLGEMMLowpMatrixMultiplyCore::validate(const ITensorInfo *a, const ITenso ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(a, 1, DataType::QASYMM8); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::S32); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(a, b); - ARM_COMPUTE_RETURN_ERROR_ON_MSG((a)->dimension(0) != (b)->dimension(1), - "The product AB is defined only if the number of columns in A is equal to the number of rows in B"); - ARM_COMPUTE_RETURN_ERROR_ON_MSG((a)->dimension(1) != (output)->dimension(1), - "The output matrix must have the same number of rows as the matrix A"); - ARM_COMPUTE_RETURN_ERROR_ON_MSG((b)->dimension(0) != (output)->dimension(0), - "The output matrix must have the same number of columns as the matrix B"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(gemm_info.is_a_reshaped(), "Matrix A already reshaped is not supported"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(gemm_info.is_b_reshaped(), "Matrix B already reshaped is not supported"); int32_t a_offset = a->quantization_info().offset; int32_t b_offset = b->quantization_info().offset; - const int m = a->dimension(1); - const int n = b->dimension(0); - const int k = a->dimension(0); - constexpr int mult_transpose1xW_width = 1; - constexpr int mult_interleave4x4_height = 1; - const int depth_output_gemm3d = gemm_info.depth_output_gemm3d(); - const GEMMReshapeInfo reshape_info(m, n, k, mult_transpose1xW_width, mult_interleave4x4_height, depth_output_gemm3d); + const int m = a->dimension(1); + const int n = b->dimension(0); + const int k = a->dimension(0); + constexpr int mult_transpose1xW_width = 1; + constexpr int mult_interleave4x4_height = 1; + bool reinterpret_input_as_3d = gemm_info.reinterpret_input_as_3d(); + const int depth_output_gemm3d = gemm_info.depth_output_gemm3d(); bool reshape_matrices = is_interleaved_transposed(m, n, k, gemm_info.reshape_b_only_on_first_run(), CLScheduler::get().target()); + // if reshape_matrices is set, force reinterpret_input_as_3d to be false as the output of CLGEMMInterleaveKernel will be 2D + if(reshape_matrices) + { + reinterpret_input_as_3d = false; + } + + const GEMMReshapeInfo reshape_info = GEMMReshapeInfo(m, n, k, mult_transpose1xW_width, mult_interleave4x4_height, depth_output_gemm3d, reinterpret_input_as_3d); + if(reshape_matrices) { TensorInfo info_a(compute_interleaved_shape(*a, mult_interleave4x4_height, gemm_info.reinterpret_input_as_3d()), 1, a->data_type()); diff --git a/tests/datasets/LargeGEMMLowpDataset.h b/tests/datasets/LargeGEMMLowpDataset.h index 5c0230e262..65cb742ead 100644 --- a/tests/datasets/LargeGEMMLowpDataset.h +++ b/tests/datasets/LargeGEMMLowpDataset.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,6 +52,26 @@ public: add_config(TensorShape(941U, 1011U), TensorShape(623U, 941U), TensorShape(623U, 1011U), -9, 1); } }; +class LargeGEMMLowpOutput3DDataset final : public GEMMLowpDataset +{ +public: + LargeGEMMLowpOutput3DDataset() + { + add_config(TensorShape(923U, 429U), TensorShape(871U, 923U), TensorShape(871U, 143U, 3U), 0, 0); + add_config(TensorShape(681U, 1025U), TensorShape(213U, 681U), TensorShape(213U, 205U, 5U), -3, 2); + add_config(TensorShape(364U, 3025U), TensorShape(96U, 364U), TensorShape(96U, 605U, 5U), 2, 3); + } +}; +class LargeGEMMLowpInputOutput3DDataset final : public GEMMLowpDataset +{ +public: + LargeGEMMLowpInputOutput3DDataset() + { + add_config(TensorShape(923U, 143U, 3U), TensorShape(871U, 923U), TensorShape(871U, 143U, 3U), 0, 0); + add_config(TensorShape(681U, 205U, 5U), TensorShape(213U, 681U), TensorShape(213U, 205U, 5U), -2, 5); + add_config(TensorShape(364U, 605U, 5U), TensorShape(96U, 364U), TensorShape(96U, 605U, 5U), 2, 4); + } +}; } // namespace datasets } // namespace test } // namespace arm_compute diff --git a/tests/datasets/SmallGEMMLowpDataset.h b/tests/datasets/SmallGEMMLowpDataset.h index b6651bdb42..40f0c718c6 100644 --- a/tests/datasets/SmallGEMMLowpDataset.h +++ b/tests/datasets/SmallGEMMLowpDataset.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,6 +52,32 @@ public: add_config(TensorShape(32U, 72U), TensorShape(17U, 32U), TensorShape(17U, 72U), -9, 1); } }; +class SmallGEMMLowpOutput3DDataset final : public GEMMLowpDataset +{ +public: + SmallGEMMLowpOutput3DDataset() + { + add_config(TensorShape(21U, 14U), TensorShape(34U, 21U), TensorShape(34U, 7U, 2U), 0, 0); + add_config(TensorShape(31U, 1U), TensorShape(23U, 31U), TensorShape(23U, 1U, 1U), -2, 13); + add_config(TensorShape(38U, 12U), TensorShape(21U, 38U), TensorShape(21U, 4U, 3U), 0, 4); + add_config(TensorShape(32U, 1U), TensorShape(17U, 32U), TensorShape(17U, 1U, 1U), -2, 1); + add_config(TensorShape(16U, 16U), TensorShape(8U, 16U), TensorShape(8U, 8U, 2U), 5, 9); + add_config(TensorShape(16U, 16U, 5U), TensorShape(8U, 16U, 5U), TensorShape(8U, 8U, 2U, 5U), -7, 2); + } +}; +class SmallGEMMLowpInputOutput3DDataset final : public GEMMLowpDataset +{ +public: + SmallGEMMLowpInputOutput3DDataset() + { + add_config(TensorShape(21U, 14U, 13U), TensorShape(34U, 21U), TensorShape(34U, 14U, 13U), 0, 0); + add_config(TensorShape(31U, 1U, 3U), TensorShape(23U, 31U), TensorShape(23U, 1U, 3U), 0, 0); + add_config(TensorShape(38U, 12U, 2U), TensorShape(21U, 38U), TensorShape(21U, 12U, 2U), -2, 13); + add_config(TensorShape(32U, 1U, 4U, 3U), TensorShape(17U, 32U), TensorShape(17U, 1U, 4U, 3U), 0, 4); + add_config(TensorShape(16U, 16U, 3U, 2U), TensorShape(8U, 16U), TensorShape(8U, 16U, 3U, 2U), -2, 0); + add_config(TensorShape(16U, 16U, 5U, 3U), TensorShape(8U, 16U), TensorShape(8U, 16U, 5U, 3U), -9, 1); + } +}; } // namespace datasets } // namespace test } // namespace arm_compute diff --git a/tests/validate_examples/cl_gemm.cpp b/tests/validate_examples/cl_gemm.cpp index e0aefbf359..cdaa33f31a 100644 --- a/tests/validate_examples/cl_gemm.cpp +++ b/tests/validate_examples/cl_gemm.cpp @@ -285,7 +285,7 @@ public: fill(ref_src0, 0); fill(ref_src1, 1); - SimpleTensor ref_tmp_dst = reference::gemmlowp_matrix_multiply_core(ref_src0, ref_src1, offset_src0, offset_src1); + SimpleTensor ref_tmp_dst = reference::gemmlowp_matrix_multiply_core(ref_src0, ref_src1, TensorShape(N, M), offset_src0, offset_src1); if(add_bias) { diff --git a/tests/validation/CL/GEMMLowp.cpp b/tests/validation/CL/GEMMLowp.cpp index 5148a31936..42bb2123bf 100644 --- a/tests/validation/CL/GEMMLowp.cpp +++ b/tests/validation/CL/GEMMLowp.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -47,6 +47,7 @@ TEST_SUITE(CL) TEST_SUITE(GEMMLowp) TEST_SUITE(MatrixMultiplyCore) + using CLGEMMLowpMatrixMultiplyCoreFixture = GEMMLowpMatrixMultiplyCoreValidationFixture; DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, framework::dataset::concat(datasets::SmallGEMMLowpDataset(), datasets::LargeGEMMLowpDataset()), @@ -81,6 +82,33 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMLowpMatrixMultiplyCoreFixture, framework: validate(CLAccessor(_target), _reference); } +TEST_SUITE(Output3D) +using CLGEMMLowpMatrixMultiplyCoreOutput3DFixture = GEMMLowpMatrixMultiplyCoreValidationFixture; +FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMLowpMatrixMultiplyCoreOutput3DFixture, framework::DatasetMode::PRECOMMIT, datasets::SmallGEMMLowpOutput3DDataset()) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMLowpMatrixMultiplyCoreOutput3DFixture, framework::DatasetMode::NIGHTLY, datasets::LargeGEMMLowpOutput3DDataset()) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() // Output3D + +TEST_SUITE(InputOutput3D) +using CLGEMMLowpMatrixMultiplyCoreInputOutput3DFixture = GEMMLowpMatrixMultiplyCoreValidationFixture; +FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMLowpMatrixMultiplyCoreInputOutput3DFixture, framework::DatasetMode::PRECOMMIT, datasets::SmallGEMMLowpInputOutput3DDataset()) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMLowpMatrixMultiplyCoreInputOutput3DFixture, framework::DatasetMode::NIGHTLY, datasets::LargeGEMMLowpInputOutput3DDataset()) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() // InputOutput3D TEST_SUITE_END() // MatrixMultiplyCore TEST_SUITE(OutputStage) diff --git a/tests/validation/fixtures/GEMMLowpAssemblyFixture.h b/tests/validation/fixtures/GEMMLowpAssemblyFixture.h index d6b94a197d..519932f3b2 100644 --- a/tests/validation/fixtures/GEMMLowpAssemblyFixture.h +++ b/tests/validation/fixtures/GEMMLowpAssemblyFixture.h @@ -128,7 +128,7 @@ protected: fill(b, 1, 0, 255); } - return reference::gemmlowp(a, b); + return reference::gemmlowp(a, b, shape_c); } TensorType _target{}; diff --git a/tests/validation/fixtures/GEMMLowpFixture.h b/tests/validation/fixtures/GEMMLowpFixture.h index 06d6be3fa4..73cb8328ea 100644 --- a/tests/validation/fixtures/GEMMLowpFixture.h +++ b/tests/validation/fixtures/GEMMLowpFixture.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -42,7 +42,7 @@ namespace test { namespace validation { -template +template class GEMMLowpMatrixMultiplyCoreValidationFixture : public framework::Fixture { public: @@ -62,8 +62,7 @@ protected: library->fill(tensor, distribution, i); } - TensorType compute_target(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_c, - int32_t a_offset, int32_t b_offset) + TensorType compute_target(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_c, int32_t a_offset, int32_t b_offset) { // Create tensors TensorType a = create_tensor(shape_a, DataType::QASYMM8, 1); @@ -74,8 +73,9 @@ protected: b.info()->set_quantization_info(QuantizationInfo(1.0f / 255, b_offset)); // Create and configure function + // The GEMMinfo includes the values of the depth in case of reinterpreted 3d input/output FunctionType gemmlowp; - gemmlowp.configure(&a, &b, &c); + gemmlowp.configure(&a, &b, &c, GEMMInfo(false, false, false, (reinterpret_output_as_3d ? shape_c[2] : 1), reinterpret_input_as_3d)); ARM_COMPUTE_EXPECT(a.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(b.info()->is_resizable(), framework::LogLevel::ERRORS); @@ -99,18 +99,24 @@ protected: return c; } - SimpleTensor compute_reference(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_c, - int32_t a_offset, int32_t b_offset) + SimpleTensor compute_reference(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_c, int32_t a_offset, int32_t b_offset) { + TensorShape shape_a_to_use = shape_a; + if(reinterpret_input_as_3d) + { + // Collapse the second and third dimension if the input is 3D + shape_a_to_use.collapse(2U, 1U); + } + // Create reference - SimpleTensor a{ shape_a, DataType::QASYMM8, 1 }; + SimpleTensor a{ shape_a_to_use, DataType::QASYMM8, 1 }; SimpleTensor b{ shape_b, DataType::QASYMM8, 1 }; // Fill reference fill(a, 0); fill(b, 1); - return reference::gemmlowp_matrix_multiply_core(a, b, a_offset, b_offset); + return reference::gemmlowp_matrix_multiply_core(a, b, shape_c, a_offset, b_offset); } TensorType _target{}; diff --git a/tests/validation/reference/GEMMLowp.cpp b/tests/validation/reference/GEMMLowp.cpp index 8e41aef46a..9a7e409e8a 100644 --- a/tests/validation/reference/GEMMLowp.cpp +++ b/tests/validation/reference/GEMMLowp.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -98,41 +98,52 @@ void quantize_down_int32_to_uint8_scale_by_fixedpoint(const SimpleTensor *in, } // namespace template -SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, int32_t a_offset, int32_t b_offset) +SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c, int32_t a_offset, int32_t b_offset) { static_assert(std::is_same::type, int32_t>::value, "Only int32_t is allowed for the output"); - TensorShape shape(b.shape()[0], a.shape()[1]); DataType dt = std::is_same::value ? DataType::S32 : DataType::U32; - SimpleTensor c(shape, dt); + SimpleTensor c(shape_c, dt); - const int K = a.shape().x(); - const int b_width = b.shape().x(); - const int rows = c.shape().y(); //M - const int cols = c.shape().x(); //N + const int K = a.shape().x(); + const int M = a.shape().y(); + const int N = b.shape().x(); + const int D = a.shape().z(); // Number of matrices in a batch + + const int a_stride_z = K * M; + // Do not slide the matrix B along the 3rd dimension in case matrix B has less than 3 dimensions + const int b_stride_z = b.shape().num_dimensions() > 2 ? N * K : 0; + const int c_stride_z = N * M; std::vector acc; - acc.resize(cols); + acc.resize(N); - for(int i = 0; i < rows; ++i) + for(int depth = 0; depth < D; ++depth) { - for(int j = 0; j < cols; ++j) - { - acc[j] = 0; - } - for(int k = 0; k < K; ++k) + const int base_addr_a = depth * a_stride_z; + const int base_addr_b = depth * b_stride_z; + const int base_addr_c = depth * c_stride_z; + + for(int i = 0; i < M; ++i) { - const T_out tmp_a = a_offset + static_cast(a[k + i * K]); - for(int j = 0; j < b_width; ++j) + for(int j = 0; j < N; ++j) { - const T_out tmp_b = b_offset + static_cast(b[j + k * b_width]); - const T_out mult_as_int = tmp_a * tmp_b; - acc[j] += mult_as_int; + acc[j] = 0; + } + for(int k = 0; k < K; ++k) + { + const T_out tmp_a = a_offset + static_cast(a[base_addr_a + k + i * K]); + for(int j = 0; j < N; ++j) + { + const T_out tmp_b = b_offset + static_cast(b[base_addr_b + j + k * N]); + const T_out mult_as_int = tmp_a * tmp_b; + acc[j] += mult_as_int; + } + } + for(int j = 0; j < N; ++j) + { + c[base_addr_c + j + i * N] = acc[j]; } - } - for(int j = 0; j < cols; ++j) - { - c[j + i * cols] = acc[j]; } } @@ -141,9 +152,9 @@ SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, c // used to validate assembly kernels which don't know anything about offsets template -SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b) +SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c) { - return gemmlowp_matrix_multiply_core(a, b, 0, 0); + return gemmlowp_matrix_multiply_core(a, b, shape_c, 0, 0); } template @@ -198,10 +209,10 @@ template SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const int32_t max); template SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &a, const SimpleTensor &b, int32_t result_offset, int32_t result_mult_int, int32_t result_shift, int32_t min, int32_t max); -template SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, int32_t a_offset, int32_t b_offset); -template SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, int32_t a_offset, int32_t b_offset); -template SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b); -template SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b); +template SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c, int32_t a_offset, int32_t b_offset); +template SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c, int32_t a_offset, int32_t b_offset); +template SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c); +template SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c); } // namespace reference } // namespace validation } // namespace test diff --git a/tests/validation/reference/GEMMLowp.h b/tests/validation/reference/GEMMLowp.h index a3d0bebe3f..4396155b96 100644 --- a/tests/validation/reference/GEMMLowp.h +++ b/tests/validation/reference/GEMMLowp.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -38,13 +38,13 @@ namespace reference template SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &in, int32_t result_offset, int32_t result_mult_int, int32_t result_shift, int32_t min = 0, int32_t max = 0); template -SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, int32_t a_offset, int32_t b_offset); +SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c, int32_t a_offset, int32_t b_offset); template SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &in, int32_t result_offset, int32_t result_mult_int, int32_t result_shift); template -SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b); +SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c); template SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &in, const SimpleTensor &bias, int32_t result_offset, int32_t result_mult_int, int32_t result_shift, -- cgit v1.2.1