modules/core/src/opencl/gemm.cl

   1 // This file is part of OpenCV project.
   2 // It is subject to the license terms in the LICENSE file found in the top-level directory
   3 // of this distribution and at http://opencv.org/license.html.
   4
   5 // Copyright (C) 2014, Itseez, Inc., all rights reserved.
   6 // Third party copyrights are property of their respective owners.
   7
   8 #ifdef DOUBLE_SUPPORT
   9 #ifdef cl_amd_fp64
  10 #pragma OPENCL EXTENSION cl_amd_fp64:enable
  11 #elif defined (cl_khr_fp64)
  12 #pragma OPENCL EXTENSION cl_khr_fp64:enable
  13 #endif
  14 #endif
  15
  16 #define TSIZE  (int)sizeof(T)
  17 #define WTSIZE (int)sizeof(WT)
  18
  19 #define IND_A mad24(y, A_step, A_offset)
  20 #define IND_B mad24(x, WTSIZE, B_offset)
  21 #define STEP_B B_step / WTSIZE
  22
  23 #if cn==2
  24 #if kercn==2
  25 #define MUL(i, a, b)\
  26     {\
  27     sum.x += fma(a.x, b.x, - a.y * b.y);\
  28     sum.y += fma(a.x, b.y, a.y * b.x);\
  29     }
  30 #else
  31 #define MUL(i, a, b)\
  32     {\
  33     sum.x += fma(a.x, b.x, - a.y * b.y);\
  34     sum.y += fma(a.x, b.y, a.y * b.x);\
  35     sum.z += fma(a.x, b.z, - a.y * b.w);\
  36     sum.w += fma(a.x, b.w, a.y * b.z);\
  37     }
  38 #endif
  39 #else
  40 #define MUL(i, a, b) sum = fma(a, b, sum);
  41 #endif
  42
  43
  44 __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset,
  45                    __global const uchar * B_ptr, int B_step, int B_offset,
  46                    __global uchar * D_ptr, int D_step, int D_offset, int D_rows, int D_cols,
  47                    int n, T1 alpha, T1 beta)
  48 {
  49     int x = get_global_id(0);
  50     int y = get_global_id(1);
  51
  52     int lidx = get_local_id(0);
  53     int lidy = get_local_id(1);
  54
  55     __global const T* A = (__global const T*)(A_ptr + IND_A);
  56     __global const WT* B = (__global const WT*)(B_ptr + IND_B);
  57
  58     WT sum = (WT)(0);
  59
  60 #if LOCAL_SIZE == 1
  61
  62     if (x < D_cols && y < D_rows)
  63     {
  64         for (int i = 0; i < n; ++i)
  65             MUL(i, A[i], B[i*STEP_B]);
  66 #else
  67
  68     __local T  a_local[LOCAL_SIZE*LOCAL_SIZE];
  69     __local WT b_local[LOCAL_SIZE*LOCAL_SIZE];
  70
  71     int reps;
  72 #if NO_MULT
  73     reps = (n + LOCAL_SIZE-1)/LOCAL_SIZE;
  74 #else
  75     reps = n/LOCAL_SIZE;
  76 #endif
  77
  78     for (int p = 0; p < reps; ++p)
  79     {
  80         if (p * LOCAL_SIZE + lidx < n && y < D_rows)
  81             a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)];
  82         if (p * LOCAL_SIZE + lidy < n && x < D_cols)
  83             b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B];
  84
  85         barrier(CLK_LOCAL_MEM_FENCE);
  86
  87         if (x < D_cols && y < D_rows)
  88         {
  89             for (int i = 0; i < LOCAL_SIZE
  90 #if NO_MULT
  91                 && p * LOCAL_SIZE + i < n
  92 #endif
  93                 ; ++i)
  94                 MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]);
  95         }
  96
  97         barrier(CLK_LOCAL_MEM_FENCE);
  98     }
  99
 100     if (x < D_cols && y < D_rows)
 101     {
 102 #endif
 103         __global WT* D = (__global WT*)(D_ptr + mad24(y, D_step, mad24(x, WTSIZE, D_offset)));
 104 #if HAVE_C
 105         D[0] = mad(alpha, sum, D[0]*beta);
 106 #else
 107         D[0] = alpha * sum;
 108 #endif
 109     }
 110 }