Speeding-up Fast Fourier Transform Mixed-Radix on Mobile ARM Mali GPU by means of OpenCL - Part 2

February 1, 2016

24 minute read time.

Back to Speeding-up Fast Fourier Transform Mixed-Radix on Mobile ARM Mali GPU by means of OpenCL - Part 1

Read Speeding-up Fast Fourier Transform Mixed-Radix on Mobile ARM Mali GPU by means of OpenCL - Part 3

Here we are for the second part of our blog series about the OpenCL™ implementation of Complex to Complex Fast Fourier Transform based on the method mixed-radix on Mobile ARM® Mali™ GPUs.

Whilst in the first article - Speeding-up Fast Fourier Transform Mixed-Radix on Mobile ARM Mali GPU by means of OpenCL - Part 1 - we presented the basic mathematical background, in the following we are going to play with the 3 main computation blocks behind the FFT mixed-radix which will be analyzed in a step-by-step approach from the point of view of both theory and of its efficient and simple implementation with OpenCL.

The development platform used to carry out our performance analysis will be the Firefly-RK3288 which is an all-in-one high performance multi-core computing platform based on the Rockchip RK3288 SoC featuring an ARM^® Cortex^®-A17 @1.8GHz and ARM^® Mali™-T760 (Mali GPU Driver: r6p0-02rel0).

For all performance evaluations, DVFS has been disabled and the GPU clock has been set at the maximum frequency of 600MHz.

Further information about the development platform can be found at:

Implementation

Let's take a look at how we store our complex values. The input and output data are sorted into a floating point complex array with 2*N elements where the real and imaginary parts are placed alternately. Two different buffers with 2*N elements have been used to store the input and output values.

The FFT mixed-radix pipeline that we are going to study is composed mainly of 3 computation blocks:

Digit-reverse
Twiddle factor multiplication
Radix computation

Indexing scheme

The only real complexity of FFT mixed-radix is the indexing scheme. The indexing scheme expresses the relation between the index n and k between 2 generic stages X and Y. This relation is fundamental because it allows both to know which input must be processed by each radix basic element and to correctly perform the stages of digit reverse and twiddle factors multiplication.

Given 2 generic stages X and Y:

n addresses the values in the stage X: n = nx + ny * Nx
k addresses the values in the stage Y: k = ky + kx * Ny

with:

nx = n % Nx, scans the columns - nx [0, Nx - 1]
ny = floor(n / Nx), scans the rows - ny [0, N / Nx - 1]
kx = floor(k / Ny), scans the rows - kx [0, N / Ny - 1]
ky = (k % Ny), scans the columns - ky [0, Ny - 1]

From the above picture we can easily see what nx, ny, kx and ky are. For what concerns Nx and Ny, in the case of just 2 radix stages they are respectively the radix order of stage X and Y. In the case of more than 2 radix stages, Ny is still the radix order of stage Y but Nx becomes the radix products from the first radix stage to the radix stage X.

Nx is also known as span or butterfly span. The span is the offset between each complex input of radix basic element and it was introduced in the first article in reference to the radix-2 FFT algorithm.

i.e.

Let's assume we have N = 8 x 4 x 4 x 3 x 2 (5 stages).

If the stages X and Y were stage 2 and stage 3, Nx would be Nx = 8 x 4 x 4 = 128 and Ny = 3

With a pipeline made up of M radix stages, we would like to have a relation between the index k and n between 2 generic stages, in particular we would like to have something like:

The general expressions for computing these mapping functions are:

where Ni is:

/**
 * @brief Map index k to index n
 *
 * @param[in] k  Index k
 * @param[in] Nx It is the span
 * @param[in] Ny It is the radix order of stage Y
 *
 * @return       Index n
 */
uint map_k_to_n(uint k, uint Nx, uint Ny)
{
    uint Ni = Nx * Ny;
    uint ky = k % Ny;    // Remainder
    uint kx = k / Ny;    // Integer part
    uint n = (kx % Nx) + (kx / Nx) * Ni + ky * Nx;
    return n;
}

