From d78bc3c3219bd52a9ae99177e300267ff49a3269 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Fri, 22 Aug 2014 15:05:29 +0400 Subject: [PATCH 1/5] naive implementation --- modules/core/src/matmul.cpp | 82 ++++++++++++++++++++++++++++++++- modules/core/src/opencl/gemm.cl | 60 ++++++++++++++++++++++++ 2 files changed, 140 insertions(+), 2 deletions(-) create mode 100644 modules/core/src/opencl/gemm.cl diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index e3e1720ef0..325d15149e 100644 --- a/modules/core/src/matmul.cpp +++ b/modules/core/src/matmul.cpp @@ -693,7 +693,7 @@ static void GEMMStore_64fc( const Complexd* c_data, size_t c_step, #ifdef HAVE_CLAMDBLAS -static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, +static bool ocl_gemm_amdblas( InputArray matA, InputArray matB, double alpha, InputArray matC, double beta, OutputArray matD, int flags ) { int type = matA.type(), esz = CV_ELEM_SIZE(type); @@ -775,6 +775,79 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, #endif +#ifdef HAVE_OPENCL + +static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, + InputArray matC, double beta, OutputArray matD, int flags ) +{ + int depth = matA.depth(), cn = matA.channels(); + int type = CV_MAKETYPE(depth, cn); + + CV_Assert( type == matB.type() && (type == CV_32FC1 || type == CV_64FC1 || type == CV_32FC2 || type == CV_64FC2) ); + + const ocl::Device & dev = ocl::Device::getDefault(); + bool doubleSupport = dev.doubleFPConfig() > 0; + + if ((!doubleSupport && depth == CV_64F)) + return false; + + bool haveC = matC.kind() != cv::_InputArray::NONE; + Size sizeA = matA.size(), sizeB = matB.size(), sizeC = haveC ? matC.size() : Size(0, 0); + bool atrans = (flags & GEMM_1_T) != 0, btrans = (flags & GEMM_2_T) != 0, ctrans = (flags & GEMM_3_T) != 0; + + if (atrans) + sizeA = Size(sizeA.height, sizeA.width); + if (btrans) + sizeB = Size(sizeB.height, sizeB.width); + if (haveC && ctrans) + sizeC = Size(sizeC.height, sizeC.width); + + Size sizeD(sizeB.width, sizeA.height); + + 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, + haveC ? "-D HAVE_C" : "", + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); + + ocl::Kernel k("gemm", cv::ocl::core::gemm_oclsrc, opts); + if (k.empty()) + return false; + + matD.create(sizeD, type); + + UMat A = matA.getUMat(), B = matB.getUMat(), D = matD.getUMat(); + + if (atrans) + A = A.t(); + + if (btrans) + B = B.t(); + + if (haveC) + ctrans ? transpose(matC, D) : matC.copyTo(D); + else + D.setTo(Scalar::all(0)); + + if (depth == CV_64F) + k.args(ocl::KernelArg::ReadOnlyNoSize(A), + ocl::KernelArg::ReadOnlyNoSize(B), + ocl::KernelArg::ReadWrite(D), + sizeA.width, alpha, beta); + else + k.args(ocl::KernelArg::ReadOnlyNoSize(A), + ocl::KernelArg::ReadOnlyNoSize(B), + ocl::KernelArg::ReadWrite(D), + sizeA.width, (float)alpha, (float)beta); + + size_t globalsize[2] = { sizeD.width, sizeD.height}; + return k.run(2, globalsize, NULL, false); +} + +#endif + } void cv::gemm( InputArray matA, InputArray matB, double alpha, @@ -783,7 +856,12 @@ void cv::gemm( InputArray matA, InputArray matB, double alpha, #ifdef HAVE_CLAMDBLAS CV_OCL_RUN(ocl::haveAmdBlas() && matA.dims() <= 2 && matB.dims() <= 2 && matC.dims() <= 2 && _matD.isUMat() && matA.cols() > 20 && matA.rows() > 20 && matB.cols() > 20, // since it works incorrect for small sizes - ocl_gemm(matA, matB, alpha, matC, beta, _matD, flags)) + ocl_gemm_amdblas(matA, matB, alpha, matC, beta, _matD, flags)) +#endif + +#ifdef HAVE_OPENCL + CV_OCL_RUN(_matD.isUMat() && matA.dims() <= 2 && matB.dims() <= 2 && matC.dims() <= 2, + ocl_gemm(matA, matB, alpha, matC, beta, _matD, flags)) #endif const int block_lin_size = 128; diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl new file mode 100644 index 0000000000..10e0eea6f3 --- /dev/null +++ b/modules/core/src/opencl/gemm.cl @@ -0,0 +1,60 @@ +// 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 IND_A mad24(y, A_step, A_offset) +#define STEP_A 1 + +#define IND_B mad24(x, TSIZE, B_offset) +#define STEP_B B_step / TSIZE + +#if cn==2 +#define MUL(i, 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(i, a, b) sum += a * b +#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); + + if (x < D_cols && y < D_rows) + { + __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); + + for (int i = 0; i < n; ++i) + MUL(i, A[i*STEP_A], B[i*STEP_B]); + + __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); +#else + D[0] = alpha * sum; +#endif + } +} From 2d89df18041713c14962da572cc2948123a02efd Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Wed, 27 Aug 2014 14:58:01 +0400 Subject: [PATCH 2/5] use local memory --- modules/core/src/matmul.cpp | 8 +++++--- modules/core/src/opencl/gemm.cl | 36 ++++++++++++++++++++++++--------- modules/ts/src/ocl_test.cpp | 2 +- 3 files changed, 33 insertions(+), 13 deletions(-) diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index 325d15149e..1db4346efe 100644 --- a/modules/core/src/matmul.cpp +++ b/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 diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index 10e0eea6f3..b2437de7e1 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/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); + 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) + { + 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]; + + barrier(CLK_LOCAL_MEM_FENCE); + + 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 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); - - for (int i = 0; i < n; ++i) - MUL(i, A[i*STEP_A], B[i*STEP_B]); - __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); diff --git a/modules/ts/src/ocl_test.cpp b/modules/ts/src/ocl_test.cpp index d429d4bc8c..3a3c08cf0f 100644 --- a/modules/ts/src/ocl_test.cpp +++ b/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 From c5a2879ce04582811070c2fd29e364abb9d5dc0a Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Thu, 4 Sep 2014 12:36:23 +0400 Subject: [PATCH 3/5] use vectors --- modules/core/src/matmul.cpp | 40 ++++++++++++--------- modules/core/src/opencl/gemm.cl | 62 +++++++++++++++++++++++++-------- modules/ts/src/ocl_test.cpp | 2 +- 3 files changed, 71 insertions(+), 33 deletions(-) diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index 1db4346efe..6d2adc8690 100644 --- a/modules/core/src/matmul.cpp +++ b/modules/core/src/matmul.cpp @@ -782,7 +782,6 @@ 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) ); @@ -808,14 +807,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 -D LOCAL_SIZE=%d %s %s", - ocl::typeToStr(type), ocl::typeToStr(depth), cn, block_size, - haveC ? "-D HAVE_C" : "", - doubleSupport ? " -D DOUBLE_SUPPORT" : ""); - - ocl::Kernel k("gemm", cv::ocl::core::gemm_oclsrc, opts); - if (k.empty()) - return false; + int max_wg_size = (int)dev.maxWorkGroupSize(); + int block_size = (max_wg_size / (32*cn) < 32) ? (max_wg_size / (16*cn) < 16) ? (max_wg_size / (8*cn) < 8) ? 1 : 8 : 16 : 32; matD.create(sizeD, type); @@ -832,24 +825,37 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, else D.setTo(Scalar::all(0)); + int vectorWidths[] = { 4, 4, 2, 2, 1, 4, cn, -1 }; + + int kercn = ocl::checkOptimalVectorWidth(vectorWidths, B, D); + + String opts = format("-D T=%s -D T1=%s -D WT=%s -D cn=%d -D kercn=%d -D LOCAL_SIZE=%d %s %s %s", + ocl::typeToStr(type), ocl::typeToStr(depth), ocl::typeToStr(CV_MAKETYPE(depth, kercn)), + cn, kercn, block_size, + (sizeA.width % block_size !=0) ? "-D NO_MULT" : "", + haveC ? "-D HAVE_C" : "", + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); + + ocl::Kernel k("gemm", cv::ocl::core::gemm_oclsrc, opts); + if (k.empty()) + return false; + if (depth == CV_64F) k.args(ocl::KernelArg::ReadOnlyNoSize(A), - ocl::KernelArg::ReadOnlyNoSize(B), - ocl::KernelArg::ReadWrite(D), + ocl::KernelArg::ReadOnlyNoSize(B, cn, kercn), + ocl::KernelArg::ReadWrite(D, cn, kercn), sizeA.width, alpha, beta); else k.args(ocl::KernelArg::ReadOnlyNoSize(A), - ocl::KernelArg::ReadOnlyNoSize(B), - ocl::KernelArg::ReadWrite(D), + ocl::KernelArg::ReadOnlyNoSize(B, cn, kercn), + ocl::KernelArg::ReadWrite(D, cn, kercn), sizeA.width, (float)alpha, (float)beta); - size_t globalsize[2] = { sizeD.width, sizeD.height}; + size_t globalsize[2] = { sizeD.width * cn / kercn, sizeD.height}; size_t localsize[2] = { block_size, block_size}; - return k.run(2, globalsize, localsize, false); + return k.run(2, globalsize, block_size!=1 ? localsize : NULL, false); } - #endif - } void cv::gemm( InputArray matA, InputArray matB, double alpha, diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index b2437de7e1..ddd18adaf2 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/modules/core/src/opencl/gemm.cl @@ -13,21 +13,30 @@ #endif #endif -#define TSIZE (int)sizeof(T) +#define TSIZE (int)sizeof(T) +#define WTSIZE (int)sizeof(WT) #define IND_A mad24(y, A_step, A_offset) -#define STEP_A 1 - -#define IND_B mad24(x, TSIZE, B_offset) -#define STEP_B B_step / TSIZE +#define IND_B mad24(x, WTSIZE, B_offset) +#define STEP_B B_step / WTSIZE #if cn==2 +#if kercn==2 #define MUL(i, 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(i, 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(i, a, b) sum = fma(a, b, sum); #endif @@ -44,22 +53,44 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, 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); + __global const WT* B = (__global const WT*)(B_ptr + IND_B); - T sum = (T)(0); - __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; - __local T b_local[LOCAL_SIZE*LOCAL_SIZE]; + WT sum = (WT)(0); - for (int p = 0; p < (n+LOCAL_SIZE-1)/LOCAL_SIZE; ++p) +#if LOCAL_SIZE == 1 + + if (x < D_cols && y < D_rows) { - 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]; + for (int i = 0; i < n; ++i) + MUL(i, A[i], B[i*STEP_B]); +#else + + __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; + __local WT b_local[LOCAL_SIZE*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, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)]; + if (p * LOCAL_SIZE + lidy < n && x < D_cols) + b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B]; barrier(CLK_LOCAL_MEM_FENCE); if (x < D_cols && y < D_rows) { - for (int i = 0; i < LOCAL_SIZE && p * LOCAL_SIZE + i < n; ++i) + for (int i = 0; i < LOCAL_SIZE +#if NO_MULT + && p * LOCAL_SIZE + i < n +#endif + ; ++i) MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); } @@ -68,11 +99,12 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, if (x < D_cols && y < D_rows) { - __global T* D = (__global T*)(D_ptr + mad24(y, D_step, mad24(x, TSIZE, D_offset))); +#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 } -} +} \ No newline at end of file diff --git a/modules/ts/src/ocl_test.cpp b/modules/ts/src/ocl_test.cpp index 3a3c08cf0f..d429d4bc8c 100644 --- a/modules/ts/src/ocl_test.cpp +++ b/modules/ts/src/ocl_test.cpp @@ -48,7 +48,7 @@ namespace ocl { using namespace cv; -int test_loop_times = 10; // TODO Read from command line / environment +int test_loop_times = 1; // TODO Read from command line / environment #ifdef HAVE_OPENCL From 65b8a1cb3756ba908440efaa4b18eed4037e25c6 Mon Sep 17 00:00:00 2001 From: ElenaGvozdeva Date: Thu, 16 Oct 2014 10:24:44 +0300 Subject: [PATCH 4/5] Some small fixes --- modules/core/src/matmul.cpp | 7 ++----- modules/core/src/opencl/gemm.cl | 18 +++++++++--------- 2 files changed, 11 insertions(+), 14 deletions(-) diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index 6d2adc8690..fec2ecb8b8 100644 --- a/modules/core/src/matmul.cpp +++ b/modules/core/src/matmul.cpp @@ -788,7 +788,7 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, const ocl::Device & dev = ocl::Device::getDefault(); bool doubleSupport = dev.doubleFPConfig() > 0; - if ((!doubleSupport && depth == CV_64F)) + if (!doubleSupport && depth == CV_64F) return false; bool haveC = matC.kind() != cv::_InputArray::NONE; @@ -804,7 +804,7 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, Size sizeD(sizeB.width, sizeA.height); - CV_Assert( matB.type() == type && (!haveC || matC.type() == type) ); + CV_Assert( !haveC || matC.type() == type ); CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) ); int max_wg_size = (int)dev.maxWorkGroupSize(); @@ -822,11 +822,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha, if (haveC) ctrans ? transpose(matC, D) : matC.copyTo(D); - else - D.setTo(Scalar::all(0)); int vectorWidths[] = { 4, 4, 2, 2, 1, 4, cn, -1 }; - int kercn = ocl::checkOptimalVectorWidth(vectorWidths, B, D); String opts = format("-D T=%s -D T1=%s -D WT=%s -D cn=%d -D kercn=%d -D LOCAL_SIZE=%d %s %s %s", diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index ddd18adaf2..0961628a49 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/modules/core/src/opencl/gemm.cl @@ -22,13 +22,13 @@ #if cn==2 #if kercn==2 -#define MUL(i, a, b)\ +#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(i, a, b)\ +#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);\ @@ -37,7 +37,7 @@ } #endif #else -#define MUL(i, a, b) sum = fma(a, b, sum); +#define MUL(a, b) sum = fma(a, b, sum); #endif @@ -62,7 +62,7 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, if (x < D_cols && y < D_rows) { for (int i = 0; i < n; ++i) - MUL(i, A[i], B[i*STEP_B]); + MUL(A[i], B[i*STEP_B]); #else __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; @@ -86,14 +86,14 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, if (x < D_cols && y < D_rows) { - for (int i = 0; i < LOCAL_SIZE #if NO_MULT - && p * LOCAL_SIZE + i < n + 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 - ; ++i) - MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); + MUL(a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); } - barrier(CLK_LOCAL_MEM_FENCE); } From d88fdd0378c280763172652931b8d06abe3ebbb5 Mon Sep 17 00:00:00 2001 From: ElenaGvozdeva Date: Tue, 28 Oct 2014 15:18:31 +0300 Subject: [PATCH 5/5] use LOCAL_SIZE+1 --- modules/core/src/opencl/gemm.cl | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index 0961628a49..fc050547be 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/modules/core/src/opencl/gemm.cl @@ -20,6 +20,8 @@ #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)\ @@ -65,8 +67,8 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, MUL(A[i], B[i*STEP_B]); #else - __local T a_local[LOCAL_SIZE*LOCAL_SIZE]; - __local WT b_local[LOCAL_SIZE*LOCAL_SIZE]; + __local T a_local[LOCAL_SIZE_ODD*LOCAL_SIZE]; + __local WT b_local[LOCAL_SIZE_ODD*LOCAL_SIZE]; int reps; #if NO_MULT @@ -78,9 +80,9 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, for (int p = 0; p < reps; ++p) { if (p * LOCAL_SIZE + lidx < n && y < D_rows) - a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)]; + 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, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B]; + b_local[mad24(lidy, LOCAL_SIZE_ODD, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B]; barrier(CLK_LOCAL_MEM_FENCE); @@ -92,7 +94,7 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset, #else for (int i = 0; i < LOCAL_SIZE; ++i) #endif - MUL(a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); + MUL(a_local[mad24(lidy, LOCAL_SIZE_ODD, i)], b_local[mad24(i, LOCAL_SIZE_ODD, lidx)]); } barrier(CLK_LOCAL_MEM_FENCE); }