From 0b76f7dd12240dc7a546c202ee80a7942d9898cd Mon Sep 17 00:00:00 2001 From: Gian Marco Iodice Date: Thu, 8 Apr 2021 17:20:00 +0100 Subject: Add support for cl_image in CLDirectConvolutionLayer - The cl_image object can be used for the weights - cl_image can only work for f32/f16 - Fix the implicit padding on the first dimension X Resolves COMPMID-4341 Change-Id: I04e0901c69e7765c42afceca38c4a840645b9123 Signed-off-by: Gian Marco Iodice Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5393 Reviewed-by: Giorgio Arena Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/winograd_output_transform.cl | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) (limited to 'src/core/CL/cl_kernels/winograd_output_transform.cl') diff --git a/src/core/CL/cl_kernels/winograd_output_transform.cl b/src/core/CL/cl_kernels/winograd_output_transform.cl index 674a138d48..6bd90604e5 100644 --- a/src/core/CL/cl_kernels/winograd_output_transform.cl +++ b/src/core/CL/cl_kernels/winograd_output_transform.cl @@ -637,7 +637,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( #if defined(HAS_BIAS) TILE(DATA_TYPE, 1, N0, b); - T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 0, b); + T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); // c = c + bias[broadcasted] T_ADD_BROADCAST_X(DATA_TYPE, 4, N0, out, b, out); @@ -718,7 +718,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( #if defined(HAS_BIAS) TILE(DATA_TYPE, 1, N0, b); - T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 0, b); + T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); // c = c + bias[broadcasted] T_ADD_BROADCAST_X(DATA_TYPE, 16, N0, out, b, out); @@ -1070,7 +1070,7 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc( #if defined(HAS_BIAS) TILE(DATA_TYPE, 1, N0, b); - T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 0, b); + T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); // c = c + bias[broadcasted] T_ADD_BROADCAST_X(DATA_TYPE, 4, N0, out, b, out); @@ -1162,7 +1162,7 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc( #if defined(HAS_BIAS) TILE(DATA_TYPE, 1, N0, b); - T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 0, b); + T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); // c = c + bias[broadcasted] T_ADD_BROADCAST_X(DATA_TYPE, 16, N0, out, b, out); -- cgit v1.2.1