/**
 * @brief Map index n to index k
 *
 * @param[in] n  Index n
 * @param[in] Nx It is the span
 * @param[in] Ny It is the radix order of stage Y
 *
 * @return       Index k
 */
uint map_n_to_k(uint n, uint Nx, uint Ny)
{
    uint Ni = Nx * Ny;
    uint k = (n * Ny) % Ni + (n / Nx) % Ny + Ni * (n / Ni);
    return k;
}

Every time we compute a radix stage, it is good to keep the span Nx update.

/* Init Nx to 1 */
uint Nx = 1;

/* Scan each radix stage */
for(uint s = 0; s < n_stages; ++s)
{
     /* Get radix order of stage s */
     uint Ny = radix[s];
     
     /* Body for computing twiddle factor multiplication and radix computation */
     ...
     
     /* Update Nx */
     Nx *= Ny;
}

Digit-reverse

Once we have introduced the indexing scheme, we are ready to analyze the main computation blocks of our pipeline. Let's start with the digit-reverse.

The digit-reverse is the first stage of our pipeline which places the input elements in a specific order called "digit-reverse order".

The digit-reverse order would be exactly the order of the output elements if we left the input elements in linear order (0, 1, 2,...).

Since we know the relation between the index n and k, a possible way of knowing the digit-reverse order may be to iterate the mapping function map_n_to_k() from the first to the last stage in order to know how the input index n would be mapped at the end.

/**
 * @brief This function computes the digit reverse index for each input
 *
 * @param[in]  stage             It contains the radix order for each radix stage
 * @param[out] idx_digit_reverse It contains the digit-reverse order index
 * @param[in]  n_stages          Total number of radix stages
 * @param[in]  N                 Number of input
 */
void digit_reverse(float* stage, uint* idx_digit_reverse, uint n_stages, uint N)
{
    /* Scan elements */
    for(uint n = 0; n < N; ++n)
    {
        uint k = n;
        uint Nx = stage[0];

        /* Scan stages */
        for (uint s = 1; s < n_stages; ++s)
        {
            /* radix of stage s-th */
            uint Ny = stage[s];
            uint Ni = Ny * Nx;

            /* Update k index */
            k = (k * Ny) % Ni + (k / Nx) % Ny + Ni * (k / Ni);

            /* Update Nx */
            Nx *= Ny;
        }

        /* K is the index of digit-reverse */
        idx_digit_reverse[n] = k;
    }
}

Once we know the digit-reverse index, we can implement a CL kernel that places the input values in the required order.

/**
 * @brief This kernel stores the input in digit-reverse order
 *
 * @param[in]  input            It contains the input complex values in linear order 
 * @param[out] output           It contains the output complex values in digit-reverse order
 * @param[in]  idx_bit_reverse  It contains the digit-reverse order index
 */
kernel void digit_reverse(global float2* input, global float2* output, global uint* idx_digit_reverse)
{
    /* Each work-item stores a single complex values */
    const uint n = get_global_id(0);

    /* Get digit-reverse index */
    const uint idx = idx_digit_reverse[n];

    /* Get complex value */
    float2 val = (float2)(input[idx]);

    /* Store complex value */
    output[n] = val;
}

The digit-reverse kernel uses 2 different cl_buffers:

one for loading the input complex values
one for storing the output complex values in digit-reverse order.

Twiddle Factor Multiplication

Twiddle factor multiplication is the intermediate stage between 2 generic radix stages X and Y. In particular it multiplies all the N output complex elements of radix stage X by a specific trigonometric complex constant before passing those values to the next stage Y.

This trigonometric complex constant -the "twiddle" factor - depends on nx and ky:

Even the twiddle factor multiplication is highly parallelizable and can be implemented by means of OpenCL in the following manner:

#define M_2PI_F 6.283185482025146484375f

#define TWIDDLE_FACTOR_MULTIPLICATION(phi, input)                 \
{                                                                 \
    float2 w, tmp;                                                \
    w.x = cos(phi);                                               \
    w.y = sin(phi);                                               \          
    tmp.x = (w.x * input.x) - (w.y * input.y);                    \
    tmp.y = (w.x * input.y) + (w.y * input.x);                    \
    input = tmp;                                                  \
}

