From 8be9148814b88e5b0cabd5a4d2b1f4ff470a8c1c Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Tue, 26 Mar 2019 17:23:28 +0000 Subject: COMPMID-1959: Implements 2D FFT on OpenCL Change-Id: I73cf3984a5463acc854c8a59dc2bd9a5234cd99c Signed-off-by: Georgios Pinitas Reviewed-on: https://review.mlplatform.org/c/936 Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins Reviewed-by: Gian Marco Iodice --- src/core/CL/cl_kernels/fft.cl | 1111 +++++++++++++++++++++---- src/core/CL/cl_kernels/fft_digit_reverse.cl | 148 ++++ src/core/CL/cl_kernels/fft_scale.cl | 78 ++ src/core/CL/cl_kernels/pixelwise_mul_float.cl | 52 +- src/core/CL/cl_kernels/reduction_operation.cl | 17 +- 5 files changed, 1226 insertions(+), 180 deletions(-) create mode 100644 src/core/CL/cl_kernels/fft_digit_reverse.cl create mode 100644 src/core/CL/cl_kernels/fft_scale.cl (limited to 'src/core/CL/cl_kernels') diff --git a/src/core/CL/cl_kernels/fft.cl b/src/core/CL/cl_kernels/fft.cl index 5f1ef2483b..0027fd5b66 100644 --- a/src/core/CL/cl_kernels/fft.cl +++ b/src/core/CL/cl_kernels/fft.cl @@ -23,48 +23,6 @@ */ #include "helpers.h" -/** Computes the digit reverse stage - * - * @param[in] src_ptr Pointer to the source tensor. Supported data types: 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_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z 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. 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_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) - * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[in] idx_ptr Pointer to the index tensor. Supported data types: U32 - * @param[in] idx_stride_x Stride of the index tensor in X dimension (in bytes) - * @param[in] idx_step_x idx_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] idx_offset_first_element_in_bytes The offset of the first element in the index tensor - */ -__kernel void digit_reverse( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), - VECTOR_DECLARATION(idx)) -{ - // Get tensor pointers - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(src); - Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); - Vector idx = CONVERT_TO_VECTOR_STRUCT(idx); - - const unsigned int iidx = *((__global uint *)(idx.ptr)); - - // Load data - float2 data = vload2(0, (__global float *)tensor3D_offset(&src, iidx, get_global_id(1), get_global_id(2))); - - // Store result - vstore2(data, 0, (__global float *)dst.ptr); -} - /** Calculates and applies the twiddle factor to a given input. * * @param[in] phi The angle. @@ -252,7 +210,7 @@ __kernel void digit_reverse( c7 = s4 + t1; \ } -/** Computes the first stage of a radix-2 DFT. +/** Computes the first stage of a radix-2 DFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -264,14 +222,14 @@ __kernel void digit_reverse( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image */ kernel void fft_radix_2_first_stage_axis_0( TENSOR3D_DECLARATION(input) @@ -289,17 +247,66 @@ kernel void fft_radix_2_first_stage_axis_0( Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); #endif /* IN_PLACE */ - // Load eight complex input values + // Load two complex input values float4 data = vload4(0, (__global float *)input.ptr); // Compute DFT N = 2 DFT_2(data.s01, data.s23); - // Store eight complex output values + // Store two complex output values vstore4(data, 0, (__global float *)output.ptr); } -/** Computes the first stage of a radix-3 DFT. +/** Computes the first stage of a radix-2 DFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + */ +kernel void fft_radix_2_first_stage_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ +) +{ + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); +#endif /* IN_PLACE */ + + // Load two complex input values + float2 data1 = vload2(0, (__global float *)input.ptr); + float2 data2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 1, 0)); + + // Compute DFT N = 2 + DFT_2(data1, data2); + + // Store two complex output values + vstore2(data1, 0, (__global float *)output.ptr); + vstore2(data2, 0, (__global float *)tensor3D_offset(&output, 0, 1, 0)); +} + +/** Computes the first stage of a radix-3 DFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -311,14 +318,14 @@ kernel void fft_radix_2_first_stage_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image */ kernel void fft_radix_3_first_stage_axis_0( TENSOR3D_DECLARATION(input) @@ -336,19 +343,70 @@ kernel void fft_radix_3_first_stage_axis_0( Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); #endif /* IN_PLACE */ - // Load eight complex input values + // Load three complex input values float4 data0 = vload4(0, (__global float *)input.ptr); float2 data1 = vload2(0, (__global float *)tensor3D_offset(&input, 2, 0, 0)); // Compute DFT N = 3 DFT_3(data0.s01, data0.s23, data1.s01); - // Store eight complex output values + // Store three complex output values vstore4(data0, 0, (__global float *)output.ptr); vstore2(data1, 0, (__global float *)tensor3D_offset(&output, 2, 0, 0)); } -/** Computes the first stage of a radix-4 DFT. +/** Computes the first stage of a radix-3 DFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + */ +kernel void fft_radix_3_first_stage_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ +) +{ + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); +#endif /* IN_PLACE */ + + // Load three complex input values + float2 data0 = vload2(0, (__global float *)input.ptr); + float2 data1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 1, 0)); + float2 data2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2, 0)); + + // Compute DFT N = 3 + DFT_3(data0, data1, data2); + + // Store three complex output values + vstore2(data0, 0, (__global float *)output.ptr); + vstore2(data1, 0, (__global float *)tensor3D_offset(&output, 0, 1, 0)); + vstore2(data2, 0, (__global float *)tensor3D_offset(&output, 0, 2, 0)); +} + +/** Computes the first stage of a radix-4 DFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -360,14 +418,14 @@ kernel void fft_radix_3_first_stage_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image */ kernel void fft_radix_4_first_stage_axis_0( TENSOR3D_DECLARATION(input) @@ -385,17 +443,70 @@ kernel void fft_radix_4_first_stage_axis_0( Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); #endif /* IN_PLACE */ - // Load eight complex input values + // Load four complex input values float8 data = vload8(0, (__global float *)input.ptr); // Compute DFT N = 4 DFT_4(data.s01, data.s23, data.s45, data.s67); - // Store eight complex output values + // Store four complex output values vstore8(data, 0, (__global float *)output.ptr); } -/** Computes the first stage of a radix-5 DFT. +/** Computes the first stage of a radix-4 DFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + */ +kernel void fft_radix_4_first_stage_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ +) +{ + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); +#endif /* IN_PLACE */ + + // Load four complex input values + float2 data0 = vload2(0, (__global float *)input.ptr); + float2 data1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 1, 0)); + float2 data2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2, 0)); + float2 data3 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 3, 0)); + + // Compute DFT N = 4 + DFT_4(data0, data1, data2, data3); + + // Store four complex output values + vstore2(data0, 0, (__global float *)output.ptr); + vstore2(data1, 0, (__global float *)tensor3D_offset(&output, 0, 1, 0)); + vstore2(data2, 0, (__global float *)tensor3D_offset(&output, 0, 2, 0)); + vstore2(data3, 0, (__global float *)tensor3D_offset(&output, 0, 3, 0)); +} + +/** Computes the first stage of a radix-5 DFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -407,14 +518,14 @@ kernel void fft_radix_4_first_stage_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image */ kernel void fft_radix_5_first_stage_axis_0( TENSOR3D_DECLARATION(input) @@ -432,19 +543,74 @@ kernel void fft_radix_5_first_stage_axis_0( Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); #endif /* IN_PLACE */ - // Load eight complex input values + // Load five complex input values float8 data0 = vload8(0, (__global float *)input.ptr); float2 data1 = vload2(0, (__global float *)tensor3D_offset(&input, 4, 0, 0)); // Compute DFT N = 5 DFT_5(data0.s01, data0.s23, data0.s45, data0.s67, data1.s01); - // Store eight complex output values + // Store five complex output values vstore8(data0, 0, (__global float *)output.ptr); vstore2(data1, 0, (__global float *)tensor3D_offset(&output, 4, 0, 0)); } -/** Computes the first stage of a radix-7 DFT. +/** Computes the first stage of a radix-5 DFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + */ +kernel void fft_radix_5_first_stage_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ +) +{ + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); +#endif /* IN_PLACE */ + + // Load five complex input values + float2 data0 = vload2(0, (__global float *)input.ptr); + float2 data1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 1, 0)); + float2 data2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2, 0)); + float2 data3 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 3, 0)); + float2 data4 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 4, 0)); + + // Compute DFT N = 5 + DFT_5(data0, data1, data2, data3, data4); + + // Store five complex output values + vstore2(data0, 0, (__global float *)output.ptr); + vstore2(data1, 0, (__global float *)tensor3D_offset(&output, 0, 1, 0)); + vstore2(data2, 0, (__global float *)tensor3D_offset(&output, 0, 2, 0)); + vstore2(data3, 0, (__global float *)tensor3D_offset(&output, 0, 3, 0)); + vstore2(data4, 0, (__global float *)tensor3D_offset(&output, 0, 4, 0)); +} + +/** Computes the first stage of a radix-7 DFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -456,14 +622,14 @@ kernel void fft_radix_5_first_stage_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image */ kernel void fft_radix_7_first_stage_axis_0( TENSOR3D_DECLARATION(input) @@ -481,7 +647,7 @@ kernel void fft_radix_7_first_stage_axis_0( Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); #endif /* IN_PLACE */ - // Load eight complex input values + // Load seven complex input values float8 data0 = vload8(0, (__global float *)input.ptr); float4 data1 = vload4(0, (__global float *)tensor3D_offset(&input, 4, 0, 0)); float2 data2 = vload2(0, (__global float *)tensor3D_offset(&input, 6, 0, 0)); @@ -489,13 +655,72 @@ kernel void fft_radix_7_first_stage_axis_0( // Compute DFT N = 7 DFT_7(data0.s01, data0.s23, data0.s45, data0.s67, data1.s01, data1.s23, data2.s01); - // Store eight complex output values + // Store seven complex output values vstore8(data0, 0, (__global float *)output.ptr); vstore4(data1, 0, (__global float *)tensor3D_offset(&output, 4, 0, 0)); vstore2(data2, 0, (__global float *)tensor3D_offset(&output, 6, 0, 0)); } -/** Computes the first stage of a radix-8 DFT. +/** Computes the first stage of a radix-7 DFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + */ +kernel void fft_radix_7_first_stage_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ +) +{ + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); +#endif /* IN_PLACE */ + + // Load seven complex input values + float2 data0 = vload2(0, (__global float *)input.ptr); + float2 data1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 1, 0)); + float2 data2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2, 0)); + float2 data3 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 3, 0)); + float2 data4 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 4, 0)); + float2 data5 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 5, 0)); + float2 data6 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 6, 0)); + + // Compute DFT N = 7 + DFT_7(data0, data1, data2, data3, data4, data5, data6); + + // Store seven complex output values + vstore2(data0, 0, (__global float *)output.ptr); + vstore2(data1, 0, (__global float *)tensor3D_offset(&output, 0, 1, 0)); + vstore2(data2, 0, (__global float *)tensor3D_offset(&output, 0, 2, 0)); + vstore2(data3, 0, (__global float *)tensor3D_offset(&output, 0, 3, 0)); + vstore2(data4, 0, (__global float *)tensor3D_offset(&output, 0, 4, 0)); + vstore2(data5, 0, (__global float *)tensor3D_offset(&output, 0, 5, 0)); + vstore2(data6, 0, (__global float *)tensor3D_offset(&output, 0, 6, 0)); +} + +/** Computes the first stage of a radix-8 DFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -507,14 +732,14 @@ kernel void fft_radix_7_first_stage_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image */ kernel void fft_radix_8_first_stage_axis_0( TENSOR3D_DECLARATION(input) @@ -542,7 +767,68 @@ kernel void fft_radix_8_first_stage_axis_0( vstore16(data, 0, (__global float *)output.ptr); } -/** Computes a stage of a radix-2 FFT. +/** Computes the first stage of a radix-8 DFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + */ +kernel void fft_radix_8_first_stage_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ +) +{ + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); +#endif /* IN_PLACE */ + + // Load eight complex input values + float2 data0 = vload2(0, (__global float *)input.ptr); + float2 data1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 1, 0)); + float2 data2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2, 0)); + float2 data3 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 3, 0)); + float2 data4 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 4, 0)); + float2 data5 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 5, 0)); + float2 data6 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 6, 0)); + float2 data7 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 7, 0)); + + // Compute DFT N = 8 + DFT_8(data0, data1, data2, data3, data4, data5, data6, data7); + + // Store eight complex output values + vstore2(data0, 0, (__global float *)output.ptr); + vstore2(data1, 0, (__global float *)tensor3D_offset(&output, 0, 1, 0)); + vstore2(data2, 0, (__global float *)tensor3D_offset(&output, 0, 2, 0)); + vstore2(data3, 0, (__global float *)tensor3D_offset(&output, 0, 3, 0)); + vstore2(data4, 0, (__global float *)tensor3D_offset(&output, 0, 4, 0)); + vstore2(data5, 0, (__global float *)tensor3D_offset(&output, 0, 5, 0)); + vstore2(data6, 0, (__global float *)tensor3D_offset(&output, 0, 6, 0)); + vstore2(data7, 0, (__global float *)tensor3D_offset(&output, 0, 7, 0)); +} + +/** Computes a stage of a radix-2 FFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -554,14 +840,14 @@ kernel void fft_radix_8_first_stage_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage * @param[in] Ni Nx * Ny. * @param[in] exp_const Exponent constant @@ -612,7 +898,7 @@ kernel void fft_radix_2_axis_0( vstore2(c1, 0, (__global float *)tensor3D_offset(&output, Nx, 0, 0)); } -/** Computes a stage of a radix-3 FFT. +/** Computes a stage of a radix-2 FFT on axis 1. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -624,19 +910,19 @@ kernel void fft_radix_2_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage * @param[in] Ni Nx * Ny. * @param[in] exp_const Exponent constant */ -kernel void fft_radix_3_axis_0( +kernel void fft_radix_2_axis_1( TENSOR3D_DECLARATION(input) #ifndef IN_PLACE , @@ -645,8 +931,8 @@ kernel void fft_radix_3_axis_0( , uint Nx, uint Ni, float exp_const) { - // Each work-item computes a single radix-3 - uint kx = get_global_id(0); + // Each work-item computes a single radix-2 + uint kx = get_global_id(1); // Compute nx uint nx = kx % Nx; @@ -656,36 +942,33 @@ kernel void fft_radix_3_axis_0( // Get tensor pointers Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input); - input.ptr += n * input.stride_x + get_global_id(1) * input.stride_y + get_global_id(2) * input.stride_z; + input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z; #ifdef IN_PLACE Tensor3D output = input; #else /* IN_PLACE */ Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); - output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z; + output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z; #endif /* IN_PLACE */ - // Load three complex input values + // Load two complex input values float2 c0 = vload2(0, (__global float *)input.ptr); - float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, Nx, 0, 0)); - float2 c2 = vload2(0, (__global float *)tensor3D_offset(&input, 2 * Nx, 0, 0)); + float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, Nx, 0)); // Compute phi float phi = (float)nx * exp_const; // Multiply by twiddle factor TWIDDLE_FACTOR_MULTIPLICATION(phi, c1); - TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2); - // Compute DFT N = 3 - DFT_3(c0, c1, c2); + // Compute DFT N = 2 + DFT_2(c0, c1); - // Store three complex output values + // Store two complex output values vstore2(c0, 0, (__global float *)output.ptr); - vstore2(c1, 0, (__global float *)tensor3D_offset(&output, Nx, 0, 0)); - vstore2(c2, 0, (__global float *)tensor3D_offset(&output, 2 * Nx, 0, 0)); + vstore2(c1, 0, (__global float *)tensor3D_offset(&output, 0, Nx, 0)); } -/** Computes a stage of a radix-4 FFT. +/** Computes a stage of a radix-3 FFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -697,19 +980,19 @@ kernel void fft_radix_3_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage * @param[in] Ni Nx * Ny. * @param[in] exp_const Exponent constant */ -kernel void fft_radix_4_axis_0( +kernel void fft_radix_3_axis_0( TENSOR3D_DECLARATION(input) #ifndef IN_PLACE , @@ -718,7 +1001,7 @@ kernel void fft_radix_4_axis_0( , uint Nx, uint Ni, float exp_const) { - // Each work-item computes a single radix-4 + // Each work-item computes a single radix-3 uint kx = get_global_id(0); // Compute nx @@ -737,7 +1020,153 @@ kernel void fft_radix_4_axis_0( output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z; #endif /* IN_PLACE */ - // Load four complex input values + // Load three complex input values + float2 c0 = vload2(0, (__global float *)input.ptr); + float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, Nx, 0, 0)); + float2 c2 = vload2(0, (__global float *)tensor3D_offset(&input, 2 * Nx, 0, 0)); + + // Compute phi + float phi = (float)nx * exp_const; + + // Multiply by twiddle factor + TWIDDLE_FACTOR_MULTIPLICATION(phi, c1); + TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2); + + // Compute DFT N = 3 + DFT_3(c0, c1, c2); + + // Store three complex output values + vstore2(c0, 0, (__global float *)output.ptr); + vstore2(c1, 0, (__global float *)tensor3D_offset(&output, Nx, 0, 0)); + vstore2(c2, 0, (__global float *)tensor3D_offset(&output, 2 * Nx, 0, 0)); +} + +/** Computes a stage of a radix-3 FFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage + * @param[in] Ni Nx * Ny. + * @param[in] exp_const Exponent constant + */ +kernel void fft_radix_3_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ + , + uint Nx, uint Ni, float exp_const) +{ + // Each work-item computes a single radix-3 + uint kx = get_global_id(1); + + // Compute nx + uint nx = kx % Nx; + + // Compute n index + uint n = nx + (kx / Nx) * Ni; + + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input); + input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z; +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); + output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z; +#endif /* IN_PLACE */ + + // Load three complex input values + float2 c0 = vload2(0, (__global float *)input.ptr); + float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, Nx, 0)); + float2 c2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2 * Nx, 0)); + + // Compute phi + float phi = (float)nx * exp_const; + + // Multiply by twiddle factor + TWIDDLE_FACTOR_MULTIPLICATION(phi, c1); + TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2); + + // Compute DFT N = 3 + DFT_3(c0, c1, c2); + + // Store three complex output values + vstore2(c0, 0, (__global float *)output.ptr); + vstore2(c1, 0, (__global float *)tensor3D_offset(&output, 0, Nx, 0)); + vstore2(c2, 0, (__global float *)tensor3D_offset(&output, 0, 2 * Nx, 0)); +} + +/** Computes a stage of a radix-4 FFT on axis 0. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage + * @param[in] Ni Nx * Ny. + * @param[in] exp_const Exponent constant + */ +kernel void fft_radix_4_axis_0( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ + , + uint Nx, uint Ni, float exp_const) +{ + // Each work-item computes a single radix-4 + uint kx = get_global_id(0); + + // Compute nx + uint nx = kx % Nx; + + // Compute n index + uint n = nx + (kx / Nx) * Ni; + + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input); + input.ptr += n * input.stride_x + get_global_id(1) * input.stride_y + get_global_id(2) * input.stride_z; +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); + output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z; +#endif /* IN_PLACE */ + + // Load four complex input values float2 c0 = vload2(0, (__global float *)input.ptr); float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, Nx, 0, 0)); float2 c2 = vload2(0, (__global float *)tensor3D_offset(&input, 2 * Nx, 0, 0)); @@ -761,7 +1190,7 @@ kernel void fft_radix_4_axis_0( vstore2(c3, 0, (__global float *)tensor3D_offset(&output, 3 * Nx, 0, 0)); } -/** Computes a stage of a radix-5 FFT. +/** Computes a stage of a radix-4 FFT on axis 1. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -773,14 +1202,90 @@ kernel void fft_radix_4_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage + * @param[in] Ni Nx * Ny. + * @param[in] exp_const Exponent constant + */ +kernel void fft_radix_4_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ + , + uint Nx, uint Ni, float exp_const) +{ + // Each work-item computes a single radix-4 + uint kx = get_global_id(1); + + // Compute nx + uint nx = kx % Nx; + + // Compute n index + uint n = nx + (kx / Nx) * Ni; + + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input); + input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z; +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); + output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z; +#endif /* IN_PLACE */ + + // Load four complex input values + float2 c0 = vload2(0, (__global float *)input.ptr); + float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, Nx, 0)); + float2 c2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2 * Nx, 0)); + float2 c3 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 3 * Nx, 0)); + + // Compute phi + float phi = (float)nx * exp_const; + + // Multiply by twiddle factor + TWIDDLE_FACTOR_MULTIPLICATION(phi, c1); + TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2); + TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3); + + // Compute DFT N = 4 + DFT_4(c0, c1, c2, c3); + + // Store four complex output values + vstore2(c0, 0, (__global float *)output.ptr); + vstore2(c1, 0, (__global float *)tensor3D_offset(&output, 0, Nx, 0)); + vstore2(c2, 0, (__global float *)tensor3D_offset(&output, 0, 2 * Nx, 0)); + vstore2(c3, 0, (__global float *)tensor3D_offset(&output, 0, 3 * Nx, 0)); +} + +/** Computes a stage of a radix-5 FFT on axis 0. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage * @param[in] Ni Nx * Ny. * @param[in] exp_const Exponent constant @@ -840,7 +1345,86 @@ kernel void fft_radix_5_axis_0( vstore2(c4, 0, (__global float *)tensor3D_offset(&output, 4 * Nx, 0, 0)); } -/** Computes a stage of a radix-7 FFT. +/** Computes a stage of a radix-5 FFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage + * @param[in] Ni Nx * Ny. + * @param[in] exp_const Exponent constant + */ +kernel void fft_radix_5_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ + , + uint Nx, uint Ni, float exp_const) +{ + // Each work-item computes a single radix-5 + uint kx = get_global_id(1); + + // Compute nx + uint nx = kx % Nx; + + // Compute n index + uint n = nx + (kx / Nx) * Ni; + + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input); + input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z; +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); + output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z; +#endif /* IN_PLACE */ + + // Load five complex input values + float2 c0 = vload2(0, (__global float *)input.ptr); + float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, Nx, 0)); + float2 c2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2 * Nx, 0)); + float2 c3 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 3 * Nx, 0)); + float2 c4 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 4 * Nx, 0)); + + // Compute phi + float phi = (float)nx * exp_const; + + // Multiply by twiddle factor + TWIDDLE_FACTOR_MULTIPLICATION(phi, c1); + TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2); + TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3); + TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4); + + // Compute DFT N = 5 + DFT_5(c0, c1, c2, c3, c4); + + // Store five complex output values + vstore2(c0, 0, (__global float *)output.ptr); + vstore2(c1, 0, (__global float *)tensor3D_offset(&output, 0, Nx, 0)); + vstore2(c2, 0, (__global float *)tensor3D_offset(&output, 0, 2 * Nx, 0)); + vstore2(c3, 0, (__global float *)tensor3D_offset(&output, 0, 3 * Nx, 0)); + vstore2(c4, 0, (__global float *)tensor3D_offset(&output, 0, 4 * Nx, 0)); +} + +/** Computes a stage of a radix-7 FFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -852,14 +1436,14 @@ kernel void fft_radix_5_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage * @param[in] Ni Nx * Ny. * @param[in] exp_const Exponent constant @@ -925,7 +1509,92 @@ kernel void fft_radix_7_axis_0( vstore2(c6, 0, (__global float *)tensor3D_offset(&output, 6 * Nx, 0, 0)); } -/** Computes a stage of a radix-8 FFT. +/** Computes a stage of a radix-7 FFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage + * @param[in] Ni Nx * Ny. + * @param[in] exp_const Exponent constant + */ +kernel void fft_radix_7_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ + , + uint Nx, uint Ni, float exp_const) +{ + // Each work-item computes a single radix-7 + uint kx = get_global_id(1); + + // Compute nx + uint nx = kx % Nx; + + // Compute n index + uint n = nx + (kx / Nx) * Ni; + + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input); + input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z; +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); + output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z; +#endif /* IN_PLACE */ + + // Load seven complex input values + float2 c0 = vload2(0, (__global float *)input.ptr); + float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, Nx, 0)); + float2 c2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2 * Nx, 0)); + float2 c3 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 3 * Nx, 0)); + float2 c4 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 4 * Nx, 0)); + float2 c5 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 5 * Nx, 0)); + float2 c6 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 6 * Nx, 0)); + + // Compute phi + float phi = (float)nx * exp_const; + + // Multiply by twiddle factor + TWIDDLE_FACTOR_MULTIPLICATION(phi, c1); + TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2); + TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3); + TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4); + TWIDDLE_FACTOR_MULTIPLICATION(5 * phi, c5); + TWIDDLE_FACTOR_MULTIPLICATION(6 * phi, c6); + + // Compute DFT N = 7 + DFT_7(c0, c1, c2, c3, c4, c5, c6); + + // Store seven complex output values + vstore2(c0, 0, (__global float *)output.ptr); + vstore2(c1, 0, (__global float *)tensor3D_offset(&output, 0, Nx, 0)); + vstore2(c2, 0, (__global float *)tensor3D_offset(&output, 0, 2 * Nx, 0)); + vstore2(c3, 0, (__global float *)tensor3D_offset(&output, 0, 3 * Nx, 0)); + vstore2(c4, 0, (__global float *)tensor3D_offset(&output, 0, 4 * Nx, 0)); + vstore2(c5, 0, (__global float *)tensor3D_offset(&output, 0, 5 * Nx, 0)); + vstore2(c6, 0, (__global float *)tensor3D_offset(&output, 0, 6 * Nx, 0)); +} + +/** Computes a stage of a radix-8 FFT on axis 0. * * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time * @@ -937,14 +1606,14 @@ kernel void fft_radix_7_axis_0( * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor - * @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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage * @param[in] Ni Nx * Ny. * @param[in] exp_const Exponent constant @@ -1011,4 +1680,92 @@ kernel void fft_radix_8_axis_0( vstore2(c5, 0, (__global float *)tensor3D_offset(&output, 5 * Nx, 0, 0)); vstore2(c6, 0, (__global float *)tensor3D_offset(&output, 6 * Nx, 0, 0)); vstore2(c7, 0, (__global float *)tensor3D_offset(&output, 7 * Nx, 0, 0)); +} + +/** Computes a stage of a radix-8 FFT on axis 1. + * + * @note In order to perform the FFT function "in-place", the pre-processor -DIN_PLACE must be passed at compile time + * + * @param[in,out] input_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in,out] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in,out] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in,out] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in,out] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in,out] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in,out] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in,out] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr (Optional) Pointer to the destination image. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination image in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination image in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] output_step_z (Optional) output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination image + * @param[in] Nx The butterfly span. Products of radix order of previous radix's stage + * @param[in] Ni Nx * Ny. + * @param[in] exp_const Exponent constant + */ +kernel void fft_radix_8_axis_1( + TENSOR3D_DECLARATION(input) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(output) +#endif /* not IN_PLACE */ + , + uint Nx, uint Ni, float exp_const) +{ + // Each work-item computes a single radix-8 + uint kx = get_global_id(1); + + // Compute nx + uint nx = kx % Nx; + + // Compute n index + uint n = nx + (kx / Nx) * Ni; + + // Get tensor pointers + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input); + input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z; +#ifdef IN_PLACE + Tensor3D output = input; +#else /* IN_PLACE */ + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); + output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z; +#endif /* IN_PLACE */ + + // Load eight complex input values + float2 c0 = vload2(0, (__global float *)input.ptr); + float2 c1 = vload2(0, (__global float *)tensor3D_offset(&input, 0, Nx, 0)); + float2 c2 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 2 * Nx, 0)); + float2 c3 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 3 * Nx, 0)); + float2 c4 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 4 * Nx, 0)); + float2 c5 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 5 * Nx, 0)); + float2 c6 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 6 * Nx, 0)); + float2 c7 = vload2(0, (__global float *)tensor3D_offset(&input, 0, 7 * Nx, 0)); + + // Compute phi + float phi = (float)nx * exp_const; + + // Multiply by twiddle factor + TWIDDLE_FACTOR_MULTIPLICATION(phi, c1); + TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2); + TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3); + TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4); + TWIDDLE_FACTOR_MULTIPLICATION(5 * phi, c5); + TWIDDLE_FACTOR_MULTIPLICATION(6 * phi, c6); + TWIDDLE_FACTOR_MULTIPLICATION(7 * phi, c7); + + // Compute DFT N = 8 + DFT_8(c0, c1, c2, c3, c4, c5, c6, c7); + + // Store eight complex output values + vstore2(c0, 0, (__global float *)output.ptr); + vstore2(c1, 0, (__global float *)tensor3D_offset(&output, 0, Nx, 0)); + vstore2(c2, 0, (__global float *)tensor3D_offset(&output, 0, 2 * Nx, 0)); + vstore2(c3, 0, (__global float *)tensor3D_offset(&output, 0, 3 * Nx, 0)); + vstore2(c4, 0, (__global float *)tensor3D_offset(&output, 0, 4 * Nx, 0)); + vstore2(c5, 0, (__global float *)tensor3D_offset(&output, 0, 5 * Nx, 0)); + vstore2(c6, 0, (__global float *)tensor3D_offset(&output, 0, 6 * Nx, 0)); + vstore2(c7, 0, (__global float *)tensor3D_offset(&output, 0, 7 * Nx, 0)); } \ No newline at end of file diff --git a/src/core/CL/cl_kernels/fft_digit_reverse.cl b/src/core/CL/cl_kernels/fft_digit_reverse.cl new file mode 100644 index 0000000000..040c2846bd --- /dev/null +++ b/src/core/CL/cl_kernels/fft_digit_reverse.cl @@ -0,0 +1,148 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers.h" + +#if defined(VEC_SIZE) +/** Computes the digit reverse stage on axis X + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: 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_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_z * number of elements along Z 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. 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_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] idx_ptr Pointer to the index tensor. Supported data types: U32 + * @param[in] idx_stride_x Stride of the index tensor in X dimension (in bytes) + * @param[in] idx_step_x idx_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] idx_offset_first_element_in_bytes The offset of the first element in the index tensor + */ +__kernel void fft_digit_reverse_axis_0( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + VECTOR_DECLARATION(idx)) +{ + // Get tensor pointers + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(src); + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); + Vector idx = CONVERT_TO_VECTOR_STRUCT(idx); + + const unsigned int iidx = *((__global uint *)(idx.ptr)); + + // Load data +#if VEC_SIZE == 1 + float data = *((__global float *)tensor3D_offset(&src, iidx, get_global_id(1), get_global_id(2))); +#elif VEC_SIZE == 2 + float2 data = vload2(0, (__global float *)tensor3D_offset(&src, iidx, get_global_id(1), get_global_id(2))); +#else // VEC_SIZE == 1 +#error "vec_size of 1 and 2 are supported" +#endif // VEC_SIZE == 1 + + // Create result +#if VEC_SIZE == 1 + float2 res = { data, 0 }; +#elif VEC_SIZE == 2 + float2 res = data; +#else // VEC_SIZE == 1 +#error "vec_size of 1 and 2 are supported" +#endif // VEC_SIZE == 1 + + // Store result +#if defined(CONJ) + vstore2((float2)(res.s0, -res.s1), 0, (__global float *)dst.ptr); +#else // defined(CONJ) + vstore2(res, 0, (__global float *)dst.ptr); +#endif // defined(CONJ) +} + +/** Computes the digit reverse stage on axis Y + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: 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_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_z * number of elements along Z 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. 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_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] idx_ptr Pointer to the index tensor. Supported data types: U32 + * @param[in] idx_stride_x Stride of the index tensor in X dimension (in bytes) + * @param[in] idx_step_x idx_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] idx_offset_first_element_in_bytes The offset of the first element in the index tensor + */ +__kernel void fft_digit_reverse_axis_1( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + VECTOR_DECLARATION(idx)) +{ + // Get tensor pointers + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(src); + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); + Vector idx = CONVERT_TO_VECTOR_STRUCT_NO_STEP(idx); + + const unsigned int iidx = *((__global uint *)vector_offset(&idx, (int)(get_global_id(1)))); + + // Load data +#if VEC_SIZE == 1 + float data = *((__global float *)tensor3D_offset(&src, get_global_id(0), iidx, get_global_id(2))); +#elif VEC_SIZE == 2 + float2 data = vload2(0, (__global float *)tensor3D_offset(&src, get_global_id(0), iidx, get_global_id(2))); +#else // VEC_SIZE == 1 +#error "vec_size of 1 and 2 are supported" +#endif // VEC_SIZE == 1 + + // Create result +#if VEC_SIZE == 1 + float2 res = { data, 0 }; +#elif VEC_SIZE == 2 + float2 res = data; +#else // VEC_SIZE == 1 +#error "vec_size of 1 and 2 are supported" +#endif // VEC_SIZE == 1 + + // Store result +#if defined(CONJ) + vstore2((float2)(res.s0, -res.s1), 0, (__global float *)dst.ptr); +#else // defined(CONJ) + vstore2(res, 0, (__global float *)dst.ptr); +#endif // defined(CONJ) +} +#endif // defined(VEC_SIZE) \ No newline at end of file diff --git a/src/core/CL/cl_kernels/fft_scale.cl b/src/core/CL/cl_kernels/fft_scale.cl new file mode 100644 index 0000000000..bf78a26eb8 --- /dev/null +++ b/src/core/CL/cl_kernels/fft_scale.cl @@ -0,0 +1,78 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers.h" + +/** Computes the fft scale stage + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: 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_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_z * number of elements along Z 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 (Optional) Pointer to the destination tensor. Supported data types: same as @p src_ptr + * @param[in] dst_stride_x (Optional) Stride of the destination tensor in X dimension (in bytes) + * @param[in] dst_step_x (Optional) dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y (Optional) Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y (Optional) dst_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] dst_stride_z (Optional) Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z (Optional) dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes (Optional) The offset of the first element in the destination tensor + * @param[in] scale Scale to apply to the complex value + */ +__kernel void fft_scale_conj( + TENSOR3D_DECLARATION(src) +#ifndef IN_PLACE + , + TENSOR3D_DECLARATION(dst) +#endif /* not IN_PLACE */ + , + float scale) +{ + // Get tensor pointers + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); +#if defined(IN_PLACE) + Tensor3D dst = src; +#else /* IN_PLACE */ + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); +#endif /* IN_PLACE */ + + // Store result +#if VEC_SIZE == 1 + *((__global float *)dst.ptr) = (*(__global float *)src.ptr) / scale; +#elif VEC_SIZE == 2 + // Load data + float2 data = vload2(0, (__global float *)src.ptr); + data /= scale; +#if defined(CONJ) + vstore2((float2)(data.s0, -data.s1), 0, (__global float *)dst.ptr); +#else // defined(CONJ) + vstore2(data, 0, (__global float *)dst.ptr); +#endif // defined(CONJ) +#else // VEC_SIZE == 1 +#error "vec_size of 1 and 2 are supported" +#endif // VEC_SIZE == 1 +} \ No newline at end of file diff --git a/src/core/CL/cl_kernels/pixelwise_mul_float.cl b/src/core/CL/cl_kernels/pixelwise_mul_float.cl index 9fa540e946..d0e04b2ffe 100644 --- a/src/core/CL/cl_kernels/pixelwise_mul_float.cl +++ b/src/core/CL/cl_kernels/pixelwise_mul_float.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -94,4 +94,52 @@ __kernel void pixelwise_mul_float( // Store result vstore16(res, 0, (__global DATA_TYPE_OUT *)out.ptr); } -#endif /* defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_RES) && defined(DATA_TYPE_OUT) */ \ No newline at end of file +#endif /* defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_RES) && defined(DATA_TYPE_OUT) */ + +/** Performs a pixelwise multiplication of complex float values + * + * @param[in] in1_ptr Pointer to the source image. Supported data types: F32 + * @param[in] in1_stride_x Stride of the source image in X dimension (in bytes) + * @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: same as @p in1_ptr + * @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: same as @p in1_ptr + * @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 + */ +__kernel void pixelwise_mul_complex( + TENSOR3D_DECLARATION(in1), + TENSOR3D_DECLARATION(in2), + TENSOR3D_DECLARATION(out)) +{ + // Get pixels pointer + Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1); + Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2); + Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out); + + // Load data + float2 vin1 = vload2(0, (__global float *)in1.ptr); + float2 vin2 = vload2(0, (__global float *)in2.ptr); + + // Perform complex multiplication + float2 res = { vin1.x *vin2.x - vin1.y * vin2.y, vin1.x *vin2.y + vin2.x * vin1.y }; + + // Store result + vstore2(res, 0, (__global float *)out.ptr); +} diff --git a/src/core/CL/cl_kernels/reduction_operation.cl b/src/core/CL/cl_kernels/reduction_operation.cl index b4ede25296..2651123cf5 100644 --- a/src/core/CL/cl_kernels/reduction_operation.cl +++ b/src/core/CL/cl_kernels/reduction_operation.cl @@ -307,6 +307,10 @@ __kernel void reduction_operation_z( VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); +#if defined(COMPLEX) + VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) + res1 = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 8, 0, 0)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); +#endif // defined(COMPLEX) #if defined(SUM_SQUARE) res *= res; #endif // defined(SUM_SQUARE) @@ -320,6 +324,11 @@ __kernel void reduction_operation_z( VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, z)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); +#if defined(COMPLEX) + VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) + in1 = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 8, 0, z)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); +#endif // defined(COMPLEX) + #if defined(ARG_MAX) uint16 cond_conv = CONVERT(isgreater(in, res), uint16); indx = select(indx, z, cond_conv); @@ -334,8 +343,11 @@ __kernel void reduction_operation_z( #endif // defined(SUM_SQUARE) #if defined(PROD) res *= in; -#else //!defined(PROD) +#else //!defined(PROD) res += in; +#if defined(COMPLEX) + res1 += in1; +#endif // defined(COMPLEX) #endif //defined(PROD) #endif // defined(ARG_MAX) || defined(ARG_MIN) } @@ -348,6 +360,9 @@ __kernel void reduction_operation_z( res /= DEPTH; #endif // defined(MEAN) vstore16(CONVERT(res, VEC_DATA_TYPE(DATA_TYPE, 16)), 0, (__global DATA_TYPE *)output.ptr); +#if defined(COMPLEX) + vstore16(CONVERT(res1, VEC_DATA_TYPE(DATA_TYPE, 16)), 0, (__global DATA_TYPE *)tensor3D_offset(&output, 8, 0, 0)); +#endif // defined(COMPLEX) #endif // defined(ARG_MAX) || defined(ARG_MIN) } #endif /* defined(DEPTH) */ -- cgit v1.2.1