From 97e258000f99c6a5f872f4e8968eaf6a93de2cf7 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Thu, 25 Mar 2021 12:37:45 +0000 Subject: Check biases pointer before referencing in CLDirectConvolutionLayer The biases input can be nullptr, hence we need to check before referencing. A test is also added to ensure a successful configure and run of Direct Convolution when there is no bias. Resolves: COMPMID-4315 Change-Id: I23223efd6ced81215aff490221fb4606945c139b Signed-off-by: Michele Di Giorgio Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5322 Tested-by: Arm Jenkins Reviewed-by: James Conroy Reviewed-by: Georgios Pinitas Comments-Addressed: Arm Jenkins --- src/core/CL/CLKernelLibrary.cpp | 3 +- src/core/CL/cl_kernels/direct_convolution.cl | 3 +- .../CL/functions/CLDirectConvolutionLayer.cpp | 4 +- tests/validation/CL/DirectConvolutionLayer.cpp | 43 ++++++++++++++++++++++ tests/validation/NEON/DirectConvolutionLayer.cpp | 43 ++++++++++++++++++++++ 5 files changed, 92 insertions(+), 4 deletions(-) diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 726efa3575..2652884912 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -901,12 +901,13 @@ std::string CLKernelLibrary::get_program_name(const std::string &kernel_name) co void CLKernelLibrary::init(std::string kernel_path, cl::Context context, cl::Device device) { _compile_context = CLCompileContext(context, device); - _kernel_path = kernel_path; + _kernel_path = kernel_path + "/"; } void CLKernelLibrary::set_kernel_path(const std::string &kernel_path) { _kernel_path = std::move(kernel_path); + _kernel_path += "/"; } cl::Context &CLKernelLibrary::context() diff --git a/src/core/CL/cl_kernels/direct_convolution.cl b/src/core/CL/cl_kernels/direct_convolution.cl index 1de3737965..dde024faa4 100644 --- a/src/core/CL/cl_kernels/direct_convolution.cl +++ b/src/core/CL/cl_kernels/direct_convolution.cl @@ -105,8 +105,9 @@ __kernel void direct_convolution_nhwc( TENSOR4D(src, SRC_TENSOR_TYPE), TENSOR4D(dst, DST_TENSOR_TYPE), - TENSOR4D(wei, WEI_TENSOR_TYPE), + TENSOR4D(wei, WEI_TENSOR_TYPE) #if defined(HAS_BIAS) + , VECTOR_DECLARATION(bia) #endif // defined(HAS_BIAS) ) diff --git a/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp b/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp index d60d11aa5f..74867ff64f 100644 --- a/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp +++ b/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp @@ -58,7 +58,7 @@ void CLDirectConvolutionLayer::configure(ICLTensor *input, const ICLTensor *weig void CLDirectConvolutionLayer::configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output, const PadStrideInfo &conv_info, const ActivationLayerInfo &act_info) { - ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); + ARM_COMPUTE_ERROR_ON_NULLPTR(input, weights, output); _impl->src = input; _impl->weights = weights; @@ -66,7 +66,7 @@ void CLDirectConvolutionLayer::configure(const CLCompileContext &compile_context _impl->dst = output; _impl->op = std::make_unique(); - _impl->op->configure(compile_context, _impl->src->info(), _impl->weights->info(), _impl->biases->info(), _impl->dst->info(), conv_info, act_info); + _impl->op->configure(compile_context, input->info(), weights->info(), (biases != nullptr) ? biases->info() : nullptr, output->info(), conv_info, act_info); } Status 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 4671d8c3ec..c01234020f 100644 --- a/tests/validation/CL/DirectConvolutionLayer.cpp +++ b/tests/validation/CL/DirectConvolutionLayer.cpp @@ -87,6 +87,49 @@ const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo TEST_SUITE(CL) TEST_SUITE(DirectConvolutionLayer) +/** Check whether the configuration of a Direct Convolution layer with no + * bias leads to a successful execution. + */ +TEST_CASE(NoBias, framework::DatasetMode::PRECOMMIT) +{ + const auto src_shape = TensorShape(27U, 13U, 2U); + const auto weights_shape = TensorShape(3U, 3U, 2U, 4U); + const auto bias_shape = TensorShape(4U); + const auto dst_shape = TensorShape(25U, 11U, 4U); + constexpr auto dt = DataType::F32; + + auto src = create_tensor(src_shape, dt); + auto weights = create_tensor(weights_shape, dt); + auto dst = create_tensor(dst_shape, dt); + + const auto conv_info = PadStrideInfo(1, 1, 0, 0); + + // Create Direct Convolution function + CLDirectConvolutionLayer conv{}; + conv.configure(&src, &weights, nullptr, &dst, conv_info); + + src.allocator()->allocate(); + weights.allocator()->allocate(); + dst.allocator()->allocate(); + + library->fill_tensor_value(CLAccessor(src), 1.f); + library->fill_tensor_value(CLAccessor(weights), 1.f); + + conv.run(); + + // Compute reference to compare + SimpleTensor ref_src{ src_shape, dt }; + SimpleTensor ref_weights{ weights_shape, dt }; + SimpleTensor ref_bias{ bias_shape, dt }; + library->fill_tensor_value(ref_src, 1.f); + library->fill_tensor_value(ref_weights, 1.f); + // No bias + library->fill_tensor_value(ref_bias, 0.f); + auto ref_dst = reference::convolution_layer(ref_src, ref_weights, ref_bias, dst_shape, conv_info); + + validate(CLAccessor(dst), ref_dst); +} + // *INDENT-OFF* // clang-format off DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(zip(zip( diff --git a/tests/validation/NEON/DirectConvolutionLayer.cpp b/tests/validation/NEON/DirectConvolutionLayer.cpp index 6c47fa1cf8..ffffe7e3d5 100644 --- a/tests/validation/NEON/DirectConvolutionLayer.cpp +++ b/tests/validation/NEON/DirectConvolutionLayer.cpp @@ -129,6 +129,49 @@ const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo TEST_SUITE(NEON) TEST_SUITE(DirectConvolutionLayer) +/** Check whether the configuration of a Direct Convolution layer with no + * bias leads to a successful execution. + */ +TEST_CASE(NoBias, framework::DatasetMode::PRECOMMIT) +{ + const auto src_shape = TensorShape(27U, 13U, 2U); + const auto weights_shape = TensorShape(3U, 3U, 2U, 4U); + const auto bias_shape = TensorShape(4U); + const auto dst_shape = TensorShape(25U, 11U, 4U); + constexpr auto dt = DataType::F32; + + auto src = create_tensor(src_shape, dt); + auto weights = create_tensor(weights_shape, dt); + auto dst = create_tensor(dst_shape, dt); + + const auto conv_info = PadStrideInfo(1, 1, 0, 0); + + // Create Direct Convolution function + NEDirectConvolutionLayer conv{}; + conv.configure(&src, &weights, nullptr, &dst, conv_info); + + src.allocator()->allocate(); + weights.allocator()->allocate(); + dst.allocator()->allocate(); + + library->fill_tensor_value(Accessor(src), 1.f); + library->fill_tensor_value(Accessor(weights), 1.f); + + conv.run(); + + // Compute reference to compare + SimpleTensor ref_src{ src_shape, dt }; + SimpleTensor ref_weights{ weights_shape, dt }; + SimpleTensor ref_bias{ bias_shape, dt }; + library->fill_tensor_value(ref_src, 1.f); + library->fill_tensor_value(ref_weights, 1.f); + // No bias + library->fill_tensor_value(ref_bias, 0.f); + auto ref_dst = reference::convolution_layer(ref_src, ref_weights, ref_bias, dst_shape, conv_info); + + validate(Accessor(dst), ref_dst); +} + // *INDENT-OFF* // clang-format off DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(zip(zip( -- cgit v1.2.1