/**
 * @brief This kernel computes the twiddle factor multiplication between 2 generic radix stages X and Y
 * 
 * @param[in, out] input It contains the input and output complex values 
 * @param[in]      Nx    It is the span
 * @param[in]      Ny    It is the radix order of stage Y
 * @param[in]      Ni    Nx * Ny
 */
kernel void twiddle_factor_multiplication(global float2* input, const uint Nx, const uint Ny, const uint Ni)
{
    /* Each work-item computes a single complex input */
    const uint n = get_global_id(0);

    /* Compute nx */
    uint nx = n % Nx;

    /* Compute k index */
    uint k = (n * Ny) % Ni + (n / Nx) % Ny + Ni * (n / Ni);

    /* Compute ky */
    uint ky = k % Ny;

    /* Compute angle of twiddle factor */
    float phi = (-M_2PI_F * nx * ky) / (float)Ni;

    /* Multiply by twiddle factor */
    TWIDDLE_FACTOR_MULTIPLICATION(phi, input[n]);
}

The above kernel uses the same cl_buffer for loading and storing.

An interesting optimization concerns the computation of ky. As we know,

so:

since in the above expression A and B are multiples of Ny, the terms 1 and 3 will be 0 therefore:

but for the identity property of the modulo operator we have:

and then we can compute ky as:

At this point we can simplify our OpenCL kernel in the following manner:

kernel void twiddle_factor_multiplication(global float2* input, const uint Nx, const uint Ny, const uint Ni)
{
    /* Each work-item computes a single complex input */
    const uint n = get_global_id(0);

    /* Compute nx */
    uint nx = n % Nx;

    /* Compute ky */
    uint ky = (n / Nx) % Ny;               // <-

    /* Compute angle of twiddle factor */
    float phi = (-M_2PI_F * nx * ky) / (float)Ni;

    /* Multiply by twiddle factor */
    TWIDDLE_FACTOR_MULTIPLICATION(phi, input[n]);
}

Another important and simple optimization that we can introduce concerns the computation of the angle for the twiddle factor multiplication. As we can observe, the term (-M_2PI_F / (float)Ni) is a constant for all work-items and can therefore be passed as a argument to the CL kernel.

/**
 * @brief This kernel computes the twiddle factor multiplication between 2 generic radix stages X and Y
 * 
 * @param[in, out] input     It contains the input and output complex values 
 * @param[in]      Nx        It is the span
 * @param[in]      Ny        It is the radix order of stage Y
 * @param[in]      Ni        Nx * Ny
 * @param[in]      exp_const (-M_2PI_F / (float)Ni)
 */
kernel void twiddle_factor_multiplication(global float2* input, const uint Nx, const uint Ny, const uint Ni, const float exp_const)
{
    /* Each work-item computes a single complex input */
    const uint n = get_global_id(0);

    /* Compute nx */
    uint nx = n % Nx;

    /* Compute ky */
    uint ky = (n / Nx) % Ny;

    /* Compute angle of twiddle factor */
    float phi = (float)(nx * ky) * exp_const;                    // <-

    /* Multiply by twiddle factor */
    TWIDDLE_FACTOR_MULTIPLICATION(phi, input[n]);
}

With these 2 simple optimizations, we are able to improve the performance of this kernel by a factor of ~1.6/1.7x as illustrated in the following graph.

Radix computation

The radix stage is the main computation block of the pipeline composed of N / S radix-S basic elements.

Each radix-S consists of S inputs and S outputs and computes an optimized DFT of length S.

Since there is not a dependency between each radix basic element of the same stage, even the radix stage is an embarrassingly parallel problem.

In our implementation each radix-S of the same radix stage is performed by a single work-item. As in the first stage the radix basic element takes the input in linear order whilst for the following stages accordingly with the span Nx, we are going to use 2 different kernels for the radix basic element:

One kernel for the first stage which takes the input in linear order
One kernel for the other radix stages

Both kernels use a single cl_buffer for loading and storing.

In the following we are going to describe the implementation of radix 2/3/4/5/7.