From ea7de7babc319e2fa31c5e1c986e48d6c5370689 Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Thu, 10 Dec 2020 16:49:39 +0000 Subject: Enable FFT for FP16 Resolves: COMPMID-4051 Change-Id: I0c0bf97212dd281c19d5081e6247e7dc0c23cd6b Signed-off-by: Giorgio Arena Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4687 Tested-by: Arm Jenkins Reviewed-by: Gian Marco Iodice Comments-Addressed: Arm Jenkins --- src/core/CL/cl_kernels/fft_scale.cl | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) (limited to 'src/core/CL/cl_kernels/fft_scale.cl') 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 -- cgit v1.2.1