From 6fdfaa856b1eb69d6afbef5727b99756912fc6fa Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Wed, 29 Nov 2017 14:27:24 +0000 Subject: COMPMID-713: Address failures in OCLGrind for CLDirectConvolution -Changes way of clamping in the kernel side. -Fills padding with quantized values Change-Id: I94d17c341fd637fbb24390722162b551b62d16cb Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/111114 Reviewed-by: Anthony Barbier Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com --- .../CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl | 4 +--- src/runtime/CL/functions/CLDirectConvolutionLayer.cpp | 8 +++++++- tests/validation/CL/DirectConvolutionLayer.cpp | 2 +- tests/validation/CPP/ConvolutionLayer.cpp | 3 +-- tests/validation/fixtures/DirectConvolutionLayerFixture.h | 4 ++-- 5 files changed, 12 insertions(+), 9 deletions(-) diff --git a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl index cbe826639d..d0cf032ac1 100644 --- a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl +++ b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl @@ -243,9 +243,7 @@ __kernel void direct_convolution_1x1_3x3_5x5_quantized( pixels0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(pixels0, output_multiplier, output_shift, 8); pixels0 = pixels0 + output_offset; - pixels0 = max(pixels0, 0); - pixels0 = min(pixels0, 255); - vstore8(convert_uchar8(pixels0), 0, (__global uchar *)dst.ptr); + vstore8(convert_uchar8_sat(pixels0), 0, (__global uchar *)dst.ptr); } #endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) diff --git a/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp b/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp index 8ef7068b88..759692505b 100644 --- a/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp +++ b/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp @@ -23,6 +23,7 @@ */ #include "arm_compute/runtime/CL/functions/CLDirectConvolutionLayer.h" +#include "arm_compute/core/CL/ICLTensor.h" #include "arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h" #include "arm_compute/core/PixelValue.h" #include "arm_compute/core/Utils.h" @@ -45,7 +46,12 @@ void CLDirectConvolutionLayer::configure(ICLTensor *input, const ICLTensor *weig _direct_conv_kernel.configure(input, weights, biases, output, conv_info); // Configure border handler - _input_border_handler.configure(input, _direct_conv_kernel.border_size(), BorderMode::CONSTANT, PixelValue(0)); + PixelValue &&zero_value(0.f); + if(is_data_type_quantized_asymmetric(input->info()->data_type())) + { + zero_value = PixelValue(static_cast(input->info()->quantization_info().offset)); + } + _input_border_handler.configure(input, _direct_conv_kernel.border_size(), BorderMode::CONSTANT, zero_value); } Error CLDirectConvolutionLayer::validate(const ITensorInfo *input, const ITensorInfo *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info) diff --git a/tests/validation/CL/DirectConvolutionLayer.cpp b/tests/validation/CL/DirectConvolutionLayer.cpp index 08d53d5af3..e6a196ae46 100644 --- a/tests/validation/CL/DirectConvolutionLayer.cpp +++ b/tests/validation/CL/DirectConvolutionLayer.cpp @@ -196,7 +196,7 @@ using CLDirectConvolutionLayerQuantizedFixture = DirectConvolutionValidationQuan TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerQuantizedFixture, framework::DatasetMode::ALL, combine(combine(data, framework::dataset::make("DataType", DataType::QASYMM8)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255, 127) }))) + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255, 10) }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qasymm8); diff --git a/tests/validation/CPP/ConvolutionLayer.cpp b/tests/validation/CPP/ConvolutionLayer.cpp index a767912879..3884a930a2 100644 --- a/tests/validation/CPP/ConvolutionLayer.cpp +++ b/tests/validation/CPP/ConvolutionLayer.cpp @@ -210,8 +210,7 @@ void convolution3d(const SimpleTensor &in, const SimpleTensor acc = asymm_rounding_divide_by_pow2(asymm_int_mult(acc, output_multiplier), output_shift); acc += output_offset; - acc = std::max(acc, 0); - acc = std::min(acc, 255); + acc = clamp(acc, 0, 255); // Store the result *out_ptr = acc; diff --git a/tests/validation/fixtures/DirectConvolutionLayerFixture.h b/tests/validation/fixtures/DirectConvolutionLayerFixture.h index 1ec4d31304..b78f13acd1 100644 --- a/tests/validation/fixtures/DirectConvolutionLayerFixture.h +++ b/tests/validation/fixtures/DirectConvolutionLayerFixture.h @@ -73,7 +73,7 @@ protected: { case DataType::QASYMM8: { - std::uniform_int_distribution distribution(0, 10); + std::uniform_int_distribution distribution(0, 50); library->fill(tensor, distribution, i); break; } @@ -86,7 +86,7 @@ protected: } case DataType::S32: { - std::uniform_int_distribution distribution(-1000, 1000); + std::uniform_int_distribution distribution(-5, 5); library->fill(tensor, distribution, i); break; } -- cgit v1.2.1