diff options
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)); } |