__kernel void kernel0(__global const uchar4 *srcBuffer,
                      __global uchar4 *dstBuffer,
                      const uint bufferWidth,
                      __constant float4 *factorValues,
                      const uint iterationsCount,
                      const uint iterationsInnerCount,
                      __local uchar4 *localBuffer)
{
    const uint groupWidth = get_local_size(0);

    uint localIndex = get_local_id(0) + get_local_id(1) * groupWidth * 2;
    const uint localLoadIndex = localIndex + groupWidth;
    const uint localCopyIndex = localIndex;

    const int bufferSize = bufferWidth * get_global_size(1);

    int srcLoadIndex = get_global_id(1) * bufferWidth +
                       get_group_id(0) * groupWidth * iterationsCount +
                       get_local_id(0) - groupWidth / 2;
    int dstStoreIndex = get_global_id(1) * bufferWidth +
                        get_group_id(0) * groupWidth * iterationsCount +
                        get_local_id(0);

    localBuffer[localLoadIndex] = srcBuffer[(srcLoadIndex + bufferSize) % bufferSize];
    srcLoadIndex += groupWidth;
    barrier(CLK_LOCAL_MEM_FENCE);

    for (uint iteration = 0; iteration < iterationsCount; iteration++)
    {
        localBuffer[localCopyIndex] = localBuffer[localLoadIndex];
        barrier(CLK_LOCAL_MEM_FENCE);

        localBuffer[localLoadIndex] = srcBuffer[(srcLoadIndex + bufferSize) % bufferSize];
        barrier(CLK_LOCAL_MEM_FENCE);

        int dataIndex = localIndex;
        float4 data = convert_float4(localBuffer[dataIndex]);
        float4 dataNext;

        float4 result = 0.0f;
        int factorIndex = 0;
        float4 factor;

        for (uint x = 0; x < iterationsInnerCount; x++)
        {
            dataIndex++;
            dataNext = convert_float4(localBuffer[dataIndex]);
            factor = factorValues[factorIndex++];
            result = mad(data, factor.s0, result);
            result = mad((float4)(data.yzw, dataNext.x), factor.s1, result);
            result = mad((float4)(data.zw, dataNext.xy), factor.s2, result);
            result = mad((float4)(data.w, dataNext.xyz), factor.s3, result);
            data = dataNext;

            dataIndex++;
            dataNext = convert_float4(localBuffer[dataIndex]);
            factor = factorValues[factorIndex++];
            result = mad(data, factor.s0, result);
            result = mad((float4)(data.yzw, dataNext.x), factor.s1, result);
            result = mad((float4)(data.zw, dataNext.xy), factor.s2, result);
            result = mad((float4)(data.w, dataNext.xyz), factor.s3, result);
            data = dataNext;
        }

        factor = factorValues[factorIndex];
        result = mad(data, factor.s0, result);

        dstBuffer[dstStoreIndex] = convert_uchar4_sat(result);

        srcLoadIndex += groupWidth;
        dstStoreIndex += groupWidth;
    }
}
