diff options
author | Georgios Pinitas <georgios.pinitas@arm.com> | 2017-06-23 18:03:44 +0100 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-09-17 14:15:39 +0100 |
commit | e5f8fd64a46ce61dac61fd50095a27cfd94930b5 (patch) | |
tree | 2b43d580730cfa4ae018d06660403622d08bd00b /src/core/CL/cl_kernels/softmax_layer.cl | |
parent | 443c8b979518ad494b9e59648cbd061ebf37cba9 (diff) | |
download | ComputeLibrary-e5f8fd64a46ce61dac61fd50095a27cfd94930b5.tar.gz |
COMPMID-423: Port CLSoftmaxLayer to QS8
Change-Id: I759b7585656d018d7c864425118cd3ec2ca9b0eb
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/78908
Reviewed-by: Michele DiGiorgio <michele.digiorgio@arm.com>
Reviewed-by: Moritz Pflanzer <moritz.pflanzer@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/softmax_layer.cl')
-rw-r--r-- | src/core/CL/cl_kernels/softmax_layer.cl | 85 |
1 files changed, 55 insertions, 30 deletions
diff --git a/src/core/CL/cl_kernels/softmax_layer.cl b/src/core/CL/cl_kernels/softmax_layer.cl index 632b4a5374..a29aea4fae 100644 --- a/src/core/CL/cl_kernels/softmax_layer.cl +++ b/src/core/CL/cl_kernels/softmax_layer.cl @@ -23,14 +23,36 @@ */ #include "helpers.h" +#if defined(FIXED_POINT_POSITION) + +#include "fixed_point.h" +#define MAX_OP(x, y, type, size) MAX_OP_EXPAND(x, y, type, size) +#define ADD_OP(x, y, type, size) ADD_SAT_OP_EXPAND((x), (y), type, size) +#define SUB_OP(x, y, type, size) SUB_SAT_OP_EXPAND((x), (y), type, size) +#define DIV_OP(x, y, type, size) DIV_SAT_OP_EXPAND((x), (y), type, size, FIXED_POINT_POSITION) +#define EXP_OP(x, type, size) EXP_OP_EXPAND((x), type, size, FIXED_POINT_POSITION) + +#define MIN_VAL_EXPAND(type) type##_MIN +#define MIN_VAL(type) MIN_VAL_EXPAND(type) +#define MINVAL MIN_VAL(DATA_TYPE) +#define SELECT_DATA_TYPE EXPAND(DATA_TYPE) + +#else + +#define MAX_OP(x, y, type, size) max((x), (y)) +#define ADD_OP(x, y, type, size) ((x) + (y)) +#define SUB_OP(x, y, type, size) ((x) - (y)) +#define DIV_OP(x, y, type, size) ((x) / (y)) +#define EXP_OP(x, type, size) exp((x)) + #if defined USE_F16 -#define MINVAL HALF_MIN +#define MINVAL -HALF_MAX #define SELECT_DATA_TYPE short -#define DATA_TYPE half #else -#define MINVAL FLT_MIN +#define MINVAL -FLT_MAX #define SELECT_DATA_TYPE int -#define DATA_TYPE float +#endif + #endif __constant VEC_DATA_TYPE(DATA_TYPE, 16) type_min = (VEC_DATA_TYPE(DATA_TYPE, 16))(MINVAL); @@ -39,16 +61,16 @@ __constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, /** Identifies the maximum value across the 1st dimension. * * @note Datatype must be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short - * @note In case F16 is used -DUSE_HALF must be passed otherwise the kernel will default to used F32. + * @note Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 * @note In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed. * - * @param[in] src_ptr Pointer to the source tensor slice. Supported data types: F16, F32 + * @param[in] src_ptr Pointer to the source tensor slice. Supported data types: QS8/F16/F32 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor - * @param[out] dst_ptr Pointer to the destination tensor slice. Supported data types: F16, F32 + * @param[out] dst_ptr Pointer to the destination tensor slice. Supported data types: same as @p src_ptr * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) @@ -74,7 +96,7 @@ __kernel void softmax_layer_max( { VEC_DATA_TYPE(DATA_TYPE, 16) data = vload16(0, (__global DATA_TYPE *)offset(&src, i << 4, 0)); - max_val = max(data, max_val); + max_val = MAX_OP(data, max_val, DATA_TYPE, 16); } #if defined NON_MULTIPLE_OF_16 @@ -83,14 +105,14 @@ __kernel void softmax_layer_max( data = vload16(0, (__global DATA_TYPE *)offset(&src, width4 << 4, 0)); VEC_DATA_TYPE(SELECT_DATA_TYPE, 16) widx = CONVERT(((uint16)(width4 << 4) + idx16) < width, VEC_DATA_TYPE(SELECT_DATA_TYPE, 16)); - max_val = max(max_val, select(type_min, data, widx)); + max_val = MAX_OP(max_val, select(type_min, data, widx), DATA_TYPE, 16); #endif // Perform max reduction - max_val.s01234567 = max(max_val.s01234567, max_val.s89ABCDEF); - max_val.s0123 = max(max_val.s0123, max_val.s4567); - max_val.s01 = max(max_val.s01, max_val.s23); - max_val.s0 = max(max_val.s0, max_val.s1); + max_val.s01234567 = MAX_OP(max_val.s01234567, max_val.s89ABCDEF, DATA_TYPE, 8); + max_val.s0123 = MAX_OP(max_val.s0123, max_val.s4567, DATA_TYPE, 4); + max_val.s01 = MAX_OP(max_val.s01, max_val.s23, DATA_TYPE, 2); + max_val.s0 = MAX_OP(max_val.s0, max_val.s1, DATA_TYPE, 1); // Store result *((__global DATA_TYPE *)dst.ptr) = max_val.s0; @@ -100,28 +122,28 @@ __kernel void softmax_layer_max( * then gets the exponent of each element as sums all elements across each row. * * @note Datatype must be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short - * @note In case F16 is used -DUSE_HALF must be passed otherwise the kernel will default to used F32. + * @note Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 * @note In case the input is not multiple of 16 -DNON_MULTIPLE_OF_16 must be passed. * - * @param[in] src_ptr Pointer to the source tensor slice. Supported data types: F16, F32 + * @param[in] src_ptr Pointer to the source tensor slice. Supported data types: QS8/F16/F32 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor - * @param[in] max_ptr Pointer to the max values tensor slice. Supported data types: F16, F32 + * @param[in] max_ptr Pointer to the max values tensor slice. Supported data types: same as @p src_ptr * @param[in] max_stride_x Stride of the max values tensor in X dimension (in bytes) * @param[in] max_step_x max_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] max_stride_y Stride of the max values tensor in Y dimension (in bytes) * @param[in] max_step_y max_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] max_offset_first_element_in_bytes The offset of the first element in the max values tensor - * @param[out] dst_ptr Pointer to the destination tensor slice. Supported data types: F16, F32 + * @param[out] dst_ptr Pointer to the destination tensor slice. Supported data types: same as @p src_ptr * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] sum_ptr Pointer to the sum values tensor slice. Supported data types: F16, F32 + * @param[out] sum_ptr Pointer to the sum values tensor slice. Supported data types: same as @p src_ptr * @param[in] sum_stride_x Stride of the sum values tensor in X dimension (in bytes) * @param[in] sum_step_x sum_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] sum_stride_y Stride of the sum values tensor in Y dimension (in bytes) @@ -154,28 +176,30 @@ __kernel void softmax_layer_shift_exp_sum( { VEC_DATA_TYPE(DATA_TYPE, 16) data = vload16(0, (__global DATA_TYPE *)offset(&src, i << 4, 0)); - data = exp(data - max_val); + data = SUB_OP(data, max_val, DATA_TYPE, 16); + data = EXP_OP(data, DATA_TYPE, 16); vstore16(data, 0, (__global DATA_TYPE *)offset(&dst, i << 4, 0)); - sum1D += data; + sum1D = ADD_OP(sum1D, data, DATA_TYPE, 16); } #if defined NON_MULTIPLE_OF_16 // Handle non multiple of 16 VEC_DATA_TYPE(DATA_TYPE, 16) data = vload16(0, (__global DATA_TYPE *)offset(&src, width4 << 4, 0)); - data = exp(data - max_val); + data = SUB_OP(data, max_val, DATA_TYPE, 16); + data = EXP_OP(data, DATA_TYPE, 16); VEC_DATA_TYPE(SELECT_DATA_TYPE, 16) widx = CONVERT(((uint16)(width4 << 4) + idx16) < width, VEC_DATA_TYPE(SELECT_DATA_TYPE, 16)); data = select(0, data, widx); vstore16(data, 0, (__global DATA_TYPE *)offset(&dst, width4 << 4, 0)); - sum1D += data; + sum1D = ADD_OP(sum1D, data, DATA_TYPE, 16); #endif // Perform min/max reduction - sum1D.s01234567 = sum1D.s01234567 + sum1D.s89ABCDEF; - sum1D.s0123 = sum1D.s0123 + sum1D.s4567; - sum1D.s01 = sum1D.s01 + sum1D.s23; - sum1D.s0 = sum1D.s0 + sum1D.s1; + sum1D.s01234567 = ADD_OP(sum1D.s01234567, sum1D.s89ABCDEF, DATA_TYPE, 8); + sum1D.s0123 = ADD_OP(sum1D.s0123, sum1D.s4567, DATA_TYPE, 4); + sum1D.s01 = ADD_OP(sum1D.s01, sum1D.s23, DATA_TYPE, 2); + sum1D.s0 = ADD_OP(sum1D.s0, sum1D.s1, DATA_TYPE, 1); // Calculate and store result *((__global DATA_TYPE *)sum.ptr) = sum1D.s0; @@ -184,20 +208,21 @@ __kernel void softmax_layer_shift_exp_sum( /** Divides all the values of the input tensor by the sum calculated from softmax_layer_shift_exp_sum kernel. * * @note Datatype must be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short + * @note Fixed point position must be given as a preprocessor argument using -DFIXED_POINT_POSITION=pos. e.g. DFIXED_POINT_POSITION=4 * - * @param[in] src_ptr Pointer to the source tensor slice. Supported data types: F16, F32 + * @param[in] src_ptr Pointer to the source tensor slice. Supported data types: QS8/F16/F32 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor - * @param[in] sum_ptr Pointer to the sum values tensor slice. Supported data types: F16, F32 + * @param[in] sum_ptr Pointer to the sum values tensor slice. Supported data types: same as @p src_ptr * @param[in] sum_stride_x Stride of the sum values tensor in X dimension (in bytes) * @param[in] sum_step_x sum_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] sum_stride_y Stride of the sum values tensor in Y dimension (in bytes) * @param[in] sum_step_y sum_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] sum_offset_first_element_in_bytes The offset of the first element in the sum values tensor - * @param[out] dst_ptr Pointer to the destination tensor slice. Supported data types: F16, F32 + * @param[out] dst_ptr Pointer to the destination tensor slice. Supported data types: same as @p src_ptr * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) @@ -217,5 +242,5 @@ __kernel void softmax_layer_norm( DATA_TYPE sum_val = *((__global DATA_TYPE *)offset(&sum, 0, get_global_id(1))); VEC_DATA_TYPE(DATA_TYPE, 16) data = vload16(0, (__global DATA_TYPE *)offset(&src, 0, 0)); - vstore16(data / sum_val, 0, (__global DATA_TYPE *)offset(&dst, 0, 0)); + vstore16(DIV_OP(data, sum_val, DATA_TYPE, 16), 0, (__global DATA_TYPE *)offset(&dst, 0, 0)); } |