{"pooling_layer_MxN_quantized_nchw", "pooling_layer_quantized.cl"},
{"quantization_layer", "quantization_layer.cl"},
{"reduce_min_max", "reduce_operation.cl"},
- {"reduce_mean", "reduce_operation.cl"},
- {"reduce_sum", "reduce_operation.cl"},
+ {"reduce_sum_mean", "reduce_operation.cl"},
{"remap_nearest_neighbour", "remap.cl"},
{"remap_bilinear", "remap.cl"},
{"reshape_layer", "reshape_layer.cl"},
*((__global DATA_TYPE *)out.ptr) = value;
}
-/** Perform reduce mean
+/** Perform reduce sum/mean
*
* @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
* @attention Output tensor depth should be given as a preprocessor argument using -DDEPTH_OUT=size. e.g. -DDEPTH_OUT=16
* @param[in] axis Axis through which reduction occurs
* @param[in] dim Dimension across the axis to be reduced.
*/
-__kernel void reduce_mean(TENSOR4D_DECLARATION(input),
- TENSOR4D_DECLARATION(output),
- const int axis,
- const int dim)
-{
- Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT(input, 0);
- Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT(output, DEPTH_OUT);
-
- int indices[4] =
- {
- get_global_id(0),
- get_global_id(1),
- get_global_id(2) % DEPTH_OUT,
- get_global_id(2) / DEPTH_OUT,
- };
-
- DATA_TYPE sum_value = (DATA_TYPE)0;
- for(int i = 0; i < dim; ++i)
- {
- indices[axis] = i;
- sum_value += *((__global DATA_TYPE *)tensor4D_offset(&in, indices[0], indices[1], indices[2], indices[3]));
- }
-
- *((__global DATA_TYPE *)out.ptr) = sum_value / CONVERT(dim, DATA_TYPE);
-}
-
-/** Perform reduce sum
- *
- * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
- * @attention Output tensor depth should be given as a preprocessor argument using -DDEPTH_OUT=size. e.g. -DDEPTH_OUT=16
- *
- * @param[in] input_ptr Pointer to the source image. Supported data types: U8/S8/U16/S16/F16/U32/S32/F32
- * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
- * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
- * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
- * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
- * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
- * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
- * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image
- * @param[in] input_stride_w Stride of the source tensor in W dimension (in bytes)
- * @param[in] input_step_w output_stride_w * number of elements along W processed per workitem(in bytes)
- * @param[out] output_ptr Pointer to the destination image. Supported data types: same as @p input_ptr
- * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
- * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
- * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
- * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
- * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
- * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
- * @param[in] output_stride_w Stride of the source tensor in W dimension (in bytes)
- * @param[in] output_step_w output_stride_w * number of elements along W processed per workitem(in bytes)
- * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
- * @param[in] axis Axis through which reduction occurs
- * @param[in] dim Dimension across the axis to be reduced.
- */
-__kernel void reduce_sum(TENSOR4D_DECLARATION(input),
- TENSOR4D_DECLARATION(output),
- const int axis,
- const int dim)
+__kernel void reduce_sum_mean(TENSOR4D_DECLARATION(input),
+ TENSOR4D_DECLARATION(output),
+ const int axis,
+ const int dim)
{
Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT(input, 0);
Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT(output, DEPTH_OUT);
sum_value += *((__global DATA_TYPE *)tensor4D_offset(&in, indices[0], indices[1], indices[2], indices[3]));
}
+#if OP_CODE == 2 // REDUCE_SUM
*((__global DATA_TYPE *)out.ptr) = sum_value;
+#elif OP_CODE == 1 // REDUCE_MEAN
+ *((__global DATA_TYPE *)out.ptr) = sum_value / CONVERT(dim, DATA_TYPE);
+#endif
}
#endif // defined(DATA_TYPE) && defined(DEPTH_OUT) && defined(OP_CODE)