diff --git a/arm_compute/runtime/NEON/functions/NEConvolutionLayer.h b/arm_compute/runtime/NEON/functions/NEConvolutionLayer.h index 00d97aa4a8d..ec95b40ec4e 100644 --- a/arm_compute/runtime/NEON/functions/NEConvolutionLayer.h +++ b/arm_compute/runtime/NEON/functions/NEConvolutionLayer.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021, 2023-2025 Arm Limited. + * Copyright (c) 2018-2021, 2023-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -104,6 +104,7 @@ class NEConvolutionLayer : public IFunction * |QASYMM8 |QSYMM8_PER_CHANNEL |S32 |QASYMM8 | * |QASYMM8_SIGNED |QASYMM8_SIGNED |S32 |QASYMM8_SIGNED | * |QASYMM8_SIGNED |QASYMM8_SIGNED |F32 |F32 | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F16 |F16 | * |QASYMM8_SIGNED |QSYMM8_PER_CHANNEL |S32 |QASYMM8_SIGNED | * * @param[in] input Source tensor. 3 lower dimensions represent a single input [width, height, IFM], @@ -112,10 +113,11 @@ class NEConvolutionLayer : public IFunction * @param[in] weights Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. * Data type supported: Same as @p input, also could be QSYMM8_PER_CHANNEL or QASYMM8_SIGNED if input is QASYMM8/QASYMM8_SIGNED. * @param[in] biases Biases tensor. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. - * Data type supported: Same as @p input, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type and - * for F32 dequantization the bias must be F32. + * Data type supported: Same as @p input, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * For floating-point dequantization the bias must match @p output. * @param[out] output Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with NEWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. @@ -142,12 +144,12 @@ class NEConvolutionLayer : public IFunction * Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32. * @param[in] weights Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. * Data type supported:Same as @p input, also could be QSYMM8_PER_CHANNEL or QASYMM8_SIGNED if input is QASYMM8/QASYMM8_SIGNED. - * Data type supported: Same as @p input, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. * @param[in] biases Biases tensor. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. - * Data type supported: Same as @p input, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type and - * for F32 dequantization the bias must be F32. + * Data type supported: Same as @p input, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * For floating-point dequantization the bias must match @p output. * @param[in] output Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with NEWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. @@ -177,7 +179,8 @@ class NEConvolutionLayer : public IFunction * @param[in] weights Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. * Data type supported:Same as @p input, also could be QSYMM8_PER_CHANNEL or QASYMM8_SIGNED if input is QASYMM8/QASYMM8_SIGNED. * @param[in] output Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with NEWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. diff --git a/arm_compute/runtime/NEON/functions/NEGEMMConv2d.h b/arm_compute/runtime/NEON/functions/NEGEMMConv2d.h index 96acc3f82ff..d9c47210d2f 100644 --- a/arm_compute/runtime/NEON/functions/NEGEMMConv2d.h +++ b/arm_compute/runtime/NEON/functions/NEGEMMConv2d.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, 2024-2025 Arm Limited. + * Copyright (c) 2020-2021, 2024-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -75,6 +75,8 @@ class NEGEMMConv2d : public IFunction * |:--------------|:--------------|:--------------|:--------------| * |QASYMM8 |QASYMM8 |S32 |QASYMM8 | * |QASYMM8_SIGNED |QASYMM8_SIGNED |S32 |QASYMM8_SIGNED | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F32 |F32 | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F16 |F16 | * |F16 |F16 |F16 |F16 | * |F32 |F32 |F32 |F32 | * |BFLOAT16 |BFLOAT16 |BFLOAT16 |BFLOAT16 | @@ -86,8 +88,10 @@ class NEGEMMConv2d : public IFunction * Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * For floating-point dequantization the bias must match @p output. * @param[out] output Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] info Convolution layer descriptor */ void @@ -101,8 +105,10 @@ class NEGEMMConv2d : public IFunction * Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * For floating-point dequantization the bias must match @p output. * @param[in] output Destination tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] info Contains padding and stride information described in @ref PadStrideInfo. * * @return a status diff --git a/arm_compute/runtime/NEON/functions/NEGEMMConvolutionLayer.h b/arm_compute/runtime/NEON/functions/NEGEMMConvolutionLayer.h index 4a5714a4470..0ea53cb061b 100644 --- a/arm_compute/runtime/NEON/functions/NEGEMMConvolutionLayer.h +++ b/arm_compute/runtime/NEON/functions/NEGEMMConvolutionLayer.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2025 Arm Limited. + * Copyright (c) 2017-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -79,6 +79,7 @@ class NEGEMMConvolutionLayer : public IFunction * |QASYMM8 |QSYMM8_PER_CHANNEL |S32 |QASYMM8 | * |QASYMM8_SIGNED |QASYMM8_SIGNED |S32 |QASYMM8_SIGNED | * |QASYMM8_SIGNED |QASYMM8_SIGNED |F32 |F32 | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F16 |F16 | * |QASYMM8_SIGNED |QSYMM8_PER_CHANNEL |S32 |QASYMM8_SIGNED | * * @param[in] input Source tensor. 3 lower dimensions represent a single input [width, height, IFM], @@ -88,9 +89,10 @@ class NEGEMMConvolutionLayer : public IFunction * Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. - * For F32 dequantization the bias must be F32. + * For floating-point dequantization the bias must match @p output. * @param[out] output Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with NEWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. @@ -119,9 +121,10 @@ class NEGEMMConvolutionLayer : public IFunction * Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. - * For F32 dequantization the bias must be F32. + * For floating-point dequantization the bias must match @p output. * @param[in] output Destination tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with NEWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. diff --git a/arm_compute/runtime/experimental/operators/CpuGemmConv2d.h b/arm_compute/runtime/experimental/operators/CpuGemmConv2d.h index a320c66fb07..4a72b881f75 100644 --- a/arm_compute/runtime/experimental/operators/CpuGemmConv2d.h +++ b/arm_compute/runtime/experimental/operators/CpuGemmConv2d.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2025 Arm Limited. + * Copyright (c) 2021-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -78,6 +78,7 @@ class CpuGemmConv2d : public IOperator * |QASYMM8 |QSYMM8_PER_CHANNEL |S32 |QASYMM8 | * |QASYMM8_SIGNED |QASYMM8_SIGNED |S32 |QASYMM8_SIGNED | * |QASYMM8_SIGNED |QASYMM8_SIGNED |F32 |F32 | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F16 |F16 | * |QASYMM8_SIGNED |QSYMM8_PER_CHANNEL |S32 |QASYMM8_SIGNED | * * @param[in] src Source tensor info. 3 lower dimensions represent a single input [width, height, IFM], @@ -87,9 +88,10 @@ class CpuGemmConv2d : public IOperator * Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. - * For F32 dequantization the bias must be F32. + * For floating-point dequantization the bias must match @p dst. * @param[out] dst Destination tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with CpuWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. diff --git a/arm_compute/runtime/experimental/operators/CpuGemmDirectConv2d.h b/arm_compute/runtime/experimental/operators/CpuGemmDirectConv2d.h index 6e43ae9fd0c..402cfdd8c6b 100644 --- a/arm_compute/runtime/experimental/operators/CpuGemmDirectConv2d.h +++ b/arm_compute/runtime/experimental/operators/CpuGemmDirectConv2d.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025 Arm Limited. + * Copyright (c) 2024-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -71,6 +71,8 @@ class CpuGemmDirectConv2d : public IOperator * |:--------------|:--------------|:--------------|:--------------| * |QASYMM8 |QASYMM8 |S32 |QASYMM8 | * |QASYMM8_SIGNED |QASYMM8_SIGNED |S32 |QASYMM8_SIGNED | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F32 |F32 | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F16 |F16 | * |F16 |F16 |F16 |F16 | * |F32 |F32 |F32 |F32 | * |BFLOAT16 |BFLOAT16 |BFLOAT16 |BFLOAT16 | @@ -82,8 +84,10 @@ class CpuGemmDirectConv2d : public IOperator * Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * For floating-point dequantization the bias must match @p dst. * @param[in] dst Destination tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] info Contains padding and stride information described in @ref PadStrideInfo. */ void configure(const ITensorInfo *src, diff --git a/docs/user_guide/operator_list.dox b/docs/user_guide/operator_list.dox index 283a45e8fcc..2753e97dbda 100644 --- a/docs/user_guide/operator_list.dox +++ b/docs/user_guide/operator_list.dox @@ -611,6 +611,7 @@ where N = batches, C = channels, H = height, W = width, D = depth QASYMM8QSYMM8_PER_CHANNELS32QASYMM8 QASYMM8_SIGNEDQASYMM8_SIGNEDS32QASYMM8_SIGNED QASYMM8_SIGNEDQASYMM8_SIGNEDF32F32 + QASYMM8_SIGNEDQASYMM8_SIGNEDF16F16 QASYMM8_SIGNEDQSYMM8_PER_CHANNELS32QASYMM8_SIGNED @@ -1687,6 +1688,8 @@ where N = batches, C = channels, H = height, W = width, D = depth src0src1src2dst QASYMM8QASYMM8S32QASYMM8 QASYMM8_SIGNEDQASYMM8_SIGNEDS32QASYMM8_SIGNED + QASYMM8_SIGNEDQASYMM8_SIGNEDF32F32 + QASYMM8_SIGNEDQASYMM8_SIGNEDF16F16 F16F16F16F16 F32F32F32F32 BFLOAT16BFLOAT16BFLOAT16BFLOAT16 @@ -1714,6 +1717,7 @@ where N = batches, C = channels, H = height, W = width, D = depth QASYMM8QSYMM8_PER_CHANNELS32QASYMM8 QASYMM8_SIGNEDQASYMM8_SIGNEDS32QASYMM8_SIGNED QASYMM8_SIGNEDQASYMM8_SIGNEDF32F32 + QASYMM8_SIGNEDQASYMM8_SIGNEDF16F16 QASYMM8_SIGNEDQSYMM8_PER_CHANNELS32QASYMM8_SIGNED diff --git a/src/core/NEON/kernels/arm_gemm/gemm_interleaved.hpp b/src/core/NEON/kernels/arm_gemm/gemm_interleaved.hpp index 5a2b83f527a..69ecac3dc9e 100644 --- a/src/core/NEON/kernels/arm_gemm/gemm_interleaved.hpp +++ b/src/core/NEON/kernels/arm_gemm/gemm_interleaved.hpp @@ -72,6 +72,33 @@ class kernel_and_merge { Tab *acc_buff); }; +template +auto run_dequantized_integrated_kernel( + strategy &strat, const Tlo *a_ptr, const Tro *b_panel, Tr *c_ptr, int ldc, + unsigned int m_size, unsigned int n_size, int kern_k, const int32_t *offset_col_bias, + const DequantizeFloat &dq, const Tr *offset_bias, const Activation &act, bool accumulate, + Tab *acc_buff, const int32_t *row_sum, int) + -> decltype(strat.kernel(a_ptr, b_panel, c_ptr, ldc, m_size, n_size, kern_k, + offset_col_bias, dq, offset_bias, act, accumulate, acc_buff, + row_sum, kern_k), + void()) +{ + strat.kernel(a_ptr, b_panel, c_ptr, ldc, m_size, n_size, kern_k, + offset_col_bias, dq, offset_bias, act, accumulate, acc_buff, + row_sum, kern_k); +} + +template +void run_dequantized_integrated_kernel( + strategy &strat, const Tlo *a_ptr, const Tro *b_panel, Tr *c_ptr, int ldc, + unsigned int m_size, unsigned int n_size, int kern_k, const int32_t *offset_col_bias, + const DequantizeFloat &dq, const Tr *offset_bias, const Activation &act, bool accumulate, + Tab *acc_buff, const int32_t *, long) +{ + strat.kernel(a_ptr, b_panel, c_ptr, ldc, m_size, n_size, kern_k, + offset_col_bias, dq, offset_bias, act, accumulate, acc_buff); +} + // Run a kernel and call the separate merge step template<> template @@ -281,14 +308,19 @@ void kernel_and_merge::run( offset_bias = bias + n_0; } - strat.kernel(// A and B pointers are just the packed panels. - a_ptr, b_panel, - // Provide relevant part of output array and row stride. - c_ptr ? (c_ptr + m_0 * ldc + n_0) : nullptr, ldc, - // M, N, K sizes - m_max-m_0, n_max - n_0, kern_k, - // Bias, activation, accumulation. Need to offset the bias as needed. - offset_col_bias, dq, offset_bias, act, accumulate, acc_buff); + // When b_offset != 0, row sums of A are packed at the end of the A panel + // (appended by the quantized PrepareA transform with multiplier=1). Read them + // to pass to dequantize_block_32 for per-row offset correction. + const int32_t *row_sum = nullptr; + if (dq.b_offset != 0) { + row_sum = reinterpret_cast(a_ptr + strategy::out_height() * kern_k); + } + + run_dequantized_integrated_kernel( + strat, a_ptr, b_panel, + c_ptr ? (c_ptr + m_0 * ldc + n_0) : nullptr, ldc, + m_max - m_0, n_max - n_0, kern_k, offset_col_bias, dq, offset_bias, act, + accumulate, acc_buff, row_sum, 0); } template<> @@ -300,7 +332,7 @@ void kernel_and_merge::run( strategy &strat, const Tlo *a_ptr, const Tro *b_panel, size_t, Tri *c_panel, Tr *c_ptr, int ldc, int kern_k, unsigned int m_0, unsigned int m_max, unsigned int n_0, unsigned int n_max, const Tr *bias, - const Activation &act, bool not_first_pass, const DequantizeFloat &qp, const int32_t *, + const Activation &act, bool not_first_pass, const DequantizeFloat &qp, const int32_t *col_bias, Tab *) { const int bblocks = iceildiv(n_max - n_0, strategy::out_width()); @@ -317,6 +349,11 @@ void kernel_and_merge::run( #ifdef CYCLE_PROFILING auto p=prof.ScopedProfiler(PROFILE_QUANTIZE, ((m_max-m_0) * bblocks * strategy::out_width() * sizeof(Tr))); #endif + // When b_offset != 0, row sums are packed after the A panel data + const int32_t *row_sum = (qp.b_offset != 0) + ? reinterpret_cast(a_ptr + strategy::out_height() * kern_k) + : nullptr; + for (int i=0; i::run( dequantize_block_32(qp, (n_end - n_start), (m_max - m_0), c_panel + (i * strategy::out_width() * strategy::out_height()), strategy::out_width(), c_ptr + m_0 * ldc + n_start, ldc, - bias != nullptr ? bias + n_start : nullptr, not_first_pass, act); + bias != nullptr ? bias + n_start : nullptr, not_first_pass, act, + col_bias != nullptr ? col_bias + n_start : nullptr, row_sum, kern_k); } } @@ -475,6 +513,13 @@ class GemmInterleaved : public GemmCommon { return _Nsize * _nmulti * sizeof(int32_t); } + if (std::is_same::value) { + const DequantizeFloat *dq = reinterpret_cast(&_os); + if (dq->a_offset != 0) { + return _Nsize * _nmulti * sizeof(int32_t); + } + } + return 0; } @@ -557,6 +602,12 @@ class GemmInterleaved : public GemmCommon { k_depth += sizeof(int32_t) / sizeof(Tloi); } + if (std::is_same::value && MergeStep) { + // transforms_quantized always packs row sum slots (zeros when multiplier=0, actual + // sums when b_offset != 0). Reserve space unconditionally when MergeStep is enabled. + k_depth += sizeof(int32_t) / sizeof(Tloi); + } + return k_depth; } @@ -647,6 +698,13 @@ class GemmInterleaved : public GemmCommon { return -qp->b_offset; } + if (std::is_same::value) { + const DequantizeFloat *dq = reinterpret_cast(&_os); + // Pack row sums into the A panel when b_offset is non-zero so that the + // merge step can apply the b_offset correction per output position. + return (dq->b_offset != 0) ? 1 : 0; + } + return 0; } @@ -693,6 +751,14 @@ class GemmInterleaved : public GemmCommon { return get_ktotal(args); } + // K blocking is not supported for DequantizeFloat with MergeStep when b_offset != 0, + // because row sums of A must cover the full K depth. We cannot check b_offset here + // (static function), so we conservatively disable K-blocking for all DequantizeFloat + // MergeStep cases. The working-memory cost is minimal and correctness is guaranteed. + if (std::is_same::value && MergeStep) { + return get_ktotal(args); + } + // We can't K block non-fast FP16 cases without an accumulation buffer. #if defined(__aarch64__) && (defined(FP16_KERNELS) || defined(ARM_COMPUTE_ENABLE_FP16)) if (std::is_same::value && std::is_same::value && !args._fast_mode && MergeStep) { @@ -937,7 +1003,7 @@ class GemmInterleaved : public GemmCommon { #endif // See comment above on transform_type<> class: this extracts either 'transforms' or // 'transforms_quantized' as appropriate. - typename transform_type::value>::type transforms; + typename transform_type::value || std::is_same::value)>::type transforms; if (_indirect_buf != nullptr) { transforms.PrepareA_indirect(a_panel, @@ -1027,7 +1093,7 @@ class GemmInterleaved : public GemmCommon { #endif // See comment above on transform_type<> class: this extracts either 'transforms' or // 'transforms_quantized' as appropriate. - typename transform_type::value>::type transforms; + typename transform_type::value || std::is_same::value)>::type transforms; for (unsigned int batch = batch_0; batch <= batch_end; batch++) { unsigned int first_m = (batch == batch_0) ? m_0 : 0; @@ -1060,6 +1126,10 @@ class GemmInterleaved : public GemmCommon { if(std::is_same::value) { a_panel_stride = kern_k + (sizeof(int32_t) / sizeof(Tloi)); + } else if (std::is_same::value && MergeStep) { + // transforms_quantized always packs row-sum slots (zeros when b_offset=0, + // actual sums when b_offset != 0), so the stride must include the slot. + a_panel_stride = kern_k + (sizeof(int32_t) / sizeof(Tloi)); } else { a_panel_stride = kern_k; } @@ -1212,6 +1282,20 @@ class GemmInterleaved : public GemmCommon { compute_col_sums(*qp_ptr, _Nsize, _Ksize * _Ksections, B + (i * B_multi_stride), ldb, col_bias + (i * _Nsize), _Ksize * _Ksections, i, 0); } } + + if (std::is_same::value) { + const DequantizeFloat *dq = reinterpret_cast(&_os); + if (dq->a_offset != 0) { + // Compute raw column sums of B (weight matrix) for use in a_offset correction. + // dequantize_block_32 applies: -a_offset * col_sums[n] * scale per output channel. + col_bias = reinterpret_cast(in_buffer); + for (unsigned int i = 0; i < _nmulti; ++i) { + compute_raw_col_sums(_Nsize, _Ksize * _Ksections, + B + (i * B_multi_stride), ldb, + col_bias + (i * _Nsize)); + } + } + } } // Support for transposed B is a property of the strategy::transpose type @@ -1431,4 +1515,3 @@ template using GemmInterleavedDequantized = GemmInterleaved; } // namespace arm_gemm - diff --git a/src/core/NEON/kernels/arm_gemm/gemm_s8fp32.cpp b/src/core/NEON/kernels/arm_gemm/gemm_s8fp32.cpp index 52716fe025f..99c5a59c918 100644 --- a/src/core/NEON/kernels/arm_gemm/gemm_s8fp32.cpp +++ b/src/core/NEON/kernels/arm_gemm/gemm_s8fp32.cpp @@ -49,26 +49,46 @@ namespace arm_gemm { +#if defined(ARM_COMPUTE_ENABLE_SME2) || defined(ARM_COMPUTE_ENABLE_SME) +namespace { + +bool supports_symmetric_dequant_no_merge(const GemmArgs &args, const DequantizeFloat &dq) +{ + return !args._accumulate && dq.a_offset == 0 && dq.b_offset == 0; +} + +} // namespace +#endif // defined(ARM_COMPUTE_ENABLE_SME2) || defined(ARM_COMPUTE_ENABLE_SME) + static const GemmImplementation gemm_s8fp32_methods[] = { #ifdef ARM_COMPUTE_ENABLE_SME2 { "sme2_interleaved_nomerge_s8qfp32_mopa_1VLx4VL", - [](const GemmArgs &args, const DequantizeFloat &) { return args._ci->has_sme2() && args._ci->has_sme_i8i32() && !args._accumulate; }, + [](const GemmArgs &args, const DequantizeFloat &dq) { + return args._ci->has_sme2() && args._ci->has_sme_i8i32() && + supports_symmetric_dequant_no_merge(args, dq); + }, [](const GemmArgs &args, const DequantizeFloat &) { const auto VL = sme::get_vector_length(); return args._Msize <= VL || (2*VL < args._Msize && args._Msize <= 3*VL); }, [](const GemmArgs &args, const DequantizeFloat &dq) { return new GemmInterleavedNoMergeDequantized(args, dq); } }, { "sme2_interleaved_nomerge_s8qfp32_mopa_4VLx1VL", - [](const GemmArgs &args, const DequantizeFloat &) { return args._ci->has_sme2() && args._ci->has_sme_i8i32() && !args._accumulate; }, + [](const GemmArgs &args, const DequantizeFloat &dq) { + return args._ci->has_sme2() && args._ci->has_sme_i8i32() && + supports_symmetric_dequant_no_merge(args, dq); + }, [](const GemmArgs &args, const DequantizeFloat &) { const auto VL = sme::get_vector_length(); return args._Nsize <= VL || (2*VL < args._Nsize && args._Nsize <= 3*VL); }, [](const GemmArgs &args, const DequantizeFloat &dq) { return new GemmInterleavedNoMergeDequantized(args, dq); } }, { "sme2_interleaved_nomerge_s8qfp32_mopa_2VLx2VL", - [](const GemmArgs &args, const DequantizeFloat &) { return args._ci->has_sme2() && args._ci->has_sme_i8i32() && !args._accumulate; }, + [](const GemmArgs &args, const DequantizeFloat &dq) { + return args._ci->has_sme2() && args._ci->has_sme_i8i32() && + supports_symmetric_dequant_no_merge(args, dq); + }, nullptr, [](const GemmArgs &args, const DequantizeFloat &dq) { return new GemmInterleavedNoMergeDequantized(args, dq); } }, @@ -76,21 +96,30 @@ static const GemmImplementation gemm_s8f #ifdef ARM_COMPUTE_ENABLE_SME { "sme_interleaved_nomerge_s8qfp32_mopa_1VLx4VL", - [](const GemmArgs &args, const DequantizeFloat &) { return args._ci->has_sme() && args._ci->has_sme_i8i32() && !args._accumulate; }, + [](const GemmArgs &args, const DequantizeFloat &dq) { + return args._ci->has_sme() && args._ci->has_sme_i8i32() && + supports_symmetric_dequant_no_merge(args, dq); + }, [](const GemmArgs &args, const DequantizeFloat &) { const auto VL = sme::get_vector_length(); return args._Msize <= VL || (2*VL < args._Msize && args._Msize <= 3*VL); }, [](const GemmArgs &args, const DequantizeFloat &dq) { return new GemmInterleavedNoMergeDequantized(args, dq); } }, { "sme_interleaved_nomerge_s8qfp32_mopa_4VLx1VL", - [](const GemmArgs &args, const DequantizeFloat &) { return args._ci->has_sme() && args._ci->has_sme_i8i32() && !args._accumulate; }, + [](const GemmArgs &args, const DequantizeFloat &dq) { + return args._ci->has_sme() && args._ci->has_sme_i8i32() && + supports_symmetric_dequant_no_merge(args, dq); + }, [](const GemmArgs &args, const DequantizeFloat &) { const auto VL = sme::get_vector_length(); return args._Nsize <= VL || (2*VL < args._Nsize && args._Nsize <= 3*VL); }, [](const GemmArgs &args, const DequantizeFloat &dq) { return new GemmInterleavedNoMergeDequantized(args, dq); } }, { "sme_interleaved_nomerge_s8qfp32_mopa_2VLx2VL", - [](const GemmArgs &args, const DequantizeFloat &) { return args._ci->has_sme() && args._ci->has_sme_i8i32() && !args._accumulate; }, + [](const GemmArgs &args, const DequantizeFloat &dq) { + return args._ci->has_sme() && args._ci->has_sme_i8i32() && + supports_symmetric_dequant_no_merge(args, dq); + }, nullptr, [](const GemmArgs &args, const DequantizeFloat &dq) { return new GemmInterleavedNoMergeDequantized(args, dq); } }, @@ -153,4 +182,3 @@ template std::vector get_compatible_kernels void dequantize_block_32<__fp16>(const DequantizeFloat &qp, unsigned int width, unsigned int height, const int32_t * in_ptr, unsigned int in_stride, __fp16 *out_ptr, unsigned int out_stride, - const __fp16 * bias_ptr, bool not_first_pass, const Activation &act) + const __fp16 * bias_ptr, bool not_first_pass, const Activation &act, + const int32_t * col_bias, const int32_t * row_sum, int32_t k_total) { const float32x4_t vscale = vdupq_n_f32(qp.scale); float maxval = std::numeric_limits::infinity(); @@ -57,11 +58,25 @@ void dequantize_block_32<__fp16>(const DequantizeFloat &qp, unsigned int width, const float16x8_t vmax = vdupq_n_f16(static_cast<__fp16>(maxval)); for(unsigned int row=0; row(-qp.b_offset * row_sum[row]) * qp.scale; + } + const int32_t effective_k_total = (qp.depth != 0) ? qp.depth : k_total; + if (col_bias != nullptr && row_sum != nullptr && effective_k_total != 0) { + row_offset += static_cast(qp.a_offset) * static_cast(qp.b_offset) + * static_cast(effective_k_total) * qp.scale; + } + const float32x4_t vrow_offset = vdupq_n_f32(row_offset); + unsigned int col=0; if (width >= 8) { for(; col <= (width-8); col+=8) { - const int32x4_t vin0 = vld1q_s32(in_ptr + col + (row * in_stride)); - const int32x4_t vin1 = vld1q_s32(in_ptr + col + 4 + (row * in_stride)); + const int32x4_t vin0 = vld1q_s32(row_in_ptr + col); + const int32x4_t vin1 = vld1q_s32(row_in_ptr + col + 4); float32x4_t vdeq0 = vmulq_f32(vcvtq_f32_s32(vin0), vscale); float32x4_t vdeq1 = vmulq_f32(vcvtq_f32_s32(vin1), vscale); @@ -75,8 +90,24 @@ void dequantize_block_32<__fp16>(const DequantizeFloat &qp, unsigned int width, vdeq1 = vaddq_f32(vdeq1, bin1); } + if(col_bias) { + const float32x4_t vcol_corr0 = vmulq_f32( + vcvtq_f32_s32(vld1q_s32(col_bias + col)), + vdupq_n_f32(static_cast(-qp.a_offset) * qp.scale)); + const float32x4_t vcol_corr1 = vmulq_f32( + vcvtq_f32_s32(vld1q_s32(col_bias + col + 4)), + vdupq_n_f32(static_cast(-qp.a_offset) * qp.scale)); + vdeq0 = vaddq_f32(vdeq0, vcol_corr0); + vdeq1 = vaddq_f32(vdeq1, vcol_corr1); + } + + if(row_sum) { + vdeq0 = vaddq_f32(vdeq0, vrow_offset); + vdeq1 = vaddq_f32(vdeq1, vrow_offset); + } + if(not_first_pass) { - const float16x8_t in = vld1q_f16(out_ptr + col + (row * out_stride)); + const float16x8_t in = vld1q_f16(row_out_ptr + col); const float32x4_t in0 = vcvt_f32_f16(vget_low_f16(in)); const float32x4_t in1 = vcvt_f32_f16(vget_high_f16(in)); @@ -87,21 +118,27 @@ void dequantize_block_32<__fp16>(const DequantizeFloat &qp, unsigned int width, float16x8_t vdeq16 = vcombine_f16(vcvt_f16_f32(vdeq0), vcvt_f16_f32(vdeq1)); vdeq16 = vminq_f16(vmaxq_f16(vdeq16, vmin), vmax); - vst1q_f16(out_ptr + col + (row * out_stride), vdeq16); + vst1q_f16(row_out_ptr + col, vdeq16); } } // left-over elements for(; col < width; ++col) { - const int32_t val = *(in_ptr + (row * in_stride) + col); - float res = static_cast(val * qp.scale); + const int32_t val = *(row_in_ptr + col); + float res = static_cast(val) * qp.scale; if(bias_ptr) { res += static_cast(*(bias_ptr + col)); } + if(col_bias) { + res += static_cast(-qp.a_offset * col_bias[col]) * qp.scale; + } + if(row_sum) { + res += row_offset; + } if(not_first_pass) { - res += *(out_ptr + (row * out_stride) + col); + res += *(row_out_ptr + col); } res = std::min(std::max(res, minval), maxval); - *(out_ptr + (row * out_stride) + col) = static_cast<__fp16>(res); + *(row_out_ptr + col) = static_cast<__fp16>(res); } } } diff --git a/src/core/NEON/kernels/arm_gemm/quantized.cpp b/src/core/NEON/kernels/arm_gemm/quantized.cpp index 0243ee14a3b..662cb090778 100644 --- a/src/core/NEON/kernels/arm_gemm/quantized.cpp +++ b/src/core/NEON/kernels/arm_gemm/quantized.cpp @@ -973,10 +973,28 @@ void compute_col_sums(const Requantize32 &qp, unsigned int width, unsigned int h template void compute_col_sums(const Requantize32 &qp, unsigned int width, unsigned int height, const int8_t *input, unsigned int in_stride, int32_t *col_bias, unsigned int depth, unsigned int multi, unsigned int first_col); template void compute_col_sums(const Requantize32 &qp, unsigned int width, unsigned int height, const uint8_t *input, unsigned int in_stride, int32_t *col_bias, unsigned int depth, unsigned int multi, unsigned int first_col); +template +void compute_raw_col_sums(unsigned int width, unsigned int height, + const T *input, unsigned int in_stride, int32_t *col_sums) +{ + memset(reinterpret_cast(col_sums), 0, width * sizeof(int32_t)); + for (unsigned int row = 0; row < height; ++row) + { + for (unsigned int col = 0; col < width; ++col) + { + col_sums[col] += static_cast(input[row * in_stride + col]); + } + } +} + +template void compute_raw_col_sums(unsigned int width, unsigned int height, const int8_t *input, unsigned int in_stride, int32_t *col_sums); +template void compute_raw_col_sums(unsigned int width, unsigned int height, const uint8_t *input, unsigned int in_stride, int32_t *col_sums); + template<> void dequantize_block_32(const DequantizeFloat &qp, unsigned int width, unsigned int height, const int32_t* in_ptr, unsigned int in_stride, float *out_ptr, unsigned int out_stride, - const float* bias_ptr, bool accumulate, const Activation &act) + const float* bias_ptr, bool accumulate, const Activation &act, + const int32_t *col_bias, const int32_t *row_sum, int32_t k_total) { const float32x4_t vscale = vdupq_n_f32(qp.scale); float maxval = std::numeric_limits::infinity(); @@ -1000,14 +1018,39 @@ void dequantize_block_32(const DequantizeFloat &qp, unsigned int width, u for(unsigned int row=0; row(-qp.b_offset * row_sum[row]) * qp.scale; + } + const int32_t effective_k_total = (qp.depth != 0) ? qp.depth : k_total; + if (col_bias != nullptr && row_sum != nullptr && effective_k_total != 0) { + // Cross-term: +a_offset * b_offset * K * scale + row_offset += static_cast(qp.a_offset) * static_cast(qp.b_offset) + * static_cast(effective_k_total) * qp.scale; + } + const float32x4_t vrow_offset = vdupq_n_f32(row_offset); + unsigned int col=0; if (width >= 4) { for(; col <= (width - 4); col+= 4) { const int32x4_t vin = vld1q_s32(row_in_ptr + col); float32x4_t vdeq = vmulq_f32(vcvtq_f32_s32(vin), vscale); if(bias_ptr) { - const float32x4_t bin = vld1q_f32(bias_ptr + col); - vdeq = vaddq_f32(vdeq, bin); + vdeq = vaddq_f32(vdeq, vld1q_f32(bias_ptr + col)); + } + if(col_bias) { + // a_offset correction: -a_offset * sum_b_col[n] * scale + const float32x4_t vcol_corr = vmulq_f32( + vcvtq_f32_s32(vld1q_s32(col_bias + col)), + vdupq_n_f32(static_cast(-qp.a_offset) * qp.scale)); + vdeq = vaddq_f32(vdeq, vcol_corr); + } + if(row_sum) { + vdeq = vaddq_f32(vdeq, vrow_offset); } if(accumulate) { vdeq = vaddq_f32(vdeq, vld1q_f32(row_out_ptr + col)); @@ -1019,10 +1062,16 @@ void dequantize_block_32(const DequantizeFloat &qp, unsigned int width, u // left-over elements for(; col < width; ++col) { const int32_t val = *(row_in_ptr + col); - float res = static_cast(val * qp.scale); + float res = static_cast(val) * qp.scale; if(bias_ptr) { res += static_cast(*(bias_ptr + col)); } + if(col_bias) { + res += static_cast(-qp.a_offset * col_bias[col]) * qp.scale; + } + if(row_sum) { + res += row_offset; + } if(accumulate) { res += *(row_out_ptr + col); } @@ -1035,4 +1084,3 @@ void dequantize_block_32(const DequantizeFloat &qp, unsigned int width, u } // namespace arm_gemm #endif // __aarch64__ - diff --git a/src/cpu/kernels/assembly/arm_common/internal/quantized.hpp b/src/cpu/kernels/assembly/arm_common/internal/quantized.hpp index 796da4e2e19..64985e7b9ba 100644 --- a/src/cpu/kernels/assembly/arm_common/internal/quantized.hpp +++ b/src/cpu/kernels/assembly/arm_common/internal/quantized.hpp @@ -41,6 +41,14 @@ void compute_col_sums(const Requantize32 &qp, unsigned int width, unsigned int h const T *input, unsigned int in_stride, int32_t *col_bias, unsigned int depth, unsigned int multi, unsigned int first_col); +/** Compute raw column sums of a matrix: col_sums[n] = sum_{k} input[k * in_stride + n]. + * Unlike compute_col_sums(), this does not apply any quantization offsets or scaling — + * it stores the plain integer sums for use as weight column reductions in the + * DequantizeFloat a_offset correction path. */ +template +void compute_raw_col_sums(unsigned int width, unsigned int height, + const T *input, unsigned int in_stride, int32_t *col_sums); + template void row_sums_indirect(size_t num_strings, const unsigned int *string_lengths, IndirectInputArg A_arg, size_t M, int32_t *output_ptr, const Requantize32 *qp); @@ -48,6 +56,8 @@ void row_sums_indirect(size_t num_strings, const unsigned int *string_lengths, I template void dequantize_block_32(const DequantizeFloat &qp, unsigned int width, unsigned int height, const int32_t* input, unsigned int in_stride, T *output, unsigned int out_stride, - const T *row_bias, bool not_first_pass, const Activation &act); + const T *row_bias, bool not_first_pass, const Activation &act, + const int32_t *col_bias = nullptr, const int32_t *row_sum = nullptr, + int32_t k_total = 0); } // namespace arm_gemm diff --git a/src/cpu/kernels/assembly/arm_gemm/arm_gemm.hpp b/src/cpu/kernels/assembly/arm_gemm/arm_gemm.hpp index 8bae014a757..043a14f6e15 100644 --- a/src/cpu/kernels/assembly/arm_gemm/arm_gemm.hpp +++ b/src/cpu/kernels/assembly/arm_gemm/arm_gemm.hpp @@ -240,12 +240,21 @@ struct Requantize32 struct DequantizeFloat { public: - float scale = 0; + float scale = 0; + int32_t a_offset = 0; // input quantization zero-point (subtract from each input value) + int32_t b_offset = 0; // weight quantization zero-point (subtract from each weight value) + int32_t depth = 0; // unrounded K depth used by zero-point cross-term DequantizeFloat() = default; - // Constructor - DequantizeFloat(const float scale) : scale(scale) + // Constructor without offset (symmetric quantization) + DequantizeFloat(const float scale) : scale(scale), a_offset(0), b_offset(0), depth(0) + { + } + + // Constructor with asymmetric quantization offsets + DequantizeFloat(const float scale, int32_t a_offset, int32_t b_offset) + : scale(scale), a_offset(a_offset), b_offset(b_offset), depth(0) { } }; diff --git a/src/cpu/operators/CpuConv2d.cpp b/src/cpu/operators/CpuConv2d.cpp index adc4486889c..f4037406c21 100644 --- a/src/cpu/operators/CpuConv2d.cpp +++ b/src/cpu/operators/CpuConv2d.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2021, 2023-2025 Arm Limited. + * Copyright (c) 2017-2021, 2023-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -65,7 +65,7 @@ void CpuConv2d::configure(ITensorInfo *input, ARM_COMPUTE_LOG_PARAMS(input, weights, biases, output, conv_info, weights_info, dilation, act_info, enable_fast_math, num_groups); - const Conv2dInfo info(conv_info, dilation, act_info, enable_fast_math, num_groups); + const Conv2dInfo info(conv_info, dilation, act_info, enable_fast_math, num_groups, weights_info); switch (CpuConv2d::get_convolution_method(input, weights, output, conv_info, weights_info, dilation, act_info, enable_fast_math)) { @@ -118,8 +118,18 @@ Status CpuConv2d::validate(const ITensorInfo *input, { ARM_COMPUTE_TRACE_EVENT(ARM_COMPUTE_PROF_CAT_CPU, ARM_COMPUTE_PROF_LVL_CPU, "CpuConv2d::validate"); ARM_COMPUTE_RETURN_ERROR_ON_MSG((num_groups != 1), "Grouping (num_groups != 1) is not supported on Neon"); + // TODO(COMPMID-9268): Temporary guard until TensorInfo rejects inconsistent quantization metadata + // centrally. QASYMM8_SIGNED tensors can currently carry multiple quantization scales even though + // this DequantizeFloat path only supports uniform weight scales. Remove this local check once + // COMPMID-9268 is resolved. + ARM_COMPUTE_RETURN_ERROR_ON_MSG( + input->data_type() == DataType::QASYMM8_SIGNED && weights->data_type() == DataType::QASYMM8_SIGNED && + (output->data_type() == DataType::F32 || output->data_type() == DataType::F16) && + weights->quantization_info().scale().size() > 1, + "Per-channel QASYMM8_SIGNED weight scales are not supported for dequantized floating-point convolution"); + + const Conv2dInfo info(conv_info, dilation, act_info, enable_fast_math, num_groups, weights_info); - const Conv2dInfo info(conv_info, dilation, act_info, enable_fast_math, num_groups); switch (CpuConv2d::get_convolution_method(input, weights, output, conv_info, weights_info, dilation, act_info, enable_fast_math)) { @@ -155,7 +165,18 @@ ConvolutionMethod CpuConv2d::get_convolution_method(const ITensorInfo *i bool enable_fast_math) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output, weights); - ARM_COMPUTE_UNUSED(weights_info); + + // For QASYMM8_SIGNED→floating-point with NHWC and no dilation, automatically select the single-kernel + // CpuGemmDirectConv2d path when it validates successfully. + if (input->data_type() == DataType::QASYMM8_SIGNED && + (output->data_type() == DataType::F32 || output->data_type() == DataType::F16)) + { + const Conv2dInfo info(conv_info, dilation, act_info, enable_fast_math, 1, weights_info); + if (bool(CpuGemmDirectConv2d::validate(input, weights, nullptr, output, info))) + { + return ConvolutionMethod::GEMM_CONV2D; + } + } const size_t idx_w = get_data_layout_dimension_index(input->data_layout(), DataLayoutDimension::WIDTH); const size_t idx_h = get_data_layout_dimension_index(input->data_layout(), DataLayoutDimension::HEIGHT); diff --git a/src/cpu/operators/CpuConv2d.h b/src/cpu/operators/CpuConv2d.h index 3ad69e76148..f70b1238efb 100644 --- a/src/cpu/operators/CpuConv2d.h +++ b/src/cpu/operators/CpuConv2d.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2021, 2023-2025 Arm Limited. + * Copyright (c) 2017-2021, 2023-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -89,6 +89,7 @@ class CpuConv2d : public ICpuOperator * |QASYMM8 |QSYMM8_PER_CHANNEL |S32 |QASYMM8 | * |QASYMM8_SIGNED |QASYMM8_SIGNED |S32 |QASYMM8_SIGNED | * |QASYMM8_SIGNED |QASYMM8_SIGNED |F32 |F32 | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F16 |F16 | * |QASYMM8_SIGNED |QSYMM8_PER_CHANNEL |S32 |QASYMM8_SIGNED | * * @param[in] src Source tensor info. 3 lower dimensions represent a single input [width, height, IFM], @@ -98,9 +99,10 @@ class CpuConv2d : public ICpuOperator * Data type supported: Same as @p src, also could be QSYMM8_PER_CHANNEL or QASYMM8_SIGNED if input is QASYMM8/QASYMM8_SIGNED. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Same as @p src, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. - * For F32 dequantization the bias must be F32. + * For floating-point dequantization the bias must match @p dst. * @param[out] dst Destination tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p src. + * Data types supported: Same as @p src, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with NEWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. @@ -144,7 +146,8 @@ class CpuConv2d : public ICpuOperator * @param[in] weights Weights tensor info. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. * Data type supported:Same as @p src, also could be QSYMM8_PER_CHANNEL or QASYMM8_SIGNED if input is QASYMM8/QASYMM8_SIGNED. * @param[in] dst Destination tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p src. + * Data types supported: Same as @p src, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with NEWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. diff --git a/src/cpu/operators/CpuGemmConv2d.cpp b/src/cpu/operators/CpuGemmConv2d.cpp index f51cb761cc9..8db2951a5ac 100644 --- a/src/cpu/operators/CpuGemmConv2d.cpp +++ b/src/cpu/operators/CpuGemmConv2d.cpp @@ -172,6 +172,7 @@ CpuGemmConv2d::WeightTransformMethod CpuGemmConv2d::get_wt_method(const ITensorI CpuGemmConv2d::SkipInfo CpuGemmConv2d::skip_im_col_info(const ITensorInfo *src, const ITensorInfo *weights, + const ITensorInfo *dst, const PadStrideInfo &conv_info, const Size2D &dilation, const ActivationLayerInfo &act_info) @@ -194,7 +195,7 @@ CpuGemmConv2d::SkipInfo CpuGemmConv2d::skip_im_col_info(const ITensorInfo { const bool skip_col2im = (data_layout == DataLayout::NHWC && - (bool(CpuGemmConv2d::validate_gemm3d(src, weights, act_info, conv_h, /*skip_im2col*/ true)))); + (bool(CpuGemmConv2d::validate_gemm3d(src, weights, dst, act_info, conv_h, /*skip_im2col*/ true)))); if (skip_col2im) { return {true, true}; @@ -204,7 +205,7 @@ CpuGemmConv2d::SkipInfo CpuGemmConv2d::skip_im_col_info(const ITensorInfo { const bool skip_col2im = (data_layout == DataLayout::NHWC && - (bool(CpuGemmConv2d::validate_gemm3d(src, weights, act_info, conv_h, /*skip_im2col*/ false)))); + (bool(CpuGemmConv2d::validate_gemm3d(src, weights, dst, act_info, conv_h, /*skip_im2col*/ false)))); if (skip_col2im) { return {false, true}; @@ -392,6 +393,13 @@ Status CpuGemmConv2d::validate_mm(const ITensorInfo *src, GEMMLowpOutputStageInfo output_info; if (int8_dequantize_f32_path(data_type, dst->data_type())) { + // TODO(COMPMID-9268): Temporary guard until TensorInfo rejects inconsistent quantization metadata + // centrally. QASYMM8_SIGNED tensors can currently carry multiple quantization scales even though + // this DequantizeFloat path only supports uniform weight scales. Remove this local check once + // COMPMID-9268 is resolved. + ARM_COMPUTE_RETURN_ERROR_ON_MSG( + weights->data_type() == DataType::QASYMM8_SIGNED && weights->quantization_info().scale().size() > 1, + "Per-channel QASYMM8_SIGNED weight scales are not supported for dequantized F32 GEMM convolution"); // No requant stage; offsets are handled via offset-contribution on int32 output_info.type = GEMMLowpOutputStageType::NONE; output_info.gemmlowp_offset = 0; @@ -439,20 +447,23 @@ Status CpuGemmConv2d::validate_mm(const ITensorInfo *src, Status CpuGemmConv2d::validate_gemm3d(const ITensorInfo *input_info, const ITensorInfo *weights_info, + const ITensorInfo *output_info, const ActivationLayerInfo &act_info, int gemm_3d_depth, bool skip_im2col) { - const DataType data_type = input_info->data_type(); - const unsigned int mult_y = skip_im2col ? 1U : gemm_3d_depth; - const unsigned int mult_z = skip_im2col ? gemm_3d_depth : 1U; + const DataType input_data_type = input_info->data_type(); + const DataType weights_data_type = weights_info->data_type(); + const DataType output_data_type = output_info->data_type(); + const unsigned int mult_y = skip_im2col ? 1U : gemm_3d_depth; + const unsigned int mult_z = skip_im2col ? gemm_3d_depth : 1U; // Set dummy tensor shapes for the validation - const TensorInfo dummy_input_info(TensorShape(4U, 4U * mult_y, 1U * mult_z), 1, data_type, + const TensorInfo dummy_input_info(TensorShape(4U, 4U * mult_y, 1U * mult_z), 1, input_data_type, input_info->quantization_info()); - const TensorInfo dummy_weights_info(TensorShape(4U, 4U), 1, data_type, weights_info->quantization_info()); - const TensorInfo dummy_output_info(TensorShape(4U, 4U, gemm_3d_depth), 1, data_type, - input_info->quantization_info()); + const TensorInfo dummy_weights_info(TensorShape(4U, 4U), 1, weights_data_type, weights_info->quantization_info()); + const TensorInfo dummy_output_info(TensorShape(4U, 4U, gemm_3d_depth), 1, output_data_type, + output_info->quantization_info()); return validate_mm(&dummy_input_info, &dummy_weights_info, nullptr, &dummy_output_info, act_info, false, gemm_3d_depth, skip_im2col); @@ -507,7 +518,7 @@ void CpuGemmConv2d::configure(const ITensorInfo *src, // Check if GEMM3D is supported const CpuGemmConv2d::SkipInfo skip_info = - CpuGemmConv2d::skip_im_col_info(src, weights, conv_info, dilation, act_info); + CpuGemmConv2d::skip_im_col_info(src, weights, dst, conv_info, dilation, act_info); _skip_im2col = skip_info.skip_im2col; _skip_col2im = skip_info.skip_col2im; @@ -672,7 +683,7 @@ Status CpuGemmConv2d::has_opt_impl(arm_compute::WeightFormat &expected_weight_fo kernel_height, conv_info, dilation); const CpuGemmConv2d::SkipInfo skip_info = - CpuGemmConv2d::skip_im_col_info(src, weights, conv_info, dilation, act_info); + CpuGemmConv2d::skip_im_col_info(src, weights, dst, conv_info, dilation, act_info); const bool skip_im2col = skip_info.skip_im2col; const bool skip_col2im = skip_info.skip_col2im; @@ -752,7 +763,7 @@ Status CpuGemmConv2d::validate(const ITensorInfo *src, // Check if GEMM3D is supported const CpuGemmConv2d::SkipInfo skip_info = - CpuGemmConv2d::skip_im_col_info(src, weights, conv_info, dilation, act_info); + CpuGemmConv2d::skip_im_col_info(src, weights, dst, conv_info, dilation, act_info); const bool skip_im2col = skip_info.skip_im2col, skip_col2im = skip_info.skip_col2im; ARM_COMPUTE_RETURN_ERROR_ON(weights->dimension(idx_channel) != src->dimension(idx_channel)); diff --git a/src/cpu/operators/CpuGemmConv2d.h b/src/cpu/operators/CpuGemmConv2d.h index e4e34cc7c5f..c3a4820885d 100644 --- a/src/cpu/operators/CpuGemmConv2d.h +++ b/src/cpu/operators/CpuGemmConv2d.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2024 Arm Limited. + * Copyright (c) 2021-2024, 2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -79,6 +79,8 @@ class CpuGemmConv2d : public ICpuOperator * |QASYMM8 |QASYMM8_SIGNED |S32 |QASYMM8 | * |QASYMM8 |QSYMM8_PER_CHANNEL |S32 |QASYMM8 | * |QASYMM8_SIGNED |QASYMM8_SIGNED |S32 |QASYMM8_SIGNED | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F32 |F32 | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F16 |F16 | * |QASYMM8_SIGNED |QSYMM8_PER_CHANNEL |S32 |QASYMM8_SIGNED | * * @param[in] src Source tensor info. 3 lower dimensions represent a single input [width, height, IFM], @@ -88,8 +90,10 @@ class CpuGemmConv2d : public ICpuOperator * Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * For floating-point dequantization the bias must match @p dst. * @param[out] dst Destination tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] weights_info Specifies if the weights tensor has been reshaped with CpuWeightsReshapeKernel. If this is not part of the fully connected layer the weights * tensor has also been transposed with cpu::kernels::CpuGemmTranspose1xWKernel. Data type supported: Same as @p input. @@ -161,8 +165,10 @@ class CpuGemmConv2d : public ICpuOperator * @param[in] weights Weights tensor info. Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * For floating-point dequantization the bias must match @p dst. * @param[out] dst Output tensor info. Data types supported: Same as @p input, - * except for input of QASYMM8/QASYMM8_SIGNED type where output should be of S32 type. + * except for input of QASYMM8/QASYMM8_SIGNED type where output should be of S32 type, or the + * QASYMM8_SIGNED to floating-point output path listed above. * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Only RELU, BOUNDED_RELU and LU_BOUNDED_RELU supported. * @param[in] enable_fast_math (Optional) Enable fast math computation. In case this flag were set, the function could dispatch the fastest implementation * available which may introduce a drop of accuracy as well. Default is false @@ -185,8 +191,10 @@ class CpuGemmConv2d : public ICpuOperator * @param[in] weights Weights tensor info. Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * For floating-point dequantization the bias must match @p dst. * @param[in] dst Output tensor info. Data types supported: Same as @p input, - * except for input of QASYMM8/QASYMM8_SIGNED type where output should be of S32 type. + * except for input of QASYMM8/QASYMM8_SIGNED type where output should be of S32 type, or the + * QASYMM8_SIGNED to floating-point output path listed above. * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Only RELU, BOUNDED_RELU and LU_BOUNDED_RELU supported. * @param[in] enable_fast_math (Optional) Enable fast math computation. In case this flag were set, the function could dispatch the fastest implementation * available which may introduce a drop of accuracy as well. Default is false @@ -211,6 +219,7 @@ class CpuGemmConv2d : public ICpuOperator * * @param[in] src Input tensor info. Data types supported: QASYMM8/QASYMM8_SIGNED/BFLOAT16/F16/F32. * @param[in] weights Weights tensor info. Data types supported: QASYMM8/QASYMM8_SIGNED/BFLOAT16/F16/F32. + * @param[in] dst Output tensor info. * @param[in] act_info Activation layer information in case of a fused activation. Only RELU, BOUNDED_RELU and LU_BOUNDED_RELU supported. * @param[in] gemm_3d_depth Depth of GEMM 3D * @param[in] skip_im2col Flag which specifies if im2col has to be skipped. i.e. 1x1 convolution with NHWC data layout @@ -219,6 +228,7 @@ class CpuGemmConv2d : public ICpuOperator */ static Status validate_gemm3d(const ITensorInfo *src, const ITensorInfo *weights, + const ITensorInfo *dst, const ActivationLayerInfo &act_info, int gemm_3d_depth, bool skip_im2col); @@ -233,6 +243,7 @@ class CpuGemmConv2d : public ICpuOperator * * @param[in] src Input tensor info. * @param[in] weights Weights tensor info. + * @param[in] dst Output tensor info. * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo. * @param[in] dilation Dilation, in elements, across x and y. * @param[in] act_info Activation layer information in case of a fused activation. @@ -241,6 +252,7 @@ class CpuGemmConv2d : public ICpuOperator */ static SkipInfo skip_im_col_info(const ITensorInfo *src, const ITensorInfo *weights, + const ITensorInfo *dst, const PadStrideInfo &conv_info, const Size2D &dilation, const ActivationLayerInfo &act_info); diff --git a/src/cpu/operators/CpuGemmDirectConv2d.cpp b/src/cpu/operators/CpuGemmDirectConv2d.cpp index caac7761e27..a451901a35d 100644 --- a/src/cpu/operators/CpuGemmDirectConv2d.cpp +++ b/src/cpu/operators/CpuGemmDirectConv2d.cpp @@ -44,6 +44,12 @@ using namespace arm_compute::utils::cast; namespace { +inline bool is_direct_i8_s8_float_path(const ITensorInfo *src, const ITensorInfo *dst) +{ + return src->data_type() == DataType::QASYMM8_SIGNED && + (dst->data_type() == DataType::F32 || dst->data_type() == DataType::F16); +} + GEMMLowpOutputStageInfo calculate_output_stage_metadata(const ITensorInfo *src, const ITensorInfo *weights, const ITensorInfo *dst, @@ -130,7 +136,15 @@ void CpuGemmDirectConv2d::configure(const ITensorInfo *src, // Configure assembly dispatch cpu::AsmGemmInfo asm_info = init_assembly_metadata(info, false); - if (is_data_type_quantized(src->data_type())) + if (is_direct_i8_s8_float_path(src, dst)) + { + // Provide the quantization zero-points via AsmGemmInfo so that create_arm_gemm_dequant + // can bake them into the DequantizeFloat output stage. The assembly kernel then handles + // all offset corrections (col sums, row sums, cross-term) natively. + asm_info.dequant_a_offset = src->quantization_info().uniform().offset; + asm_info.dequant_b_offset = weights->quantization_info().uniform().offset; + } + else if (is_data_type_quantized(src->data_type())) { asm_info.output_stage = calculate_output_stage_metadata(src, weights, dst, info.act_info); } @@ -189,6 +203,25 @@ Status CpuGemmDirectConv2d::validate(const ITensorInfo *src, ARM_COMPUTE_RETURN_ERROR_ON(info.dilation != Size2D(1U, 1U)); ARM_COMPUTE_RETURN_ERROR_ON(weights->num_dimensions() > 4); + const bool direct_i8_float = is_direct_i8_s8_float_path(src, dst); + + if (direct_i8_float) + { + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(src, 1, DataType::QASYMM8_SIGNED); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(weights, 1, DataType::QASYMM8_SIGNED); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(dst, 1, DataType::F32, DataType::F16); + // TODO(COMPMID-9268): Temporary guard until TensorInfo rejects inconsistent quantization metadata + // centrally. QASYMM8_SIGNED tensors can currently carry multiple quantization scales even though + // this DequantizeFloat path only supports uniform weight scales. Remove this local check once + // COMPMID-9268 is resolved. + ARM_COMPUTE_RETURN_ERROR_ON_MSG( + weights->quantization_info().scale().size() > 1, + "Per-channel QASYMM8_SIGNED weight scales are not supported for dequantized floating-point direct " + "convolution"); + // Offset correction (a_offset, b_offset) is handled inside the assembly kernel via the + // extended DequantizeFloat output stage. + } + // Validate Permute TensorInfo perm_weights; ARM_COMPUTE_RETURN_ON_ERROR(CpuPermute::validate(weights, &perm_weights, PermutationVector{3, 0, 1, 2})); @@ -204,7 +237,11 @@ Status CpuGemmDirectConv2d::validate(const ITensorInfo *src, // Validate biases if (biases != nullptr) { - if (is_data_type_quantized_asymmetric(data_type)) + if (direct_i8_float) + { + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(biases, dst); + } + else if (is_data_type_quantized_asymmetric(data_type)) { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(biases, 1, DataType::S32); } @@ -216,9 +253,17 @@ Status CpuGemmDirectConv2d::validate(const ITensorInfo *src, ARM_COMPUTE_RETURN_ERROR_ON(biases->num_dimensions() > 1); } - ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(src, dst); + if (!direct_i8_float) + { + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(src, dst); + } cpu::AsmGemmInfo asm_info = init_assembly_metadata(info, false); + if (direct_i8_float) + { + asm_info.dequant_a_offset = src->quantization_info().uniform().offset; + asm_info.dequant_b_offset = weights->quantization_info().uniform().offset; + } ARM_COMPUTE_RETURN_ON_ERROR(cpu::CpuGemmAssemblyDispatch::validate(src, weights, biases, dst, asm_info)); return Status{}; } diff --git a/src/cpu/operators/CpuGemmDirectConv2d.h b/src/cpu/operators/CpuGemmDirectConv2d.h index a7365615b94..b5832bad554 100644 --- a/src/cpu/operators/CpuGemmDirectConv2d.h +++ b/src/cpu/operators/CpuGemmDirectConv2d.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, 2023 Arm Limited. + * Copyright (c) 2021, 2023, 2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -55,6 +55,8 @@ class CpuGemmDirectConv2d : public ICpuOperator * |:--------------|:--------------|:--------------|:--------------| * |QASYMM8 |QASYMM8 |S32 |QASYMM8 | * |QASYMM8_SIGNED |QASYMM8_SIGNED |S32 |QASYMM8_SIGNED | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F32 |F32 | + * |QASYMM8_SIGNED |QASYMM8_SIGNED |F16 |F16 | * |F16 |F16 |F16 |F16 | * |F32 |F32 |F32 |F32 | * |BFLOAT16 |BFLOAT16 |BFLOAT16 |BFLOAT16 | @@ -65,9 +67,11 @@ class CpuGemmDirectConv2d : public ICpuOperator * @param[in] weights Weights tensor info. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. * Data type supported: QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL/BFLOAT16/F16/F32. * @param[in] biases Biases tensor info. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. - * Data type supported: Should match @p input data type, except for input of QASYMM8/QASYMM8_SIGNED type where biases should be of S32 type. + * Data type supported: Should match @p input data type, except for quantized output where biases + * should be S32, and QASYMM8_SIGNED to floating-point output where biases should match @p dst. * @param[in] dst Destination tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. - * Data types supported: Same as @p input. + * Data types supported: Same as @p input, except for the QASYMM8_SIGNED to floating-point output path + * listed above. * @param[in] info Contains padding and stride information described in @ref PadStrideInfo. */ void configure(const ITensorInfo *src, diff --git a/src/cpu/operators/CpuGemmLowpMatrixMultiplyCore.cpp b/src/cpu/operators/CpuGemmLowpMatrixMultiplyCore.cpp index 01fff57242c..b86dbea47d1 100644 --- a/src/cpu/operators/CpuGemmLowpMatrixMultiplyCore.cpp +++ b/src/cpu/operators/CpuGemmLowpMatrixMultiplyCore.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2025 Arm Limited. + * Copyright (c) 2021-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -58,9 +58,9 @@ namespace cpu { namespace { -inline bool int8_dequantize_f32_path(DataType src, DataType dst) +inline bool int8_dequantize_float_path(DataType src, DataType dst) { - return src == DataType::QASYMM8_SIGNED && dst == DataType::F32; + return src == DataType::QASYMM8_SIGNED && (dst == DataType::F32 || dst == DataType::F16); } cpu::AsmGemmInfo init_assembly_metadata(const GEMMInfo &info) @@ -137,13 +137,13 @@ void CpuGemmLowpMatrixMultiplyCore::configure( _reshape_b_only_on_first_run; _gemm_info = gemm_info; - // F32 dequant path? (input quantized, output float) - const bool dequantize_f32 = int8_dequantize_f32_path(a->data_type(), dst->data_type()); + // Float dequant path? (input quantized, output float) + const bool dequantize_float = int8_dequantize_float_path(a->data_type(), dst->data_type()); const ITensorInfo *a_to_use = a; // Initialize assembly kernel meta-data cpu::AsmGemmInfo asm_info = init_assembly_metadata(gemm_info); - if (dequantize_f32) + if (dequantize_float) { // We don't want arm_gemm to compute the activations because bias and offsets are added in ACL at a later step // so we disable activation in arm_gemm and run it as a post op in ACL @@ -184,6 +184,8 @@ void CpuGemmLowpMatrixMultiplyCore::configure( output_stage_corr.gemmlowp_min_bound -= offset_correction; output_stage_corr.gemmlowp_max_bound -= offset_correction; info.set_gemmlowp_output_stage(output_stage_corr); + asm_info.output_stage = output_stage_corr; + _gemm_info = info; // Update matrix a matrix_a = &_signed_a; @@ -200,8 +202,9 @@ void CpuGemmLowpMatrixMultiplyCore::configure( _fuse_output_stage = true; _mm_result_s32 = TensorInfo(dst->tensor_shape(), 1, DataType::S32); } - #ifdef __aarch64__ + ITensorInfo *dst_to_use = (_flip_signedness && _fuse_output_stage) ? &_signed_output : dst; + if (!(!b->are_values_constant() && b->tensor_shape().z() > 1)) // Disable batch matmul as optimized GeMM handles batching differently. { @@ -212,11 +215,11 @@ void CpuGemmLowpMatrixMultiplyCore::configure( case DataType::U8: case DataType::S8: { - if (!dequantize_f32 && is_data_type_quantized_asymmetric(a_to_use->data_type()) && + if (!dequantize_float && is_data_type_quantized_asymmetric(a_to_use->data_type()) && info.gemmlowp_output_stage().type == GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT) { auto c_info_to_use = c == nullptr ? nullptr : c; - _asm_glue->configure(a_to_use, b, c_info_to_use, dst, asm_info); + _asm_glue->configure(a_to_use, b, c_info_to_use, dst_to_use, asm_info); _fused_assembly_path = _asm_glue->is_configured(); } else @@ -293,12 +296,6 @@ void CpuGemmLowpMatrixMultiplyCore::configure( &_mm_result_s32, a_offset_kernel_needed ? &_vector_sum_col : nullptr, b_offset_kernel_needed ? &_vector_sum_row : nullptr, c, _flip_signedness ? &_signed_output : dst, a->dimension(0), _a_offset, _b_offset, info.gemmlowp_output_stage()); - - if (_flip_signedness) - { - _convert_from_signed_asymm = std::make_unique(); - _convert_from_signed_asymm->configure(&_signed_output, dst); - } } else { @@ -320,10 +317,15 @@ void CpuGemmLowpMatrixMultiplyCore::configure( a_to_use->dimension(0), _a_offset, _b_offset, dequantize_scale); } } + if (_fuse_output_stage && _flip_signedness) + { + _convert_from_signed_asymm = std::make_unique(); + _convert_from_signed_asymm->configure(&_signed_output, dst); + } // Configure activation const ActivationLayerInfo &activation = gemm_info.activation_info(); - _run_activation = activation.enabled() && (dequantize_f32 || !_assembly_path || + _run_activation = activation.enabled() && (dequantize_float || !_assembly_path || !cpu::CpuGemmAssemblyDispatch::is_activation_supported(activation)); if (_run_activation) { @@ -384,10 +386,17 @@ Status CpuGemmLowpMatrixMultiplyCore::validate(const ITensorInfo *a, ARM_COMPUTE_RETURN_ERROR_ON_MSG(gemm_info.pretranspose_A(), "Matrix A already pretransposed is not supported"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(gemm_info.pretranspose_B(), "Matrix B already pretransposed is not supported"); - if (int8_dequantize_f32_path(a->data_type(), output->data_type())) + if (int8_dequantize_float_path(a->data_type(), output->data_type())) { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(a, 1, DataType::QASYMM8_SIGNED); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(b, 1, DataType::QASYMM8_SIGNED); + // TODO(COMPMID-9268): Temporary guard until TensorInfo rejects inconsistent quantization metadata + // centrally. QASYMM8_SIGNED tensors can currently carry multiple quantization scales even though + // this DequantizeFloat path only supports uniform weight scales. Remove this local check once + // COMPMID-9268 is resolved. + ARM_COMPUTE_RETURN_ERROR_ON_MSG( + b->quantization_info().scale().size() > 1, + "Per-channel QASYMM8_SIGNED weight scales are not supported for dequantized output"); } // When using accumulation(in place summation), for now, the only supported DataType for output is S32. @@ -421,7 +430,7 @@ Status CpuGemmLowpMatrixMultiplyCore::validate(const ITensorInfo *a, } // Initialize assembly kernel meta-data - const AsmGemmInfo asm_info = init_assembly_metadata(info); + AsmGemmInfo asm_info = init_assembly_metadata(info); // Convert QASYMM8->QASYMM8_SIGNED const int32_t offset_correction = 128; @@ -459,10 +468,12 @@ Status CpuGemmLowpMatrixMultiplyCore::validate(const ITensorInfo *a, output_stage_corr.gemmlowp_min_bound -= offset_correction; output_stage_corr.gemmlowp_max_bound -= offset_correction; info.set_gemmlowp_output_stage(output_stage_corr); + asm_info.output_stage = output_stage_corr; // Update matrix a matrix_a_info = &signed_a; } + const ITensorInfo *output_to_use = (flip_signedness && fuse_output_stage) ? &signed_output : output; // Offset kernel is need if offset is non-zero or it may change (i.e. dynamic). bool a_offset_kernel_needed = a_offset != 0 || a->quantization_info().is_dynamic(); @@ -478,7 +489,7 @@ Status CpuGemmLowpMatrixMultiplyCore::validate(const ITensorInfo *a, if (is_data_type_quantized_asymmetric(a_to_use->data_type()) && info.gemmlowp_output_stage().type == GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT) { - run_optimised = bool(CpuGemmAssemblyDispatch::validate(a_to_use, b, c, output, asm_info)); + run_optimised = bool(CpuGemmAssemblyDispatch::validate(a_to_use, b, c, output_to_use, asm_info)); run_optimised_requantized = run_optimised; } else @@ -690,10 +701,11 @@ void CpuGemmLowpMatrixMultiplyCore::run(ITensorPack &tensors) if (is_data_type_quantized_asymmetric(a_to_use->info()->data_type()) && _gemm_info.gemmlowp_output_stage().type == GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT) { + auto output_to_use = (_flip_signedness && _fuse_output_stage) ? signed_output.get() : dst; asm_glue_tensors.add_const_tensor(TensorType::ACL_SRC_0, a_to_use); asm_glue_tensors.add_const_tensor(TensorType::ACL_SRC_1, b); asm_glue_tensors.add_const_tensor(TensorType::ACL_SRC_2, c); - asm_glue_tensors.add_tensor(TensorType::ACL_DST, dst); + asm_glue_tensors.add_tensor(TensorType::ACL_DST, output_to_use); } else { @@ -795,7 +807,7 @@ void CpuGemmLowpMatrixMultiplyCore::run(ITensorPack &tensors) } // Convert QASYMM8_SIGNED->QASYMM8 - if (!_fused_assembly_path && _fuse_output_stage && _flip_signedness) + if (_fuse_output_stage && _flip_signedness) { ITensorPack pack = {{TensorType::ACL_SRC, signed_output.get()}, {TensorType::ACL_DST, dst}}; NEScheduler::get().schedule_op(_convert_from_signed_asymm.get(), Window::DimY, @@ -856,6 +868,13 @@ void CpuGemmLowpMatrixMultiplyCore::update_quantization_parameters(const GEMMLow const bool negated_offsets) { auto lowp_os = output_info; + if (_flip_signedness && lowp_os.type != GEMMLowpOutputStageType::NONE) + { + const int32_t offset_correction = 128; + lowp_os.gemmlowp_offset -= offset_correction; + lowp_os.gemmlowp_min_bound -= offset_correction; + lowp_os.gemmlowp_max_bound -= offset_correction; + } _gemm_info.set_gemmlowp_output_stage(lowp_os); const QuantizationInfo *a_to_use = &a; @@ -871,7 +890,7 @@ void CpuGemmLowpMatrixMultiplyCore::update_quantization_parameters(const GEMMLow a_to_use = &a_signed; } - _asm_glue->update_quantization_parameters(output_info, *a_to_use, b, is_prepared, negated_offsets); + _asm_glue->update_quantization_parameters(lowp_os, *a_to_use, b, is_prepared, negated_offsets); _is_prepared = is_prepared; } } // namespace cpu diff --git a/src/cpu/operators/internal/CpuGemmAssemblyDispatch.cpp b/src/cpu/operators/internal/CpuGemmAssemblyDispatch.cpp index a8116441caa..edff52d019e 100644 --- a/src/cpu/operators/internal/CpuGemmAssemblyDispatch.cpp +++ b/src/cpu/operators/internal/CpuGemmAssemblyDispatch.cpp @@ -889,11 +889,11 @@ void create_arm_gemm_dequant(std::unique_ptr // Create arm_gemm fallback auto fallback = std::make_unique>(); - // Configure requantization info - const GEMMLowpOutputStageInfo os_info = info.output_stage; - arm_gemm::DequantizeFloat gemm_dequant_info{}; - gemm_dequant_info = arm_gemm::DequantizeFloat(d->quantization_info().uniform().scale); + gemm_dequant_info.scale = a->quantization_info().uniform().scale * b->quantization_info().uniform().scale; + gemm_dequant_info.a_offset = info.dequant_a_offset; + gemm_dequant_info.b_offset = info.dequant_b_offset; + gemm_dequant_info.depth = p.K * p.sections; fallback->configure(a, b, c, d, args, info, gemm_dequant_info); arm_gemm = std::move(fallback); @@ -1001,7 +1001,8 @@ Status CpuGemmAssemblyDispatch::has_opt_impl(arm_compute::WeightFormat &expected ARM_COMPUTE_RETURN_ERROR_ON_MSG( !(arm_gemm::has_opt_gemm(arm_gemm_expected_wf, args, {})), - "We could not find an optimized kernel for U8 input with S8 weights and U8 output"); + "We could not find an optimized kernel for U8/QASYMM8 input with QASYMM8_SIGNED weights and U8 " + "output"); } else { @@ -1020,12 +1021,32 @@ Status CpuGemmAssemblyDispatch::has_opt_impl(arm_compute::WeightFormat &expected {})), "We could not find an optimized kernel for S8/QASYMM8_SIGNED input and S32 output"); } - else + else if (a->data_type() == DataType::QASYMM8_SIGNED && d->data_type() == DataType::F32) + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG( + !(arm_gemm::has_opt_gemm(arm_gemm_expected_wf, + args, {})), + "We could not find an optimized kernel for QASYMM8_SIGNED input and F32 output"); + } +#if defined(ENABLE_FP16_KERNELS) + else if (a->data_type() == DataType::QASYMM8_SIGNED && d->data_type() == DataType::F16) + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG( + !(arm_gemm::has_opt_gemm(arm_gemm_expected_wf, + args, {})), + "We could not find an optimized kernel for QASYMM8_SIGNED input and F16 output"); + } +#endif /* defined(ENABLE_FP16_KERNELS) */ + else if (a->data_type() == DataType::QASYMM8_SIGNED) { ARM_COMPUTE_RETURN_ERROR_ON_MSG( !(arm_gemm::has_opt_gemm(arm_gemm_expected_wf, args, {})), - "We could not find an optimized kernel for S8 input and S8 output"); + "We could not find an optimized kernel for QASYMM8_SIGNED input and QASYMM8_SIGNED output"); + } + else + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG(true, "Only S32 output supported for S8 input"); } break; #endif /* __aarch64__ */ @@ -1130,6 +1151,23 @@ Status CpuGemmAssemblyDispatch::validate( a->data_type() == DataType::QASYMM8 && (d->data_type() != DataType::QASYMM8 && d->data_type() != DataType::S32 && d->data_type() != DataType::F32), "Only QASYMM8/S32/F32 output supported for QASYMM8 input"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(a->data_type() == DataType::QASYMM8 && b->data_type() == DataType::QASYMM8_SIGNED && + d->data_type() == DataType::F32, + "QASYMM8 input with QASYMM8_SIGNED weights and F32 output is not supported"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(a->data_type() == DataType::QASYMM8_SIGNED && + (d->data_type() != DataType::QASYMM8_SIGNED && + d->data_type() != DataType::S32 && d->data_type() != DataType::F32 && + d->data_type() != DataType::F16), + "Only QASYMM8_SIGNED/S32/F32/F16 output supported for QASYMM8_SIGNED input"); + // TODO(COMPMID-9268): Temporary guard until TensorInfo rejects inconsistent quantization metadata + // centrally. QASYMM8_SIGNED tensors can currently carry multiple quantization scales even though + // this DequantizeFloat path only supports uniform weight scales. Remove this local check once + // COMPMID-9268 is resolved. + ARM_COMPUTE_RETURN_ERROR_ON_MSG( + b->data_type() == DataType::QASYMM8_SIGNED && + (d->data_type() == DataType::F32 || d->data_type() == DataType::F16) && + b->quantization_info().scale().size() > 1, + "Per-channel QASYMM8_SIGNED weight scales are not supported for dequantized assembly output"); arm_compute::WeightFormat expected_weight_format = arm_compute::WeightFormat::UNSPECIFIED; const Status ret = CpuGemmAssemblyDispatch::has_opt_impl(expected_weight_format, a, b, c, d, info); if (bool(ret) && expected_weight_format != arm_compute::WeightFormat::ANY) @@ -1173,14 +1211,7 @@ void CpuGemmAssemblyDispatch::configure( case DataType::QASYMM8: if (b->data_type() == DataType::S8 || b->data_type() == DataType::QASYMM8_SIGNED) { - if (d->data_type() == DataType::F32) - { - create_arm_gemm_dequant(_arm_gemm, a, b, c, d, act, info); - } - else - { - create_arm_gemm_quant(_arm_gemm, a, b, c, d, act, info); - } + create_arm_gemm_quant(_arm_gemm, a, b, c, d, act, info); } else if (d->data_type() == DataType::S32) { @@ -1197,17 +1228,17 @@ void CpuGemmAssemblyDispatch::configure( { create_arm_gemm(_arm_gemm, a, b, c, d, act, info); } - else if (d->data_type() == DataType::F32) + else if (a->data_type() == DataType::QASYMM8_SIGNED && d->data_type() == DataType::F32) { create_arm_gemm_dequant(_arm_gemm, a, b, c, d, act, info); } #if defined(ENABLE_FP16_KERNELS) - else if (d->data_type() == DataType::F16) + else if (a->data_type() == DataType::QASYMM8_SIGNED && d->data_type() == DataType::F16) { create_arm_gemm_dequant(_arm_gemm, a, b, c, d, act, info); } #endif /* defined(ENABLE_FP16_KERNELS) */ - else + else if (a->data_type() == DataType::QASYMM8_SIGNED) { create_arm_gemm_quant(_arm_gemm, a, b, c, d, act, info); } diff --git a/src/cpu/operators/internal/CpuGemmAssemblyDispatch.h b/src/cpu/operators/internal/CpuGemmAssemblyDispatch.h index 9b4b15d0dbc..cfad8328a37 100644 --- a/src/cpu/operators/internal/CpuGemmAssemblyDispatch.h +++ b/src/cpu/operators/internal/CpuGemmAssemblyDispatch.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2025 Arm Limited. + * Copyright (c) 2018-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -29,6 +29,8 @@ #include "src/core/common/Macros.h" #include "src/cpu/ICpuOperator.h" +#include + namespace arm_compute { namespace cpu @@ -65,7 +67,9 @@ struct AsmGemmInfo * @note This flag will be silently ignored (assumed to be false) when the weight_format is a fixed format. Because * fixed format kernels do not accept weights (B) with any prior transformations */ - bool transpose_b{false}; + bool transpose_b{false}; + int32_t dequant_a_offset{0}; // input zero-point for DequantizeFloat path (handled in kernel) + int32_t dequant_b_offset{0}; // weight zero-point for DequantizeFloat path (handled in kernel) }; /** Assembly kernel glue */ diff --git a/src/runtime/NEON/functions/NEConvolutionLayer.cpp b/src/runtime/NEON/functions/NEConvolutionLayer.cpp index c4f28079852..d9061b4c355 100644 --- a/src/runtime/NEON/functions/NEConvolutionLayer.cpp +++ b/src/runtime/NEON/functions/NEConvolutionLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2021, 2023-2025 Arm Limited. + * Copyright (c) 2017-2021, 2023-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -83,7 +83,6 @@ void NEConvolutionLayer::configure(ITensor *input, ARM_COMPUTE_LOG_PARAMS(input, weights, biases, output, conv_info, weights_info, dilation, act_info, enable_fast_math, num_groups); - const Conv2dInfo info(conv_info, dilation, act_info, enable_fast_math, num_groups); switch (cpu::CpuConv2d::get_convolution_method(input->info(), weights->info(), output->info(), conv_info, weights_info, dilation, act_info, enable_fast_math)) { @@ -134,7 +133,6 @@ Status NEConvolutionLayer::validate(const ITensorInfo *input, unsigned int num_groups) { ARM_COMPUTE_TRACE_EVENT(ARM_COMPUTE_PROF_CAT_CPU, ARM_COMPUTE_PROF_LVL_CPU, "NEConvolutionLayer::validate"); - const Conv2dInfo info(conv_info, dilation, act_info, enable_fast_math, num_groups); ARM_COMPUTE_RETURN_ERROR_ON_MSG(!weights->are_values_constant(), "Dynamic weights are not supported"); ARM_COMPUTE_RETURN_ERROR_ON_DYNAMIC_SHAPE(input, weights, biases, output); diff --git a/tests/validation/NEON/ConvolutionLayer.cpp b/tests/validation/NEON/ConvolutionLayer.cpp index ac494c5cdd2..474aa243df4 100644 --- a/tests/validation/NEON/ConvolutionLayer.cpp +++ b/tests/validation/NEON/ConvolutionLayer.cpp @@ -42,6 +42,7 @@ #include "tests/NEON/Accessor.h" #include "tests/validation/fixtures/ConvolutionLayerFixture.h" #include "tests/validation/fixtures/WinogradConvolutionLayerFixture.h" +#include "tests/validation/Helpers.h" #include "tests/validation/Validation.h" namespace arm_compute @@ -306,11 +307,168 @@ FIXTURE_DATA_TEST_CASE( { validate(Accessor(_target), _reference, rel_tolerance_f32, tolerance_num_dequantize_f32, float(abs_tolerance_f32)); } -TEST_SUITE_END() // QASYMM8_SIGNED -#endif // #ifdef __aarch64__ // clang-format on // *INDENT-ON* + +TEST_SUITE(DequantizeFloat) + +const auto DequantizeFloatOffsetsDataset = + zip(make("InputQI", + {QuantizationInfo(0.25f, 0), QuantizationInfo(0.25f, -10), QuantizationInfo(0.25f, 12), + QuantizationInfo(0.25f, 0), QuantizationInfo(0.25f, 0), QuantizationInfo(0.25f, -20), + QuantizationInfo(0.01f, 100)}), + make("WeightsQI", + {QuantizationInfo(0.125f, 0), QuantizationInfo(0.125f, 0), QuantizationInfo(0.125f, 0), + QuantizationInfo(0.125f, 5), QuantizationInfo(0.125f, -8), QuantizationInfo(0.125f, 10), + QuantizationInfo(0.01f, -100)})); + +const auto DequantizeFloatBasicActivationDataset = + make("ActivationInfo", {ActivationLayerInfo(), ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::RELU)}); + +/** Validate accepts QASYMM8_SIGNED to floating-point dequantized convolutions with zero and non-zero offsets, and rejects + * mixed-sign input/weight configurations. + * Shapes use NHWC: [C, W, H] for input/output, [Cin, Kw, Kh, Cout] for weights. */ +DATA_TEST_CASE(Validate, + framework::DatasetMode::ALL, + zip(make("SrcInfo", + {TensorInfo(TensorShape(16U, 8U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(32U, 4U, 4U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 8U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 8U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 8U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 8U, 8U), 1, DataType::QASYMM8, DataLayout::NHWC)}), + make("WgtInfo", + {TensorInfo(TensorShape(16U, 3U, 3U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(32U, 1U, 1U, 16U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 3U, 3U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 3U, 3U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 3U, 3U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 3U, 3U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC)}), + make("BiasInfo", + {TensorInfo(TensorShape(8U), 1, DataType::F32), TensorInfo(TensorShape(16U), 1, DataType::F32), + TensorInfo(TensorShape(8U), 1, DataType::F32), TensorInfo(TensorShape(8U), 1, DataType::F32), + TensorInfo(TensorShape(8U), 1, DataType::F16), TensorInfo(TensorShape(8U), 1, DataType::F32)}), + make("DstInfo", + {TensorInfo(TensorShape(8U, 6U, 6U), 1, DataType::F32, DataLayout::NHWC), + TensorInfo(TensorShape(16U, 4U, 4U), 1, DataType::F32, DataLayout::NHWC), + TensorInfo(TensorShape(8U, 6U, 6U), 1, DataType::F32, DataLayout::NHWC), + TensorInfo(TensorShape(8U, 6U, 6U), 1, DataType::F32, DataLayout::NHWC), + TensorInfo(TensorShape(8U, 6U, 6U), 1, DataType::F16, DataLayout::NHWC), + TensorInfo(TensorShape(8U, 6U, 6U), 1, DataType::F32, DataLayout::NHWC)}), + make("SrcOffset", {0, 0, 5, -10, 0, 0}), + make("WgtOffset", {0, 0, 0, 3, 0, 0}), + make("Expected", {true, true, true, true, true, false})), + src_info_const, + wgt_info_const, + bias_info_const, + dst_info_const, + src_offset, + wgt_offset, + expected) +{ + TensorInfo src_info = src_info_const; + TensorInfo wgt_info = wgt_info_const; + TensorInfo bias_info = bias_info_const; + TensorInfo dst_info = dst_info_const; + src_info.set_quantization_info(QuantizationInfo(0.25f, src_offset)); + wgt_info.set_quantization_info(QuantizationInfo(0.125f, wgt_offset)); + + const Status s = NEConvolutionLayer::validate(&src_info, &wgt_info, &bias_info, &dst_info, + PadStrideInfo(1, 1, 0, 0), WeightsInfo(), Size2D(1U, 1U), + ActivationLayerInfo(), false /*fast_math*/, 1 /*num_groups*/); + const bool expected_result = expected && cpu_supports_dtypes({src_info.data_type(), wgt_info.data_type(), + bias_info.data_type(), dst_info.data_type()}); + ARM_COMPUTE_EXPECT(bool(s) == expected_result, framework::LogLevel::ERRORS); +} + +/** Verify method selection for NHWC QASYMM8_SIGNED to floating-point dequantized convolution. */ +TEST_CASE(GetConvMethod, framework::DatasetMode::ALL) +{ + const QuantizationInfo qi(0.25f, 0); + for (const DataType data_type : {DataType::F32, DataType::F16}) + { + TensorInfo src_info(TensorShape(16U, 8U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC); + TensorInfo wgt_info(TensorShape(16U, 3U, 3U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC); + TensorInfo dst_info(TensorShape(8U, 6U, 6U), 1, data_type, DataLayout::NHWC); + src_info.set_quantization_info(qi); + wgt_info.set_quantization_info(qi); + + const ConvolutionMethod m = NEConvolutionLayer::get_convolution_method( + &src_info.clone()->set_is_resizable(true), &wgt_info.clone()->set_is_resizable(true), + &dst_info.clone()->set_is_resizable(true), PadStrideInfo(1, 1, 0, 0), WeightsInfo(), Size2D(1U, 1U), + ActivationLayerInfo(), false /*fast_math*/); + const bool expected_direct_conv = cpu_supports_dtypes({src_info.data_type(), wgt_info.data_type(), data_type}); + ARM_COMPUTE_EXPECT((m == ConvolutionMethod::GEMM_CONV2D) == expected_direct_conv, framework::LogLevel::ERRORS); + } +} + +TEST_CASE(ValidateRejectsPerChannelWeightScales, framework::DatasetMode::ALL) +{ + const QuantizationInfo src_qinfo(0.25f, 0); + const QuantizationInfo weights_qinfo({0.125f, 0.25f, 0.5f, 1.f, 2.f, 4.f, 8.f, 16.f}, {0, 0, 0, 0, 0, 0, 0, 0}); + + TensorInfo src_info(TensorShape(16U, 8U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC); + TensorInfo wgt_info(TensorShape(16U, 3U, 3U, 8U), 1, DataType::QASYMM8_SIGNED, DataLayout::NHWC); + TensorInfo bias_info(TensorShape(8U), 1, DataType::F32); + TensorInfo dst_info(TensorShape(8U, 6U, 6U), 1, DataType::F32, DataLayout::NHWC); + + src_info.set_quantization_info(src_qinfo); + wgt_info.set_quantization_info(weights_qinfo); + + const Status status = NEConvolutionLayer::validate(&src_info, &wgt_info, &bias_info, &dst_info, + PadStrideInfo(1, 1, 0, 0), WeightsInfo(), Size2D(1U, 1U), + ActivationLayerInfo(), false /*fast_math*/, 1 /*num_groups*/); + ARM_COMPUTE_EXPECT(!bool(status), framework::LogLevel::ERRORS); +} + +using NEDequantizeFloatConvolutionFixture = DequantizeFloatConvolutionFixture; +#if defined(ARM_COMPUTE_ENABLE_FP16) +using NEDequantizeFloatF16ConvolutionFixture = + DequantizeFloatConvolutionFixture; +#endif // ARM_COMPUTE_ENABLE_FP16 + +/** SmallConvolutionLayerDataset covers 1x1, 2x2, 3x3, 4x4, 5x5 and 5x7 convolutions, including stride, + * padding, asymmetric padding and batch cases. */ +FIXTURE_DATA_TEST_CASE(RunSmall, + NEDequantizeFloatConvolutionFixture, + framework::DatasetMode::ALL, + combine(datasets::SmallConvolutionLayerDataset(), + make("ReshapeWeights", {true}), + make("DataLayout", {DataLayout::NHWC}), + DequantizeFloatBasicActivationDataset, + DequantizeFloatOffsetsDataset)) +{ + validate(Accessor(_target), _reference, rel_tolerance_f32, tolerance_num_dequantize_f32, float(abs_tolerance_f32)); +} + +#if defined(ARM_COMPUTE_ENABLE_FP16) +FIXTURE_DATA_TEST_CASE(RunF16NonZeroOffsets, + NEDequantizeFloatF16ConvolutionFixture, + framework::DatasetMode::ALL, + combine(datasets::SmallConvolutionLayerDataset(), + make("ReshapeWeights", {true}), + make("DataLayout", {DataLayout::NHWC}), + make("ActivationInfo", {ActivationLayerInfo()}), + zip(make("InputQI", {QuantizationInfo(0.25f, -10)}), + make("WeightsQI", {QuantizationInfo(0.125f, 5)})))) +{ + if (CPUInfo::get().has_fp16()) + { + validate(Accessor(_target), _reference, rel_tolerance_f16, 0.f, half(abs_tolerance_f16)); + } + else + { + ARM_COMPUTE_TEST_WARNING("Device does not support fp16 vector operations. Test SKIPPED."); + framework::ARM_COMPUTE_PRINT_WARNING(); + } +} +#endif // ARM_COMPUTE_ENABLE_FP16 + +TEST_SUITE_END() // DequantizeFloat +TEST_SUITE_END() // QASYMM8_SIGNED +#endif // #ifdef __aarch64__ + TEST_SUITE_END() // ConvolutionLayer /* diff --git a/tests/validation/NEON/GEMMLowp.cpp b/tests/validation/NEON/GEMMLowp.cpp index eab75dee07b..edd134eae66 100644 --- a/tests/validation/NEON/GEMMLowp.cpp +++ b/tests/validation/NEON/GEMMLowp.cpp @@ -541,20 +541,26 @@ DATA_TEST_CASE( 1, DataType::QASYMM8_SIGNED, QuantizationInfo(1.f / 255, 10)), // Invalid types + TensorInfo(TensorShape(16U, 32U), 1, DataType::QASYMM8_SIGNED, QuantizationInfo(1.f / 255, 10)), }), make("InputBInfo", { TensorInfo(TensorShape(64U, 16U), 1, DataType::QASYMM8_SIGNED, QuantizationInfo(1.f / 256, 10)), TensorInfo(TensorShape(64U, 16U), 1, DataType::QASYMM8_SIGNED, QuantizationInfo(1.f / 256, 10)), TensorInfo(TensorShape(64U, 16U), 1, DataType::QASYMM8, QuantizationInfo(1.f / 256, 10)), + TensorInfo(TensorShape(64U, 16U), + 1, + DataType::QASYMM8_SIGNED, + QuantizationInfo({1.f / 256, 2.f / 256}, {10, 10})), }), make("OutputInfo", { TensorInfo(TensorShape(64U, 32U), 1, DataType::F32), TensorInfo(TensorShape(64U, 32U), 1, DataType::F32), TensorInfo(TensorShape(64U, 32U), 1, DataType::F32), + TensorInfo(TensorShape(64U, 32U), 1, DataType::F32), }), - make("Expected", {true, true, false})), + make("Expected", {true, true, false, false})), a_info, b_info, output_info, diff --git a/tests/validation/fixtures/ConvolutionLayerFixture.h b/tests/validation/fixtures/ConvolutionLayerFixture.h index 2ccefa54575..3e2e396279e 100644 --- a/tests/validation/fixtures/ConvolutionLayerFixture.h +++ b/tests/validation/fixtures/ConvolutionLayerFixture.h @@ -1320,6 +1320,149 @@ class HasOptImplFixture : public framework::Fixture }; #endif // ARM_COMPUTE_ENABLE_FIXED_FORMAT_KERNELS +/** Fixture for testing QASYMM8_SIGNED to floating-point dequantized convolution. + * + * Supports NHWC, no dilation, arbitrary quantization offsets. + * Compares target output against a CPU reference using the same floating-point output type. + */ +template +class DequantizeFloatConvolutionFixture : public framework::Fixture +{ +public: + void setup(TensorShape input_shape_nchw, + TensorShape weights_shape_nchw, + TensorShape bias_shape, + TensorShape output_shape_nchw, + PadStrideInfo conv_info, + Size2D dilation, + bool reshape_weights, + DataLayout data_layout, + ActivationLayerInfo act_info, + QuantizationInfo input_qi = QuantizationInfo(0.25f, 0), + QuantizationInfo weights_qi = QuantizationInfo(0.125f, 0)) + { + if (std::is_same::value && !CPUInfo::get().has_fp16()) + { + return; + } + + ARM_COMPUTE_EXPECT_EQUAL(data_layout, DataLayout::NHWC, framework::LogLevel::ERRORS); + _data_layout = data_layout; + _input_qi = input_qi; + _weights_qi = weights_qi; + + _hash = static_cast(input_shape_nchw[0] + input_shape_nchw[1] + input_shape_nchw[2] + + weights_shape_nchw[0] + weights_shape_nchw[1] + + static_cast(input_qi.uniform().offset) + + static_cast(weights_qi.uniform().offset)); + + _target = compute_target(input_shape_nchw, weights_shape_nchw, bias_shape, output_shape_nchw, conv_info, + reshape_weights, dilation, act_info); + _reference = compute_reference(input_shape_nchw, weights_shape_nchw, bias_shape, output_shape_nchw, conv_info, + dilation, act_info); + } + +protected: + template + void fill(U &&tensor, int i) + { + switch (tensor.data_type()) + { + case DataType::QASYMM8_SIGNED: + { + std::uniform_int_distribution dist(-127, 127); + library->fill(tensor, dist, i); + break; + } + default: + library->fill_tensor_uniform(tensor, i); + } + } + + TensorType compute_target(TensorShape input_shape, + TensorShape weights_shape, + const TensorShape &bias_shape, + TensorShape output_shape, + const PadStrideInfo &conv_info, + bool reshape_weights, + const Size2D &dilation, + const ActivationLayerInfo act_info) + { + // Convert NCHW→NHWC + permute(input_shape, PermutationVector(2U, 0U, 1U)); + permute(weights_shape, PermutationVector(2U, 0U, 1U)); + permute(output_shape, PermutationVector(2U, 0U, 1U)); + + const int idx_w = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::WIDTH); + const int idx_h = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::HEIGHT); + + WeightsInfo weights_info(!reshape_weights, weights_shape[idx_w], weights_shape[idx_h], weights_shape[3]); + + TensorType src = create_tensor(input_shape, DataType::QASYMM8_SIGNED, 1, _input_qi, _data_layout); + TensorType wgt = + create_tensor(weights_shape, DataType::QASYMM8_SIGNED, 1, _weights_qi, _data_layout); + TensorType bias = + create_tensor(bias_shape, output_data_type(), 1, QuantizationInfo(), _data_layout); + TensorType dst = + create_tensor(output_shape, output_data_type(), 1, QuantizationInfo(), _data_layout); + + FunctionType conv; + conv.configure(&src, &wgt, &bias, &dst, conv_info, weights_info, dilation, act_info, false /*enable_fast_math*/, + 1 /*num_groups*/); + + ARM_COMPUTE_ASSERT(src.info()->is_resizable()); + src.allocator()->allocate(); + wgt.allocator()->allocate(); + bias.allocator()->allocate(); + dst.allocator()->allocate(); + + fill(AccessorType(src), 0 + _hash); + fill(AccessorType(wgt), 1 + _hash); + fill(AccessorType(bias), 2 + _hash); + + conv.run(); + return dst; + } + + SimpleTensor compute_reference(TensorShape input_shape, + TensorShape weights_shape, + const TensorShape &bias_shape, + TensorShape output_shape, + const PadStrideInfo &conv_info, + const Size2D &dilation, + const ActivationLayerInfo act_info) + { + // Reference runs in NCHW with quantized input/weights and floating-point output. + SimpleTensor src_q{input_shape, DataType::QASYMM8_SIGNED, 1, _input_qi}; + SimpleTensor wgt_q{weights_shape, DataType::QASYMM8_SIGNED, 1, _weights_qi}; + SimpleTensor bias{bias_shape, output_data_type()}; + + fill(src_q, 0 + _hash); + fill(wgt_q, 1 + _hash); + fill(bias, 2 + _hash); + + auto conv = reference::convolution_layer( + src_q, wgt_q, bias, output_shape, conv_info, dilation, 1 /*num_groups*/, QuantizationInfo()); + if (act_info.enabled()) + { + conv = reference::activation_layer(conv, act_info); + } + return conv; + } + + static DataType output_data_type() + { + return std::is_same::value ? DataType::F16 : DataType::F32; + } + + TensorType _target{}; + SimpleTensor _reference{}; + DataLayout _data_layout{DataLayout::NHWC}; + QuantizationInfo _input_qi{}; + QuantizationInfo _weights_qi{}; + int32_t _hash{0}; +}; + } // namespace validation } // namespace test } // namespace arm_compute diff --git a/tests/validation/reference/Convolution3d.h b/tests/validation/reference/Convolution3d.h index 5c3332b0667..0714dc76222 100644 --- a/tests/validation/reference/Convolution3d.h +++ b/tests/validation/reference/Convolution3d.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2021, 2023, 2025 Arm Limited. + * Copyright (c) 2017-2021, 2023, 2025-2026 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -212,7 +212,7 @@ inline void convolution3d(const SimpleTensor &in, } } - if (std::is_same::value) + if (validation::is_floating_point::value) { const float scale = input_scale * weights_scale; @@ -223,8 +223,8 @@ inline void convolution3d(const SimpleTensor &in, b = static_cast(*b_ptr) * scale; } else - { // TB == float - // New path: bias already in float units + { + // Bias already uses the floating-point output units. b = static_cast(*b_ptr); } diff --git a/tests/validation/reference/ConvolutionLayer.cpp b/tests/validation/reference/ConvolutionLayer.cpp index 08a68f8d310..ba0b5591458 100644 --- a/tests/validation/reference/ConvolutionLayer.cpp +++ b/tests/validation/reference/ConvolutionLayer.cpp @@ -127,12 +127,17 @@ SimpleTensor convolution_layer(const SimpleTensor &src, DataType dst_dt = src.data_type(); QuantizationInfo dst_qinfo = out_quant_info; - // For float output we override both + // For floating-point output we override both if (std::is_same::value) { dst_dt = DataType::F32; dst_qinfo = QuantizationInfo(); // no quantization for F32 } + else if (std::is_same::value) + { + dst_dt = DataType::F16; + dst_qinfo = QuantizationInfo(); // no quantization for F16 + } SimpleTensor dst{output_shape, dst_dt, 1, dst_qinfo}; @@ -148,6 +153,15 @@ template SimpleTensor convolution_layer(con const Size2D &dilation, unsigned int num_groups, QuantizationInfo out_quant_info); +// Dequantize i8+i8->F16 +template SimpleTensor convolution_layer(const SimpleTensor &src, + const SimpleTensor &weights, + const SimpleTensor &bias, + const TensorShape &output_shape, + const PadStrideInfo &info, + const Size2D &dilation, + unsigned int num_groups, + QuantizationInfo out_quant_info); // Dequantize u8+u8->F32 template SimpleTensor convolution_layer(const SimpleTensor &src, const SimpleTensor &weights,