diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2020-10-26 15:04:08 +0000 |
|---|---|---|
| committer | Giorgio Arena <giorgio.arena@arm.com> | 2020-11-12 12:42:51 +0000 |
| commit | 2d1a835b68eb27a800838fc2b563b12eddf2c19f (patch) | |
| tree | 228dee073d37d2ec5b5dfbdb3d0e1e512ecb2d22 /src/core/CL/cl_kernels/pooling_layer.cl | |
| parent | 00c7601b1f9c3bec1d3b1db844abb513b9012541 (diff) | |
| download | ComputeLibrary-2d1a835b68eb27a800838fc2b563b12eddf2c19f.tar.gz | |
COMPMID-3735 Remove OpenCL padding: CLSoftmaxLayerKernel
- Renamed SELECT_DATA_TYPE to SELECT_VEC_DATA_TYPE to reflect its usage with vectors. SELECT_DATA_TYPE(dt) will now return the primitive data type
- Changed the interface of VEC_OFFS and V_OFFS in order to receive the primitive data type as a parameter rather than its vector form
- Performed a general cleanup of the kernels, such as creating macro for sum and max reduces, remove reduntant macros, defines, variables, calculations, etc...
- Using VEC_SIZE and VEC_SIZE_LEFTOVER in every kernel in order to allow computation for smaller shapes without adding paddings
- Removed the actual padding from the kernel and adjusting its calculations accordingly. Added asserts for padding removal checks. Removed invalid Validate tests.
Change-Id: If5ccbd5d34e255d38c7f6bfe8740e2b80b28e264
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4277
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: SiCong Li <sicong.li@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/pooling_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/pooling_layer.cl | 6 |
1 files changed, 3 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index 680e947149..00250a08a5 100644 --- a/src/core/CL/cl_kernels/pooling_layer.cl +++ b/src/core/CL/cl_kernels/pooling_layer.cl @@ -786,7 +786,7 @@ __kernel void pooling_layer_MxN_nhwc( } #endif // defined(POOL_SIZE_X) && defined(POOL_SIZE_Y) -#define SELECT_TYPE SELECT_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE) +#define SELECT_TYPE SELECT_VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE) /** Performs pooling layer of size equal to 2. This OpenCL kernel can perform the following pooling types: * -# max, -DPOOL_MAX must be passed at compile time @@ -957,7 +957,7 @@ __kernel void pooling_layer_2x2_nhwc( // note: Batch dimension does not contribute in the offset contribution VEC_DATA_TYPE(uint, VEC_SIZE) base_index = (uint)idx_out_c; - base_index += VEC_OFFS(VEC_DATA_TYPE(uint, VEC_SIZE), VEC_SIZE); + base_index += VEC_OFFS(uint, VEC_SIZE); VEC_DATA_TYPE(uint, VEC_SIZE) index0 = base_index + (uint)x0 * DST_CHANNELS + (uint)y0 * (DST_CHANNELS * SRC_WIDTH); VEC_DATA_TYPE(uint, VEC_SIZE) index1 = base_index + (uint)x1 * DST_CHANNELS + (uint)y0 * (DST_CHANNELS * SRC_WIDTH); @@ -978,4 +978,4 @@ __kernel void pooling_layer_2x2_nhwc( STORE_VECTOR_SELECT(index, uint, idx_base_ptr, VEC_SIZE, VEC_SIZE_LEFTOVER, ((VEC_SIZE_LEFTOVER != 0) && get_global_id(0) == 0)); #endif // defined(EXTRACT_MAX_INDEX) && defined(POOL_MAX) } -#endif // defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(DST_CHANNELS) && defined(DST_HEIGHT) && defined(DST_BATCH_SIZE) && defined(SELECT_DATA_TYPE) && defined(ACC_DATA_TYPE)
\ No newline at end of file +#endif // defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(DST_CHANNELS) && defined(DST_HEIGHT) && defined(DST_BATCH_SIZE) && defined(ACC_DATA_TYPE)
\ No newline at end of file |