commit 97e258000f99c6a5f872f4e8968eaf6a93de2cf7
author Michele Di Giorgio <michele.digiorgio@arm.com> Thu Mar 25 12:37:45 2021 +0000
committer Georgios Pinitas <georgios.pinitas@arm.com> Fri Mar 26 09:23:49 2021 +0000
tree d4bd44c1350823ec9b12918e45180b3c7e7fe66e
parent 40efd5365108b97fc8d6cb93fa1d572a08a93ad5

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>

src/core/CL/CLKernelLibrary.cpp

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()

src/core/CL/cl_kernels/direct_convolution.cl

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)
 )

src/runtime/CL/functions/CLDirectConvolutionLayer.cpp

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,