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 <michele.digiorgio@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5322
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: James Conroy <james.conroy@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index 726efa3..2652884 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -901,12 +901,13 @@
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 1de3737..dde024f 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 d60d11a..74867ff 100644
--- a/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp
+++ b/src/runtime/CL/functions/CLDirectConvolutionLayer.cpp
@@ -58,7 +58,7 @@
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 @@
_impl->dst = output;
_impl->op = std::make_unique<opencl::ClDirectConvolution>();
- _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 4671d8c..c012340 100644
--- a/tests/validation/CL/DirectConvolutionLayer.cpp
+++ b/tests/validation/CL/DirectConvolutionLayer.cpp
@@ -87,6 +87,49 @@
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<CLTensor>(src_shape, dt);
+ auto weights = create_tensor<CLTensor>(weights_shape, dt);
+ auto dst = create_tensor<CLTensor>(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<float> ref_src{ src_shape, dt };
+ SimpleTensor<float> ref_weights{ weights_shape, dt };
+ SimpleTensor<float> 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<float>(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 6c47fa1..ffffe7e 100644
--- a/tests/validation/NEON/DirectConvolutionLayer.cpp
+++ b/tests/validation/NEON/DirectConvolutionLayer.cpp
@@ -129,6 +129,49 @@
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<Tensor>(src_shape, dt);
+ auto weights = create_tensor<Tensor>(weights_shape, dt);
+ auto dst = create_tensor<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<float> ref_src{ src_shape, dt };
+ SimpleTensor<float> ref_weights{ weights_shape, dt };
+ SimpleTensor<float> 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<float>(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(