diff options
author | Michalis Spyrou <michalis.spyrou@arm.com> | 2018-09-11 11:16:47 +0100 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:54:54 +0000 |
commit | f1addb665ad668dcd34e18c52e4961a7cf5e3886 (patch) | |
tree | 11c1b606ab26416dd0b39f1526e7eba60394e7c3 /src/core/CL/cl_kernels/batch_to_space.cl | |
parent | 96f6769de9d9ebe3e631b81aac9a82934a79b0c4 (diff) | |
download | ComputeLibrary-f1addb665ad668dcd34e18c52e4961a7cf5e3886.tar.gz |
COMPMID-1549 Implementing Batch to Space on OpenCL - NHWC
Change-Id: If7ae0a8b6255a10711365068d9fb153c71f09818
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/147751
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/batch_to_space.cl')
-rw-r--r-- | src/core/CL/cl_kernels/batch_to_space.cl | 109 |
1 files changed, 105 insertions, 4 deletions
diff --git a/src/core/CL/cl_kernels/batch_to_space.cl b/src/core/CL/cl_kernels/batch_to_space.cl index 3043c2cf17..8506fc3709 100644 --- a/src/core/CL/cl_kernels/batch_to_space.cl +++ b/src/core/CL/cl_kernels/batch_to_space.cl @@ -24,7 +24,7 @@ #include "helpers.h" #if defined(DATA_TYPE) && defined(BATCH_SIZE) -/** Batch to space transformation. +/** Batch to space transformation. (NCHW) * * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float @@ -54,7 +54,7 @@ * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor */ -__kernel void batch_to_space( +__kernel void batch_to_space_nchw( TENSOR3D_DECLARATION(input), const int batch_id, VECTOR_DECLARATION(block_shape), @@ -78,10 +78,64 @@ __kernel void batch_to_space( *((__global DATA_TYPE *)tensor4D_offset(&out, out_x, out_y, z, w)) = *((__global DATA_TYPE *)in.ptr); } +/** Batch to space transformation. (NHWC) + * + * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float + * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float + * @note The input tensor batch size must be passed at compile time using -DBATCH_SIZE. e.g. -DBATCH_SIZE=2 + * + * @param[in] input_ptr Pointer to the source tensor. Supported data types: U8/S8/QASYMM8/U16/S16/F16/U32/S32/F32 + * @param[in] input_stride_x Stride of the source tensor 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 tensor 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 first source tensor + * @param[in] batch_id The input tensor batch id + * @param[in] block_shape_ptr Pointer to the source tensor. Supported data types: S32 + * @param[in] block_shape_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] block_shape_step_x block_shape_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] block_shape_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] block_shape_step_y block_shape_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_offset_first_element_in_bytes The offset of the first element in the first source tensor + * @param[out] output_ptr Pointer to the destination tensor. Supported data types: same as @p input_ptr + * @param[in] output_stride_x Stride of the destination tensor 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 tensor 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_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void batch_to_space_nhwc( + TENSOR3D_DECLARATION(input), + const int batch_id, + VECTOR_DECLARATION(block_shape), + TENSOR4D_DECLARATION(output)) +{ + Tensor3D in = CONVERT_TO_TENSOR3D_STRUCT(input); + Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(output, 0); + Vector block = CONVERT_TO_VECTOR_STRUCT_NO_STEP(block_shape); + + const int block_x = *((__global int *)vector_offset(&block, 0)); + const int block_y = *((__global int *)vector_offset(&block, 1)); + + const int r = (BATCH_SIZE / (block_x * block_y)); + const int x = get_global_id(1); + const int y = get_global_id(2); + const int z = get_global_id(0); + const int w = batch_id % r; + + const int out_x = x * block_x + (batch_id / r) % block_x; + const int out_y = y * block_y + (batch_id / r) / block_x; + + *((__global DATA_TYPE *)tensor4D_offset(&out, z, out_x, out_y, w)) = *((__global DATA_TYPE *)in.ptr); +} #endif // defined(DATA_TYPE) && defined(BATCH_SIZE) #if defined(DATA_TYPE) && defined(BATCH_SIZE) && defined(BLOCK_SHAPE_X) && defined(BLOCK_SHAPE_Y) -/** Batch to space transformation. +/** Batch to space transformation. (NCHW) * * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float * @note The input tensor batch size must be passed at compile time using -DBATCH_SIZE. e.g. -DBATCH_SIZE=2 @@ -106,7 +160,7 @@ __kernel void batch_to_space( * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor */ -__kernel void batch_to_space_static( +__kernel void batch_to_space_static_nchw( TENSOR3D_DECLARATION(input), const int batch_id, TENSOR4D_DECLARATION(output)) @@ -128,4 +182,51 @@ __kernel void batch_to_space_static( *((__global DATA_TYPE *)tensor4D_offset(&out, out_x, out_y, z, w)) = *((__global DATA_TYPE *)in.ptr); } +/** Batch to space transformation. (NHWC) + * + * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float + * @note The input tensor batch size must be passed at compile time using -DBATCH_SIZE. e.g. -DBATCH_SIZE=2 + * @note The block shape x must be passed at compile time using -DBLOCK_SHAPE_X. e.g. -DBLOCK_SHAPE_X=2 + * @note The block shape y must be passed at compile time using -DBLOCK_SHAPE_Y. e.g. -DBLOCK_SHAPE_Y=2 + * + * @param[in] input_ptr Pointer to the source tensor. Supported data types: U8/S8/QASYMM8/U16/S16/F16/U32/S32/F32 + * @param[in] input_stride_x Stride of the source tensor 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 tensor 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 first source tensor + * @param[in] batch_id The input tensor batch id + * @param[out] output_ptr Pointer to the destination tensor. Supported data types: same as @p input_ptr + * @param[in] output_stride_x Stride of the destination tensor 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 tensor 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_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void batch_to_space_static_nhwc( + TENSOR3D_DECLARATION(input), + const int batch_id, + TENSOR4D_DECLARATION(output)) +{ + Tensor3D in = CONVERT_TO_TENSOR3D_STRUCT(input); + Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(output, 0); + + const int block_x = BLOCK_SHAPE_X; + const int block_y = BLOCK_SHAPE_Y; + + const int r = (BATCH_SIZE / (block_x * block_y)); + const int x = get_global_id(1); + const int y = get_global_id(2); + const int z = get_global_id(0); + const int w = batch_id % r; + + const int out_x = x * block_x + (batch_id / r) % block_x; + const int out_y = y * block_y + (batch_id / r) / block_x; + + *((__global DATA_TYPE *)tensor4D_offset(&out, z, out_x, out_y, w)) = *((__global DATA_TYPE *)in.ptr); +} #endif // defined(DATA_TYPE) && defined(BATCH_SIZE) && defined(BLOCK_SHAPE_X) && defined(BLOCK_SHAPE_Y)
\ No newline at end of file |