diff options
author | Giorgio Arena <giorgio.arena@arm.com> | 2020-12-10 16:49:39 +0000 |
---|---|---|
committer | Giorgio Arena <giorgio.arena@arm.com> | 2020-12-14 13:58:17 +0000 |
commit | ea7de7babc319e2fa31c5e1c986e48d6c5370689 (patch) | |
tree | 2303791668c67eda76dfb14d07b912af1cb54a17 /src/core/CL/cl_kernels/fft_scale.cl | |
parent | ec241b48ea7481e797285788fd68e5e1d42382bb (diff) | |
download | ComputeLibrary-ea7de7babc319e2fa31c5e1c986e48d6c5370689.tar.gz |
Enable FFT for FP16
Resolves: COMPMID-4051
Change-Id: I0c0bf97212dd281c19d5081e6247e7dc0c23cd6b
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4687
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/fft_scale.cl')
-rw-r--r-- | src/core/CL/cl_kernels/fft_scale.cl | 19 |
1 files changed, 11 insertions, 8 deletions
diff --git a/src/core/CL/cl_kernels/fft_scale.cl b/src/core/CL/cl_kernels/fft_scale.cl index 270fb78ae2..57e25ef504 100644 --- a/src/core/CL/cl_kernels/fft_scale.cl +++ b/src/core/CL/cl_kernels/fft_scale.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 Arm Limited. + * Copyright (c) 2019-2020 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -23,9 +23,10 @@ */ #include "helpers.h" +#if defined(VEC_SIZE) && defined(DATA_TYPE) /** Computes the fft scale stage * - * @param[in] src_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/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) @@ -62,17 +63,19 @@ __kernel void fft_scale_conj( // Store result #if VEC_SIZE == 1 - *((__global float *)dst.ptr) = (*(__global float *)src.ptr) / scale; + *((__global DATA_TYPE *)dst.ptr) = (*(__global DATA_TYPE *)src.ptr) / (DATA_TYPE)scale; #elif VEC_SIZE == 2 // Load data - float2 data = vload2(0, (__global float *)src.ptr); - data /= scale; + VEC_DATA_TYPE(DATA_TYPE, 2) + data = vload2(0, (__global DATA_TYPE *)src.ptr); + data /= (DATA_TYPE)scale; #if defined(CONJ) - vstore2((float2)(data.s0, -data.s1), 0, (__global float *)dst.ptr); + vstore2((VEC_DATA_TYPE(DATA_TYPE, 2))(data.s0, -data.s1), 0, (__global DATA_TYPE *)dst.ptr); #else // defined(CONJ) - vstore2(data, 0, (__global float *)dst.ptr); + vstore2(data, 0, (__global DATA_TYPE *)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 +} +#endif // defined(VEC_SIZE) && defined(DATA_TYPE)
\ No newline at end of file |