diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2021-04-16 17:03:39 +0100 |
|---|---|---|
| committer | Giorgio Arena <giorgio.arena@arm.com> | 2021-04-20 09:26:59 +0000 |
| commit | ada6cbc057ff725e57d301a99a1816ce602485b9 (patch) | |
| tree | f869994cb2b061de0bc4731d720336413b81d32a /src/core/CL/cl_kernels/tile_helpers.h | |
| parent | 031d6a97de79fc3ca3eb6fca1611f03aa9b5893b (diff) | |
| download | ComputeLibrary-ada6cbc057ff725e57d301a99a1816ce602485b9.tar.gz | |
Remove OpenCL padding: CLPixelWiseMultiplicationKernel
- Change kernel's vec_size to 16 / sizeof(output)
- Change ICLKernel.cpp to handle broadcast without padding
Resolve COMPMID-3913
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Change-Id: I03e884b250ef5784dc109bff8cf2c96b345d119f
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5450
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/tile_helpers.h')
| -rw-r--r-- | src/core/CL/cl_kernels/tile_helpers.h | 14 |
1 files changed, 1 insertions, 13 deletions
diff --git a/src/core/CL/cl_kernels/tile_helpers.h b/src/core/CL/cl_kernels/tile_helpers.h index 496f2dd664..8b6d5309e3 100644 --- a/src/core/CL/cl_kernels/tile_helpers.h +++ b/src/core/CL/cl_kernels/tile_helpers.h @@ -83,18 +83,6 @@ */ #define GET_SPATIAL_IDX(IDX, N0, PARTIAL_N0) (max((int)(get_global_id(IDX) * N0 - (N0 - PARTIAL_N0) % N0), 0)) -/** Offset (in bytes) calculation for a 1D BUFFER (cl_buffer) tensor */ -#define OFFSET1D(base, data_type, x) (base##_offset_first_element_in_bytes + x * sizeof(data_type)) - -/** Offset (in bytes) calculation for a 2D BUFFER (cl_buffer) tensor */ -#define OFFSET2D(base, data_type, x, y) (base##_offset_first_element_in_bytes + x * sizeof(data_type) + y * base##_stride_y) - -/** Offset (in bytes) calculation for a 3D BUFFER (cl_buffer) tensor */ -#define OFFSET3D(base, data_type, x, y, z) (base##_offset_first_element_in_bytes + x * sizeof(data_type) + y * base##_stride_y + z * base##_stride_z) - -/** Offset (in bytes) calculation for a 4D BUFFER (cl_buffer) tensor */ -#define OFFSET4D(base, data_type, x, y, z, w) (base##_offset_first_element_in_bytes + x * sizeof(data_type) + y * base##_stride_y + z * base##_stride_z + w * base##_stride_w) - /** Dot product integet 8bit function * * @note Performs: c += dot(a, b) @@ -184,7 +172,7 @@ LOOP_UNROLLING(int, _i, 0, HEIGHT, 1) \ { \ dst[_i].v = V_LOAD(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, ((Y) + _i * (int)(YI_MULTIPLIER)), STRIDE_Y); \ - } \ + } \ }) /** Load a tile from global memory (tensor) using an indirect Y index tile |