diff options
| author | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2021-04-08 17:20:00 +0100 |
|---|---|---|
| committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2021-04-12 17:39:32 +0000 |
| commit | 0b76f7dd12240dc7a546c202ee80a7942d9898cd (patch) | |
| tree | 7dbd9ae56483e111952a0cab4f19d2c3f25157e7 /src/core/CL/cl_kernels/winograd_output_transform.cl | |
| parent | 6dbcc0e4d2fd0c61602a1a0c4a0ac548da713087 (diff) | |
| download | ComputeLibrary-0b76f7dd12240dc7a546c202ee80a7942d9898cd.tar.gz | |
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 <gianmarco.iodice@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5393
Reviewed-by: Giorgio Arena <giorgio.arena@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/winograd_output_transform.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/winograd_output_transform.cl | 8 |
1 files changed, 4 insertions, 4 deletions
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); |