use local memory

modules/core/src/matmul.cpp

@@ -782,6 +782,7 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
 {
     int depth = matA.depth(), cn = matA.channels();
     int type = CV_MAKETYPE(depth, cn);
+    const int block_size = 16;
 
     CV_Assert( type == matB.type() && (type == CV_32FC1 || type == CV_64FC1 || type == CV_32FC2 || type == CV_64FC2) );
 
@@ -807,8 +808,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
     CV_Assert( matB.type() == type && (!haveC || matC.type() == type) );
     CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) );
 
-    String opts = format("-D T=%s -D T1=%s -D cn=%d %s %s",
-                          ocl::typeToStr(type), ocl::typeToStr(depth), cn,
+    String opts = format("-D T=%s -D T1=%s -D cn=%d -D LOCAL_SIZE=%d %s %s",
+                          ocl::typeToStr(type), ocl::typeToStr(depth), cn, block_size,
                           haveC ? "-D HAVE_C" : "",
                           doubleSupport ? " -D DOUBLE_SUPPORT" : "");
 
@@ -843,7 +844,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
                sizeA.width, (float)alpha, (float)beta);
 
     size_t globalsize[2] = { sizeD.width, sizeD.height};
-    return k.run(2, globalsize, NULL, false);
+    size_t localsize[2] = { block_size, block_size};
+    return k.run(2, globalsize, localsize, false);
 }
 
 #endif

modules/core/src/opencl/gemm.cl

@@ -28,7 +28,7 @@
     sum.y += fma(a.x, b.y, a.y * b.x);\
     }
 #else
-#define MUL(i, a, b) sum += a * b
+#define MUL(i, a, b) sum = fma(a, b, sum);
 #endif
 
 
@@ -40,16 +40,34 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset,
     int x = get_global_id(0);
     int y = get_global_id(1);
 
-    if (x < D_cols && y < D_rows)
+    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 T* B = (__global const T*)(B_ptr + IND_B);
+
+    T sum = (T)(0);
+    __local T a_local[LOCAL_SIZE*LOCAL_SIZE];
+    __local T b_local[LOCAL_SIZE*LOCAL_SIZE];
+
+    for (int p = 0; p < (n+LOCAL_SIZE-1)/LOCAL_SIZE; ++p)
     {
-        __global const T* A = (__global const T*)(A_ptr + IND_A);
-        __global const T* B = (__global const T*)(B_ptr + IND_B);
+        a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)];
+        b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B];
 
-        T sum = (T)(0);
+        barrier(CLK_LOCAL_MEM_FENCE);
 
-        for (int i = 0; i < n; ++i)
-            MUL(i, A[i*STEP_A], B[i*STEP_B]);
+        if (x < D_cols && y < D_rows)
+        {
+            for (int i = 0; i < LOCAL_SIZE && p * LOCAL_SIZE + i < n; ++i)
+                MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]);
+        }
 
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+
+    if (x < D_cols && y < D_rows)
+    {
         __global T* D = (__global T*)(D_ptr + mad24(y, D_step, mad24(x, TSIZE, D_offset)));
 #if HAVE_C
         D[0] = mad(alpha, sum, D[0]*beta);

modules/ts/src/ocl_test.cpp

@@ -48,7 +48,7 @@ namespace ocl {
 
 using namespace cv;
 
-int test_loop_times = 1; // TODO Read from command line / environment
+int test_loop_times = 10; // TODO Read from command line / environment
 
 #ifdef HAVE_OPENCL