1/* 2 * Copyright (c) 2019-2021 Arm Limited. 3 * 4 * SPDX-License-Identifier: MIT 5 * 6 * Permission is hereby granted, free of charge, to any person obtaining a copy 7 * of this software and associated documentation files (the "Software"), to 8 * deal in the Software without restriction, including without limitation the 9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or 10 * sell copies of the Software, and to permit persons to whom the Software is 11 * furnished to do so, subject to the following conditions: 12 * 13 * The above copyright notice and this permission notice shall be included in all 14 * copies or substantial portions of the Software. 15 * 16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE 22 * SOFTWARE. 23 */ 24#include "helpers.h" 25 26#if defined(VEC_SIZE) && defined(DATA_TYPE) 27/** Computes the digit reverse stage on axis X 28 * 29 * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32 30 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) 31 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) 32 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) 33 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) 34 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) 35 * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) 36 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor 37 * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr 38 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) 39 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) 40 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) 41 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) 42 * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) 43 * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) 44 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor 45 * @param[in] idx_ptr Pointer to the index tensor. Supported data types: U32 46 * @param[in] idx_stride_x Stride of the index tensor in X dimension (in bytes) 47 * @param[in] idx_step_x idx_stride_x * number of elements along X processed per workitem(in bytes) 48 * @param[in] idx_offset_first_element_in_bytes The offset of the first element in the index tensor 49 */ 50__kernel void fft_digit_reverse_axis_0( 51 TENSOR3D_DECLARATION(src), 52 TENSOR3D_DECLARATION(dst), 53 VECTOR_DECLARATION(idx)) 54{ 55 // Get tensor pointers 56 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(src); 57 Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); 58 Vector idx = CONVERT_TO_VECTOR_STRUCT(idx); 59 60 const unsigned int iidx = *((__global uint *)(idx.ptr)); 61 62 // Load data 63#if VEC_SIZE == 1 64 DATA_TYPE data = *((__global DATA_TYPE *)tensor3D_offset(&src, iidx, get_global_id(1), get_global_id(2))); 65#elif VEC_SIZE == 2 66 VEC_DATA_TYPE(DATA_TYPE, 2) 67 data = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&src, iidx, get_global_id(1), get_global_id(2))); 68#else // VEC_SIZE == 1 69#error "vec_size of 1 and 2 are supported" 70#endif // VEC_SIZE == 1 71 72 // Create result 73#if VEC_SIZE == 1 74 VEC_DATA_TYPE(DATA_TYPE, 2) 75 res = { data, 0 }; 76#elif VEC_SIZE == 2 77 VEC_DATA_TYPE(DATA_TYPE, 2) 78 res = data; 79#else // VEC_SIZE == 1 80#error "vec_size of 1 and 2 are supported" 81#endif // VEC_SIZE == 1 82 83 // Store result 84#if defined(CONJ) 85 vstore2((VEC_DATA_TYPE(DATA_TYPE, 2))(res.s0, -res.s1), 0, (__global DATA_TYPE *)dst.ptr); 86#else // defined(CONJ) 87 vstore2(res, 0, (__global DATA_TYPE *)dst.ptr); 88#endif // defined(CONJ) 89} 90 91/** Computes the digit reverse stage on axis Y 92 * 93 * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32 94 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) 95 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) 96 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) 97 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) 98 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) 99 * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) 100 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor 101 * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr 102 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) 103 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) 104 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) 105 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) 106 * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) 107 * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) 108 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor 109 * @param[in] idx_ptr Pointer to the index tensor. Supported data types: U32 110 * @param[in] idx_stride_x Stride of the index tensor in X dimension (in bytes) 111 * @param[in] idx_step_x idx_stride_x * number of elements along X processed per workitem(in bytes) 112 * @param[in] idx_offset_first_element_in_bytes The offset of the first element in the index tensor 113 */ 114__kernel void fft_digit_reverse_axis_1( 115 TENSOR3D_DECLARATION(src), 116 TENSOR3D_DECLARATION(dst), 117 VECTOR_DECLARATION(idx)) 118{ 119 // Get tensor pointers 120 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(src); 121 Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); 122 Vector idx = CONVERT_TO_VECTOR_STRUCT_NO_STEP(idx); 123 124 const unsigned int iidx = *((__global uint *)vector_offset(&idx, (int)(get_global_id(1)))); 125 126 // Load data 127#if VEC_SIZE == 1 128 DATA_TYPE data = *((__global DATA_TYPE *)tensor3D_offset(&src, get_global_id(0), iidx, get_global_id(2))); 129#elif VEC_SIZE == 2 130 VEC_DATA_TYPE(DATA_TYPE, 2) 131 data = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&src, get_global_id(0), iidx, get_global_id(2))); 132#else // VEC_SIZE == 1 133#error "vec_size of 1 and 2 are supported" 134#endif // VEC_SIZE == 1 135 136 // Create result 137#if VEC_SIZE == 1 138 VEC_DATA_TYPE(DATA_TYPE, 2) 139 res = { data, 0 }; 140#elif VEC_SIZE == 2 141 VEC_DATA_TYPE(DATA_TYPE, 2) 142 res = data; 143#else // VEC_SIZE == 1 144#error "vec_size of 1 and 2 are supported" 145#endif // VEC_SIZE == 1 146 147 // Store result 148#if defined(CONJ) 149 vstore2((VEC_DATA_TYPE(DATA_TYPE, 2))(res.s0, -res.s1), 0, (__global DATA_TYPE *)dst.ptr); 150#else // defined(CONJ) 151 vstore2(res, 0, (__global DATA_TYPE *)dst.ptr); 152#endif // defined(CONJ) 153} 154#endif // defined(VEC_SIZE) && defined(DATA_TYPE)