Home | History | Annotate | Download | only in opencl 
      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 #define LOCAL_SIZE_ODD (LOCAL_SIZE + 1)
     24 
     25 #if cn==2
     26 #if kercn==2
     27 #define MUL(a, b)\
     28     {\
     29     sum.x += fma(a.x, b.x, - a.y * b.y);\
     30     sum.y += fma(a.x, b.y, a.y * b.x);\
     31     }
     32 #else
     33 #define MUL(a, b)\
     34     {\
     35     sum.x += fma(a.x, b.x, - a.y * b.y);\
     36     sum.y += fma(a.x, b.y, a.y * b.x);\
     37     sum.z += fma(a.x, b.z, - a.y * b.w);\
     38     sum.w += fma(a.x, b.w, a.y * b.z);\
     39     }
     40 #endif
     41 #else
     42 #define MUL(a, b) sum = fma(a, b, sum);
     43 #endif
     44 
     45 
     46 __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset,
     47                    __global const uchar * B_ptr, int B_step, int B_offset,
     48                    __global uchar * D_ptr, int D_step, int D_offset, int D_rows, int D_cols,
     49                    int n, T1 alpha, T1 beta)
     50 {
     51     int x = get_global_id(0);
     52     int y = get_global_id(1);
     53 
     54     int lidx = get_local_id(0);
     55     int lidy = get_local_id(1);
     56 
     57     __global const T* A = (__global const T*)(A_ptr + IND_A);
     58     __global const WT* B = (__global const WT*)(B_ptr + IND_B);
     59 
     60     WT sum = (WT)(0);
     61 
     62 #if LOCAL_SIZE == 1
     63 
     64     if (x < D_cols && y < D_rows)
     65     {
     66         for (int i = 0; i < n; ++i)
     67             MUL(A[i], B[i*STEP_B]);
     68 #else
     69 
     70     __local T  a_local[LOCAL_SIZE_ODD*LOCAL_SIZE];
     71     __local WT b_local[LOCAL_SIZE_ODD*LOCAL_SIZE];
     72 
     73     int reps;
     74 #if NO_MULT
     75     reps = (n + LOCAL_SIZE-1)/LOCAL_SIZE;
     76 #else
     77     reps = n/LOCAL_SIZE;
     78 #endif
     79 
     80     for (int p = 0; p < reps; ++p)
     81     {
     82         if (p * LOCAL_SIZE + lidx < n && y < D_rows)
     83             a_local[mad24(lidy, LOCAL_SIZE_ODD, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)];
     84         if (p * LOCAL_SIZE + lidy < n && x < D_cols)
     85             b_local[mad24(lidy, LOCAL_SIZE_ODD, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B];
     86 
     87         barrier(CLK_LOCAL_MEM_FENCE);
     88 
     89         if (x < D_cols && y < D_rows)
     90         {
     91 #if NO_MULT
     92             int ie = min(LOCAL_SIZE, n - p * LOCAL_SIZE);
     93             for (int i = 0; i < ie; ++i)
     94 #else
     95             for (int i = 0; i < LOCAL_SIZE; ++i)
     96 #endif
     97                 MUL(a_local[mad24(lidy, LOCAL_SIZE_ODD, i)], b_local[mad24(i, LOCAL_SIZE_ODD, lidx)]);
     98         }
     99         barrier(CLK_LOCAL_MEM_FENCE);
    100     }
    101 
    102     if (x < D_cols && y < D_rows)
    103     {
    104 #endif
    105         __global WT* D = (__global WT*)(D_ptr + mad24(y, D_step, mad24(x, WTSIZE, D_offset)));
    106 #if HAVE_C
    107         D[0] = mad(alpha, sum, D[0]*beta);
    108 #else
    109         D[0] = alpha * sum;
    110 #endif
    111     }
    112 }