diff options
| author | Anthony Barbier <anthony.barbier@arm.com> | 2017-07-11 18:36:40 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-09-17 14:15:39 +0100 |
| commit | 9a7182e8b53ab77d26b56752cc95b80f4e1774f2 (patch) | |
| tree | f7798baaf2faa7f5b69678115a6d89b5f8499350 /src/core/CL/cl_kernels/pixelwise_mul_float.cl | |
| parent | be35b0ed7cb6fbf9010142aab97854c467b14c5c (diff) | |
| download | ComputeLibrary-9a7182e8b53ab77d26b56752cc95b80f4e1774f2.tar.gz | |
COMPMID-443 Use 3D tensor for pixel multiply (Needed for Normalization Layer)
Change-Id: I117688f12334e6afc705c863acdf71b0bb1fc6e8
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/80352
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/pixelwise_mul_float.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/pixelwise_mul_float.cl | 18 |
1 files changed, 12 insertions, 6 deletions
diff --git a/src/core/CL/cl_kernels/pixelwise_mul_float.cl b/src/core/CL/cl_kernels/pixelwise_mul_float.cl index 89367dc0ce..98127e0311 100644 --- a/src/core/CL/cl_kernels/pixelwise_mul_float.cl +++ b/src/core/CL/cl_kernels/pixelwise_mul_float.cl @@ -43,31 +43,37 @@ * @param[in] in1_step_x in1_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] in1_stride_y Stride of the source image in Y dimension (in bytes) * @param[in] in1_step_y in1_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] in1_stride_z Stride of the source image in Y dimension (in bytes) + * @param[in] in1_step_z in1_stride_z * number of elements along Y processed per workitem(in bytes) * @param[in] in1_offset_first_element_in_bytes The offset of the first element in the source image * @param[in] in2_ptr Pointer to the source image. Supported data types: U8, S16, F16, F32 * @param[in] in2_stride_x Stride of the source image in X dimension (in bytes) * @param[in] in2_step_x in2_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] in2_stride_y Stride of the source image in Y dimension (in bytes) * @param[in] in2_step_y in2_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] in2_stride_z Stride of the source image in Y dimension (in bytes) + * @param[in] in2_step_z in2_stride_z * number of elements along Y processed per workitem(in bytes) * @param[in] in2_offset_first_element_in_bytes The offset of the first element in the source image * @param[out] out_ptr Pointer to the destination image. Supported data types: U8, S16, F16, F32 * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) * @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes) * @param[in] out_step_y out_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] out_stride_z Stride of the destination image in Y dimension (in bytes) + * @param[in] out_step_z out_stride_z * number of elements along Y processed per workitem(in bytes) * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination image * @param[in] scale Float scaling factor. Supported data types: F32 */ __kernel void pixelwise_mul_float( - IMAGE_DECLARATION(in1), - IMAGE_DECLARATION(in2), - IMAGE_DECLARATION(out), + TENSOR3D_DECLARATION(in1), + TENSOR3D_DECLARATION(in2), + TENSOR3D_DECLARATION(out), const float scale) { // Get pixels pointer - Image in1 = CONVERT_TO_IMAGE_STRUCT(in1); - Image in2 = CONVERT_TO_IMAGE_STRUCT(in2); - Image out = CONVERT_TO_IMAGE_STRUCT(out); + Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1); + Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2); + Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out); // Load data VEC_DATA_TYPE(DATA_TYPE_RES, 16) |