From 868e54102b1761577e7783c80b66e82896ee063c Mon Sep 17 00:00:00 2001 From: steniu01 Date: Mon, 17 Jul 2017 23:16:00 +0100 Subject: COMPMID-459 Collapse CL Im2col's higher dimensions Change-Id: I0ccc39cbcf6926e6810faf3fe264c4af7adc3f7b Reviewed-on: http://mpd-gerrit.cambridge.arm.com/83070 Tested-by: Kaizen Reviewed-by: Anthony Barbier --- src/core/CL/cl_kernels/helpers.h | 64 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 64 insertions(+) (limited to 'src/core/CL/cl_kernels/helpers.h') diff --git a/src/core/CL/cl_kernels/helpers.h b/src/core/CL/cl_kernels/helpers.h index 59b81d7f06..68af64e344 100644 --- a/src/core/CL/cl_kernels/helpers.h +++ b/src/core/CL/cl_kernels/helpers.h @@ -72,6 +72,18 @@ uint name##_step_z, \ uint name##_offset_first_element_in_bytes +#define TENSOR4D_DECLARATION(name) \ + __global uchar *name##_ptr, \ + uint name##_stride_x, \ + uint name##_step_x, \ + uint name##_stride_y, \ + uint name##_step_y, \ + uint name##_stride_z, \ + uint name##_step_z, \ + uint name##_stride_w, \ + uint name##_step_w, \ + uint name##_offset_first_element_in_bytes + #define CONVERT_TO_VECTOR_STRUCT(name) \ update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x) @@ -84,6 +96,9 @@ #define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name) \ update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0) +#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \ + update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z) + #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name) \ update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, name##_step_z) @@ -97,6 +112,13 @@ #define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name) \ update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0) +#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size) \ + update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ + name##_stride_z, name##_step_z, name##_stride_w, name##_step_z, mod_size) + +#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size) \ + update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0, name##_stride_w, 0, mod_size) + /** Structure to hold Vector information */ typedef struct Vector { @@ -124,6 +146,17 @@ typedef struct Tensor3D int stride_z; /**< Stride of the image in Z dimension (in bytes) */ } Tensor3D; +/** Structure to hold 4D tensor information */ +typedef struct Tensor4D +{ + __global uchar *ptr; /**< Pointer to the starting postion of the buffer */ + int offset_first_element_in_bytes; /**< The offset of the first element in the source image */ + int stride_x; /**< Stride of the image in X dimension (in bytes) */ + int stride_y; /**< Stride of the image in Y dimension (in bytes) */ + int stride_z; /**< Stride of the image in Z dimension (in bytes) */ + int stride_w; /**< Stride of the image in W dimension (in bytes) */ +} Tensor4D; + /** Wrap vector information into an Vector structure, and make the pointer point at this workitem's data. * * @param[in] ptr Pointer to the starting postion of the buffer @@ -222,6 +255,24 @@ Tensor3D inline update_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_fi return tensor; } +Tensor4D inline update_tensor4D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z, uint stride_w, + uint step_w, + uint mod_size) +{ + Tensor4D tensor = + { + .ptr = ptr, + .offset_first_element_in_bytes = offset_first_element_in_bytes, + .stride_x = stride_x, + .stride_y = stride_y, + .stride_z = stride_z, + .stride_w = stride_w + }; + + tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + (get_global_id(2) % mod_size) * step_z + (get_global_id(2) / mod_size) * step_w; + return tensor; +} + /** Get the pointer position of a Vector * * @param[in] vec Pointer to the starting position of the buffer @@ -255,4 +306,17 @@ __global inline const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z; } +/** Get the pointer position of a Tensor4D + * + * @param[in] tensor Pointer to the starting position of the buffer + * @param[in] x Relative X position + * @param[in] y Relative Y position + * @param[in] z Relative Z position + * @param[in] w Relative W position + */ +__global inline const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w) +{ + return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + w * tensor->stride_w; +} + #endif // _HELPER_H -- cgit v1.2.1