diff options
| author | Adnan AlSinan <adnan.alsinan@arm.com> | 2022-09-22 01:21:42 +0100 |
|---|---|---|
| committer | Adnan AlSinan <adnan.alsinan@arm.com> | 2022-10-06 09:04:41 +0000 |
| commit | 3bedd2f031680f53e2982638adfe99a29dca8d06 (patch) | |
| tree | 49a3723083c847928b170a8235cb5221b2228a80 | |
| parent | db14af697b934d684d8b3d63a00ad5bea5c07bfb (diff) | |
| download | ComputeLibrary-3bedd2f031680f53e2982638adfe99a29dca8d06.tar.gz | |
Improve start-up time in gemmlowp reshaped rhs only.
- Does not Pass RESULT_MULTIPLIER, RESULT_SHIFT When PER_CHANNEL_QUANTIZATION is defined.
Resolves COMPMID-5499
Signed-off-by: Adnan AlSinan <adnan.alsinan@arm.com>
Change-Id: Ie433eaf83c003a6d5ccfeb89eb2783528dc2b48e
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/8316
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Benchmark: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/common/gemmlowp.cl | 2 | ||||
| -rw-r--r-- | src/gpu/cl/kernels/ClGemmLowpMatrixMultiplyReshapedOnlyRhsKernel.cpp | 13 |
2 files changed, 12 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/common/gemmlowp.cl b/src/core/CL/cl_kernels/common/gemmlowp.cl index 773e0333b2..62c4cd31f5 100644 --- a/src/core/CL/cl_kernels/common/gemmlowp.cl +++ b/src/core/CL/cl_kernels/common/gemmlowp.cl @@ -1748,7 +1748,7 @@ __kernel void gemmlowp_offset_contribution_quantize_down_fixedpoint(TENSOR3D_DEC #else // defined(PER_CHANNEL_QUANTIZATION) #if RESULT_SHIFT < 0 - in_s32 = ASYMM_MULT_BY_QUANT_MULTIPLIER_GREATER_THAN_ONE(in_s32, RESULT_MULTIPLIER, RESULT_SHIFT, VEC_SIZE); + in_s32 = ASYMM_MULT_BY_QUANT_MULTIPLIER_GREATER_THAN_ONE(in_s32, RESULT_MULTIPLIER, RESULT_SHIFT, VEC_SIZE); #else // RESULT_SHIFT >= 0 in_s32 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(in_s32, RESULT_MULTIPLIER, RESULT_SHIFT, VEC_SIZE); #endif // RESULT_SHIFT < 0 diff --git a/src/gpu/cl/kernels/ClGemmLowpMatrixMultiplyReshapedOnlyRhsKernel.cpp b/src/gpu/cl/kernels/ClGemmLowpMatrixMultiplyReshapedOnlyRhsKernel.cpp index 386c13eb92..90084ea97f 100644 --- a/src/gpu/cl/kernels/ClGemmLowpMatrixMultiplyReshapedOnlyRhsKernel.cpp +++ b/src/gpu/cl/kernels/ClGemmLowpMatrixMultiplyReshapedOnlyRhsKernel.cpp @@ -376,8 +376,17 @@ void ClGemmLowpMatrixMultiplyReshapedOnlyRhsKernel::configure(const CLCompileCon build_opts.add_option("-DK_OFFSET=" + support::cpp11::to_string(a_offset * b_offset * src0->dimension(0))); build_opts.add_option_if(bias != nullptr, "-DADD_BIAS"); build_opts.add_option("-DRESULT_OFFSET=" + support::cpp11::to_string(output_stage.gemmlowp_offset)); - build_opts.add_option("-DRESULT_MULTIPLIER=" + support::cpp11::to_string(output_stage.gemmlowp_multipliers[0])); - build_opts.add_option("-DRESULT_SHIFT=" + support::cpp11::to_string(output_stage.gemmlowp_shifts[0])); + // In case of _is_quantized_per_channel, RESULT_MULTIPLIER and RESULT_SHIFT are not utilized, but they are passed as a part of T_QUANTIZE8 macro. + if(!_is_quantized_per_channel) + { + build_opts.add_option("-DRESULT_MULTIPLIER=" + support::cpp11::to_string(output_stage.gemmlowp_multipliers[0])); + build_opts.add_option("-DRESULT_SHIFT=" + support::cpp11::to_string(output_stage.gemmlowp_shifts[0])); + } + else + { + build_opts.add_option("-DRESULT_MULTIPLIER=0"); + build_opts.add_option("-DRESULT_SHIFT=0"); + } build_opts.add_option_if(_is_quantized_per_channel, "-DPER_CHANNEL_QUANTIZATION"); const int min = output_stage.gemmlowp_min_bound; |