root/modules/core/src/opencl/gemm.cl

DEFINITIONS

1. gemm

// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.

// Copyright (C) 2014, Itseez, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.

#ifdef DOUBLE_SUPPORT
#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64:enable
#elif defined (cl_khr_fp64)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#endif

#define TSIZE  (int)sizeof(T)
#define WTSIZE (int)sizeof(WT)

#define IND_A mad24(y, A_step, A_offset)
#define IND_B mad24(x, WTSIZE, B_offset)
#define STEP_B B_step / WTSIZE

#define LOCAL_SIZE_ODD (LOCAL_SIZE + 1)

#if cn==2
#if kercn==2
#define MUL(a, b)\
    {\
    sum.x += fma(a.x, b.x, - a.y * b.y);\
    sum.y += fma(a.x, b.y, a.y * b.x);\
    }
#else
#define MUL(a, b)\
    {\
    sum.x += fma(a.x, b.x, - a.y * b.y);\
    sum.y += fma(a.x, b.y, a.y * b.x);\
    sum.z += fma(a.x, b.z, - a.y * b.w);\
    sum.w += fma(a.x, b.w, a.y * b.z);\
    }
#endif
#else
#define MUL(a, b) sum = fma(a, b, sum);
#endif


__kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset,
                   __global const uchar * B_ptr, int B_step, int B_offset,
                   __global uchar * D_ptr, int D_step, int D_offset, int D_rows, int D_cols,
                   int n, T1 alpha, T1 beta)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    int lidx = get_local_id(0);
    int lidy = get_local_id(1);

    __global const T* A = (__global const T*)(A_ptr + IND_A);
    __global const WT* B = (__global const WT*)(B_ptr + IND_B);

    WT sum = (WT)(0);

#if LOCAL_SIZE == 1

    if (x < D_cols && y < D_rows)
    {
        for (int i = 0; i < n; ++i)
            MUL(A[i], B[i*STEP_B]);
#else

    __local T  a_local[LOCAL_SIZE_ODD*LOCAL_SIZE];
    __local WT b_local[LOCAL_SIZE_ODD*LOCAL_SIZE];

    int reps;
#if NO_MULT
    reps = (n + LOCAL_SIZE-1)/LOCAL_SIZE;
#else
    reps = n/LOCAL_SIZE;
#endif

    for (int p = 0; p < reps; ++p)
    {
        if (p * LOCAL_SIZE + lidx < n && y < D_rows)
            a_local[mad24(lidy, LOCAL_SIZE_ODD, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)];
        if (p * LOCAL_SIZE + lidy < n && x < D_cols)
            b_local[mad24(lidy, LOCAL_SIZE_ODD, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B];

        barrier(CLK_LOCAL_MEM_FENCE);

        if (x < D_cols && y < D_rows)
        {
#if NO_MULT
            int ie = min(LOCAL_SIZE, n - p * LOCAL_SIZE);
            for (int i = 0; i < ie; ++i)
#else
            for (int i = 0; i < LOCAL_SIZE; ++i)
#endif
                MUL(a_local[mad24(lidy, LOCAL_SIZE_ODD, i)], b_local[mad24(i, LOCAL_SIZE_ODD, lidx)]);
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if (x < D_cols && y < D_rows)
    {
#endif
        __global WT* D = (__global WT*)(D_ptr + mad24(y, D_step, mad24(x, WTSIZE, D_offset)));
#if HAVE_C
        D[0] = mad(alpha, sum, D[0]*beta);
#else
        D[0] = alpha * sum;
#endif
    }