diff options
author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2021-03-25 12:37:45 +0000 |
---|---|---|
committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2021-03-26 09:23:49 +0000 |
commit | 97e258000f99c6a5f872f4e8968eaf6a93de2cf7 (patch) | |
tree | d4bd44c1350823ec9b12918e45180b3c7e7fe66e /src | |
parent | 40efd5365108b97fc8d6cb93fa1d572a08a93ad5 (diff) | |
download | ComputeLibrary-97e258000f99c6a5f872f4e8968eaf6a93de2cf7.tar.gz |
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>
Diffstat (limited to 'src')
-rw-r--r-- | src/core/CL/CLKernelLibrary.cpp | 3 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/direct_convolution.cl | 3 | ||||
-rw-r--r-- | src/runtime/CL/functions/CLDirectConvolutionLayer.cpp | 4 |
3 files changed, 6 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<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, |