Merge remote-tracking branch 'upstream/3.4' into merge-3.4

modules/core/src/intel_gpu_gemm.inl.hpp

@@ -24,11 +24,6 @@
 
 #ifdef HAVE_OPENCL
 
-#include <sstream>
-#include "opencl_kernels_core.hpp"
-#include "opencv2/core/opencl/runtime/opencl_clblas.hpp"
-#include "opencv2/core/opencl/runtime/opencl_core.hpp"
-
 namespace cv
 {
 
@@ -37,52 +32,75 @@ static bool intel_gpu_gemm(
     UMat B, Size sizeB,
     UMat D, Size sizeD,
     double alpha, double beta,
-    bool atrans, bool btrans)
+    bool atrans, bool btrans,
+    bool& isPropagatedC2D
+)
 {
     CV_UNUSED(sizeB);
 
     int M = sizeD.height, N = sizeD.width, K = ((atrans)? sizeA.height : sizeA.width);
 
-    std::string kernelName;
-    bool ret = true;
+    if (M < 4 || N < 4 || K < 4)  // vload4
+        return false;
 
-    size_t lx = 8, ly = 4;
-    size_t dx = 4, dy = 8;
+    CV_LOG_VERBOSE(NULL, 0, "M=" << M << " N=" << N << " K=" << K);
+
+    std::string kernelName;
+
+    unsigned int lx = 8, ly = 4;
+    unsigned int dx = 4, dy = 8;
 
     if(!atrans && !btrans)
     {
-
         if (M % 32 == 0 && N % 32 == 0 && K % 16 == 0)
         {
             kernelName = "intelblas_gemm_buffer_NN_sp";
         }
         else
         {
+            if (M % 2 != 0)
+                return false;
+            // vload4(0, dst_write0) - 4 cols
+            // multiply by lx: 8
+            if (N % (4*8) != 0)
+                return false;
             kernelName = "intelblas_gemm_buffer_NN";
         }
     }
     else if(atrans && !btrans)
     {
+        if (M % 32 != 0)
+            return false;
+        if (N % 32 != 0)
+            return false;
         kernelName = "intelblas_gemm_buffer_TN";
     }
     else if(!atrans && btrans)
     {
+        if (K % 4 != 0)
+            return false;
         kernelName = "intelblas_gemm_buffer_NT";
         ly = 16;
         dx = 1;
     }
     else
     {
+        if (M % 32 != 0)
+            return false;
+        if (N % 32 != 0)
+            return false;
+        if (K % 16 != 0)
+            return false;
         kernelName = "intelblas_gemm_buffer_TT";
     }
 
-    const size_t gx = (size_t)(N + dx - 1) / dx;
-    const size_t gy = (size_t)(M + dy - 1) / dy;
+    CV_LOG_DEBUG(NULL, "kernel: " << kernelName << "  (M=" << M << " N=" << N << " K=" << K << ")");
+
+    const size_t gx = divUp((size_t)N, dx);
+    const size_t gy = divUp((size_t)M, dy);
 
     size_t local[] = {lx, ly, 1};
-    size_t global[] = {(gx + lx - 1) / lx * lx, (gy + ly - 1) / ly * ly, 1};
-
-    int stride = (M * N < 1024 * 1024) ? 10000000 : 256;
+    size_t global[] = {roundUp(gx, lx), roundUp(gy, ly), 1};
 
     ocl::Queue q;
     String errmsg;
@@ -110,10 +128,13 @@ static bool intel_gpu_gemm(
                (int)(D.step / sizeof(float))
         );
 
-        ret = k.run(2, global, local, false, q);
+        bool ret = k.run(2, global, local, false, q);
+        return ret;
     }
     else
     {
+        int stride = (M * N < 1024 * 1024) ? 10000000 : 256;
+
         for(int start_index = 0; start_index < K; start_index += stride)
         {
              ocl::Kernel k(kernelName.c_str(), program);
@@ -132,12 +153,16 @@ static bool intel_gpu_gemm(
                     (int) start_index,                          // 14 start_index
                     stride);
 
-            ret = k.run(2, global, local, false, q);
-            if (!ret) return ret;
+            bool ret = k.run(2, global, local, false, q);
+            if (!ret)
+            {
+                if (start_index != 0)
+                    isPropagatedC2D = false;  // D array content is changed, need to rewrite
+                return false;
+            }
         }
+        return true;
     }
-
-    return ret;
 }
 
 } // namespace cv

modules/core/src/matmul.dispatch.cpp

@@ -42,6 +42,8 @@
 //M*/
 
 #include "precomp.hpp"
+#include <opencv2/core/utils/logger.hpp>
+
 #include "opencl_kernels_core.hpp"
 #include "opencv2/core/opencl/runtime/opencl_clblas.hpp"
 #include "opencv2/core/opencl/runtime/opencl_core.hpp"
@@ -155,10 +157,12 @@ static bool ocl_gemm_amdblas( InputArray matA, InputArray matB, double alpha,
 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);
+    int type = matA.type();
+    int depth = CV_MAT_DEPTH(type);
+    int cn = CV_MAT_CN(type);
 
-    CV_Assert_N( type == matB.type(), (type == CV_32FC1 || type == CV_64FC1 || type == CV_32FC2 || type == CV_64FC2) );
+    CV_CheckTypeEQ(type, matB.type(), "");
+    CV_CheckType(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;
@@ -170,88 +174,103 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
     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;
 
-    CV_Assert( !haveC || matC.type() == type );
+    if (haveC)
+        CV_CheckTypeEQ(type, matC.type(), "");
+
+    Size sizeD(((btrans) ? sizeB.height : sizeB.width),
+               ((atrans) ? sizeA.width : sizeA.height));
+
+    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);
+
+    CV_CheckEQ(sizeA.width, sizeB.height, "");
+    if (haveC)
+        CV_CheckEQ(sizeC, sizeD, "");
+
+    UMat A = matA.getUMat();
+    UMat B = matB.getUMat();
 
-    Size sizeD(((btrans)? sizeB.height : sizeB.width),
-               ((atrans)? sizeA.width : sizeA.height));
     matD.create(sizeD, type);
+    UMat D = matD.getUMat();
 
-    UMat A = matA.getUMat(), B = matB.getUMat(), D = matD.getUMat();
+    bool isPropagatedC2D = false;  // D content is updated with C / C.t()
 
-
-    if (!dev.intelSubgroupsSupport() || (depth == CV_64F) || cn != 1)
-    {
-        String opts;
-
-        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);
-
-        CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) );
-
-        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;
-
-        if (atrans)
-            A = A.t();
-
-        if (btrans)
-            B = B.t();
-
-        if (haveC)
-            ctrans ? transpose(matC, D) : matC.copyTo(D);
-
-        int vectorWidths[] = { 4, 4, 2, 2, 1, 4, cn, -1 };
-        int kercn = ocl::checkOptimalVectorWidth(vectorWidths, B, D);
-
-        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, cn, kercn),
-                   ocl::KernelArg::ReadWrite(D, cn, kercn),
-                   sizeA.width, alpha, beta);
-        else
-            k.args(ocl::KernelArg::ReadOnlyNoSize(A),
-                   ocl::KernelArg::ReadOnlyNoSize(B, cn, kercn),
-                   ocl::KernelArg::ReadWrite(D, cn, kercn),
-                   sizeA.width, (float)alpha, (float)beta);
-
-        size_t globalsize[2] = { (size_t)sizeD.width * cn / kercn, (size_t)sizeD.height};
-        size_t localsize[2] = { (size_t)block_size, (size_t)block_size};
-
-        return k.run(2, globalsize, block_size!=1 ? localsize : NULL, false);
-    }
-    else
+    if (dev.intelSubgroupsSupport() && (depth == CV_32F) && cn == 1)
     {
         if (haveC && beta != 0.0)
         {
             ctrans ? transpose(matC, D) : matC.copyTo(D);
+            isPropagatedC2D = true;
         }
         else
         {
             beta = 0.0;
         }
 
-        return intel_gpu_gemm(A, sizeA,
-                              B, sizeB,
-                              D, sizeD,
-                              alpha,
-                              beta,
-                              atrans, btrans);
+        bool res = intel_gpu_gemm(A, matA.size(),
+                                  B, matB.size(),
+                                  D, sizeD,
+                                  alpha,
+                                  beta,
+                                  atrans, btrans,
+                                  isPropagatedC2D);
+        if (res)
+            return true;
+        // fallback on generic OpenCL code
     }
+
+    if (sizeD.width < 8 || sizeD.height < 8)
+        return false;
+
+    String opts;
+
+    int wg_size = (int)dev.maxWorkGroupSize();
+    int sizeDmin = std::min(sizeD.width, sizeD.height);
+    wg_size = std::min(wg_size, sizeDmin * sizeDmin);
+    int block_size = (wg_size / (32*cn) < 32) ? (wg_size / (16*cn) < 16) ? (wg_size / (8*cn) < 8) ? 1 : 8 : 16 : 32;
+
+    if (atrans)
+        A = A.t();
+
+    if (btrans)
+        B = B.t();
+
+    if (haveC && !isPropagatedC2D)
+        ctrans ? transpose(matC, D) : matC.copyTo(D);
+
+    int vectorWidths[] = { 4, 4, 2, 2, 1, 4, cn, -1 };
+    int kercn = ocl::checkOptimalVectorWidth(vectorWidths, B, D);
+
+    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, cn, kercn),
+               ocl::KernelArg::ReadWrite(D, cn, kercn),
+               sizeA.width, alpha, beta);
+    else
+        k.args(ocl::KernelArg::ReadOnlyNoSize(A),
+               ocl::KernelArg::ReadOnlyNoSize(B, cn, kercn),
+               ocl::KernelArg::ReadWrite(D, cn, kercn),
+               sizeA.width, (float)alpha, (float)beta);
+
+    size_t globalsize[2] = { (size_t)sizeD.width * cn / kercn, (size_t)sizeD.height};
+    size_t localsize[2] = { (size_t)block_size, (size_t)block_size};
+
+    return k.run(2, globalsize, block_size !=1 ? localsize : NULL, false);
 }
 #endif
 

modules/core/src/ocl.cpp

@@ -76,8 +76,11 @@
 #undef CV__ALLOCATOR_STATS_LOG
 
 #define CV_OPENCL_ALWAYS_SHOW_BUILD_LOG          0
+#define CV_OPENCL_SHOW_BUILD_OPTIONS             0
+#define CV_OPENCL_SHOW_BUILD_KERNELS             0
 
 #define CV_OPENCL_SHOW_RUN_KERNELS               0
+#define CV_OPENCL_SYNC_RUN_KERNELS               0
 #define CV_OPENCL_TRACE_CHECK                    0
 
 #define CV_OPENCL_VALIDATE_BINARY_PROGRAMS       1
@@ -2155,20 +2158,22 @@ static cl_device_id selectOpenCLDevice(const char* configuration = NULL)
         platforms.resize(numPlatforms);
     }
 
-    int selectedPlatform = -1;
     if (platform.length() > 0)
     {
-        for (size_t i = 0; i < platforms.size(); i++)
+        for (std::vector<cl_platform_id>::iterator currentPlatform = platforms.begin(); currentPlatform != platforms.end();)
         {
             std::string name;
-            CV_OCL_DBG_CHECK(getStringInfo(clGetPlatformInfo, platforms[i], CL_PLATFORM_NAME, name));
+            CV_OCL_DBG_CHECK(getStringInfo(clGetPlatformInfo, *currentPlatform, CL_PLATFORM_NAME, name));
             if (name.find(platform) != std::string::npos)
             {
-                selectedPlatform = (int)i;
-                break;
+                ++currentPlatform;
+            }
+            else
+            {
+                currentPlatform = platforms.erase(currentPlatform);
             }
         }
-        if (selectedPlatform == -1)
+        if (platforms.size() == 0)
         {
             CV_LOG_ERROR(NULL, "OpenCL: Can't find OpenCL platform by name: " << platform);
             goto not_found;
@@ -2205,13 +2210,11 @@ static cl_device_id selectOpenCLDevice(const char* configuration = NULL)
             goto not_found;
         }
 
-        std::vector<cl_device_id> devices; // TODO Use clReleaseDevice to cleanup
-        for (int i = selectedPlatform >= 0 ? selectedPlatform : 0;
-                (selectedPlatform >= 0 ? i == selectedPlatform : true) && (i < (int)platforms.size());
-                i++)
+        std::vector<cl_device_id> devices;
+        for (std::vector<cl_platform_id>::iterator currentPlatform = platforms.begin(); currentPlatform != platforms.end(); ++currentPlatform)
         {
             cl_uint count = 0;
-            cl_int status = clGetDeviceIDs(platforms[i], deviceType, 0, NULL, &count);
+            cl_int status = clGetDeviceIDs(*currentPlatform, deviceType, 0, NULL, &count);
             if (!(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND))
             {
                 CV_OCL_DBG_CHECK_RESULT(status, "clGetDeviceIDs get count");
@@ -2220,7 +2223,7 @@ static cl_device_id selectOpenCLDevice(const char* configuration = NULL)
                 continue;
             size_t base = devices.size();
             devices.resize(base + count);
-            status = clGetDeviceIDs(platforms[i], deviceType, count, &devices[base], &count);
+            status = clGetDeviceIDs(*currentPlatform, deviceType, count, &devices[base], &count);
             if (!(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND))
             {
                 CV_OCL_DBG_CHECK_RESULT(status, "clGetDeviceIDs get IDs");
@@ -3679,6 +3682,8 @@ static cv::String dumpValue(size_t sz, const void* p)
 {
     if (!p)
         return "NULL";
+    if (sz == 2)
+        return cv::format("%d / %uu / 0x%04x", *(short*)p, *(unsigned short*)p, *(short*)p);
     if (sz == 4)
         return cv::format("%d / %uu / 0x%08x / %g", *(int*)p, *(int*)p, *(int*)p, *(float*)p);
     if (sz == 8)
@@ -3851,6 +3856,14 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
 }
 
 
+bool Kernel::run_(int dims, size_t _globalsize[], size_t _localsize[],
+                  bool sync, const Queue& q)
+{
+    CV_Assert(p);
+    return p->run(dims, _globalsize, _localsize, sync, NULL, q);
+}
+
+
 static bool isRaiseErrorOnReuseAsyncKernel()
 {
     static bool initialized = false;
@@ -3891,6 +3904,10 @@ bool Kernel::Impl::run(int dims, size_t globalsize[], size_t localsize[],
         return false;  // OpenCV 5.0: raise error
     }
 
+#if CV_OPENCL_SYNC_RUN_KERNELS
+    sync = true;
+#endif
+
     cl_command_queue qq = getQueue(q);
     if (haveTempDstUMats)
         sync = true;
@@ -4338,7 +4355,28 @@ struct Program::Impl
             if (!param_buildExtraOptions.empty())
                 buildflags = joinBuildOptions(buildflags, param_buildExtraOptions);
         }
+#if CV_OPENCL_SHOW_BUILD_OPTIONS
+        CV_LOG_INFO(NULL, "OpenCL program '" << sourceModule_ << "/" << sourceName_ << "' options:" << buildflags);
+#endif
         compile(ctx, src_, errmsg);
+#if CV_OPENCL_SHOW_BUILD_KERNELS
+        if (handle)
+        {
+            size_t retsz = 0;
+            char kernels_buffer[4096] = {0};
+            cl_int result = clGetProgramInfo(handle, CL_PROGRAM_KERNEL_NAMES, sizeof(kernels_buffer), &kernels_buffer[0], &retsz);
+            CV_OCL_DBG_CHECK_RESULT(result, cv::format("clGetProgramInfo(CL_PROGRAM_KERNEL_NAMES: %s/%s)", sourceModule_.c_str(), sourceName_.c_str()).c_str());
+            if (result == CL_SUCCESS && retsz < sizeof(kernels_buffer))
+            {
+                kernels_buffer[retsz] = 0;
+                CV_LOG_INFO(NULL, "OpenCL program '" << sourceModule_ << "/" << sourceName_ << "' kernels: '" << kernels_buffer << "'");
+            }
+            else
+            {
+                CV_LOG_ERROR(NULL, "OpenCL program '" << sourceModule_ << "/" << sourceName_ << "' can't retrieve kernel names!");
+            }
+        }
+#endif
     }
 
     bool compile(const Context& ctx, const ProgramSource::Impl* src_, String& errmsg)
@@ -4570,7 +4608,6 @@ struct Program::Impl
                 CV_LOG_INFO(NULL, result << ": Kernels='" << kernels_buffer << "'");
             }
 #endif
-
         }
         return handle != NULL;
     }

modules/core/src/opencl/intel_gemm.cl

@@ -392,6 +392,15 @@ __kernel void intelblas_gemm_buffer_NN(
 #define TILE_N          8
 #define SLM_BLOCK       512
 
+/*
+   A     K         B.t()   K            D     N
+    -----------       -----------        -----------
+    |         |       |         |        |         |
+  M |         |  x  N |         |  =>  M |         |
+    |         |       |         |        |         |
+    -----------       -----------        -----------
+*/
+
 __attribute__((reqd_work_group_size(8, LWG_HEIGHT, 1)))
 __kernel void intelblas_gemm_buffer_NT(
     const __global float *src0, int off0,
@@ -422,59 +431,79 @@ __kernel void intelblas_gemm_buffer_NT(
     float8 dot06 = 0.f;
     float8 dot07 = 0.f;
 
-    float4 brow0;
-    float4 brow1;
-    float4 brow2;
-    float4 brow3;
-    float4 brow4;
-    float4 brow5;
-    float4 brow6;
-    float4 brow7;
+    const int dst_row = (global_y * TILE_M);
+    __global float *dst_write0 = dst + global_x + dst_row * ldC + offd;
 
-    __global float *dst_write0 = dst + local_x * VEC_SIZE + ( group_x * TILE_N ) + ( group_y * LWG_HEIGHT * TILE_M + local_y * TILE_M) * ldC + offd;
+    const __global float *src0_read00 = src0 + off0;
+    const int a_row_base = global_y * TILE_M;
+    const int a_col_base = local_x * (TILE_K / 8);  // <= TILE_K - 4
 
-    const __global float *src0_read = src0 + local_x * ( TILE_K / 8 ) + ( group_y * LWG_HEIGHT * TILE_M + local_y * TILE_M ) * ldA + off0;
-
-    const __global float *src1_read0 = src1 + ( group_x * TILE_N ) * ldB + off1;
+    const __global float *src1_read00 = src1 + off1;
+    const int b_row_base = (group_x * TILE_N);
+    //const int b_col_base = 0;
 
     __local float slm_brow[8 * SLM_BLOCK];
-    __local float* slm_brow0;
 
     int local_index = mad24(local_y, 8, local_x) * 4;
-    int w;
-    for(int b_tile = 0; b_tile < K; b_tile += SLM_BLOCK) {
+    int w = 0;
+    for (int b_tile = 0; b_tile < K; b_tile += SLM_BLOCK)
+    {
+#define UPDATE_BROW(_row) \
+        { \
+            float4 brow; \
+            int b_row = b_row_base + _row; \
+            int b_col = b_tile + local_index; \
+            if (b_row < N && b_col <= K - 4 /*vload4*/) \
+                brow = vload4(0, src1_read00 + mad24(b_row, ldB, b_col)); \
+            else \
+                brow = (float4)0; \
+            vstore4(brow, 0, slm_brow + mad24(_row, SLM_BLOCK, local_index)); \
+        }
+
         barrier(CLK_LOCAL_MEM_FENCE);
-        vstore4(vload4(0, src1_read0 + mad24(0, ldB, local_index)), 0, slm_brow + mad24(0, SLM_BLOCK, local_index));
-        vstore4(vload4(0, src1_read0 + mad24(1, ldB, local_index)), 0, slm_brow + mad24(1, SLM_BLOCK, local_index));
-        vstore4(vload4(0, src1_read0 + mad24(2, ldB, local_index)), 0, slm_brow + mad24(2, SLM_BLOCK, local_index));
-        vstore4(vload4(0, src1_read0 + mad24(3, ldB, local_index)), 0, slm_brow + mad24(3, SLM_BLOCK, local_index));
-        vstore4(vload4(0, src1_read0 + mad24(4, ldB, local_index)), 0, slm_brow + mad24(4, SLM_BLOCK, local_index));
-        vstore4(vload4(0, src1_read0 + mad24(5, ldB, local_index)), 0, slm_brow + mad24(5, SLM_BLOCK, local_index));
-        vstore4(vload4(0, src1_read0 + mad24(6, ldB, local_index)), 0, slm_brow + mad24(6, SLM_BLOCK, local_index));
-        vstore4(vload4(0, src1_read0 + mad24(7, ldB, local_index)), 0, slm_brow + mad24(7, SLM_BLOCK, local_index));
+        UPDATE_BROW(0);
+        UPDATE_BROW(1);
+        UPDATE_BROW(2);
+        UPDATE_BROW(3);
+        UPDATE_BROW(4);
+        UPDATE_BROW(5);
+        UPDATE_BROW(6);
+        UPDATE_BROW(7);
         barrier(CLK_LOCAL_MEM_FENCE);
+#undef UPDATE_BROW
 
-        slm_brow0 = slm_brow + local_x * (TILE_K / 8);
-        w = b_tile;
-        int end_w = min(b_tile + SLM_BLOCK, K);
-        while( w + TILE_K <= end_w ) {
-            float4 arow;
+        for (int k_tile_offset = 0; k_tile_offset < SLM_BLOCK; k_tile_offset += TILE_K)
+        {
+            int a_col = a_col_base + b_tile + k_tile_offset;
 
-            brow0 = vload4(0, slm_brow0 + 0 * SLM_BLOCK);
-            brow1 = vload4(0, slm_brow0 + 1 * SLM_BLOCK);
-            brow2 = vload4(0, slm_brow0 + 2 * SLM_BLOCK);
-            brow3 = vload4(0, slm_brow0 + 3 * SLM_BLOCK);
-            brow4 = vload4(0, slm_brow0 + 4 * SLM_BLOCK);
-            brow5 = vload4(0, slm_brow0 + 5 * SLM_BLOCK);
-            brow6 = vload4(0, slm_brow0 + 6 * SLM_BLOCK);
-            brow7 = vload4(0, slm_brow0 + 7 * SLM_BLOCK);
+            if (a_col > K - 4 /*vload4*/)
+                break;
 
-#define MM_DOT_PRODUCT(_row,_dot)   \
-            arow = vload4(0, src0_read + _row * ldA);                           \
-            _dot = mad( (float8)(arow.x), (float8)(brow0.x, brow1.x, brow2.x, brow3.x, brow4.x, brow5.x, brow6.x, brow7.x), _dot ); \
-            _dot = mad( (float8)(arow.y), (float8)(brow0.y, brow1.y, brow2.y, brow3.y, brow4.y, brow5.y, brow6.y, brow7.y), _dot ); \
-            _dot = mad( (float8)(arow.z), (float8)(brow0.z, brow1.z, brow2.z, brow3.z, brow4.z, brow5.z, brow6.z, brow7.z), _dot ); \
-            _dot = mad( (float8)(arow.w), (float8)(brow0.w, brow1.w, brow2.w, brow3.w, brow4.w, brow5.w, brow6.w, brow7.w), _dot );
+            int slm_brow_col = a_col_base + k_tile_offset;  // <= SLM_BLOCK - 4
+#define READ_SLM_BROW(_row) \
+            float4 brow##_row = vload4(0, slm_brow + mad24(_row, SLM_BLOCK, slm_brow_col));
+
+            READ_SLM_BROW(0);
+            READ_SLM_BROW(1);
+            READ_SLM_BROW(2);
+            READ_SLM_BROW(3);
+            READ_SLM_BROW(4);
+            READ_SLM_BROW(5);
+            READ_SLM_BROW(6);
+            READ_SLM_BROW(7);
+#undef READ_SLM_BROW
+
+#define MM_DOT_PRODUCT(_row,_dot) \
+            { \
+                int a_row = a_row_base + _row; \
+                if (a_row < M) { \
+                    float4 arow = vload4(0, src0_read00 + mad24(a_row, ldA, a_col)); \
+                    _dot = mad( (float8)(arow.x), (float8)(brow0.x, brow1.x, brow2.x, brow3.x, brow4.x, brow5.x, brow6.x, brow7.x), _dot ); \
+                    _dot = mad( (float8)(arow.y), (float8)(brow0.y, brow1.y, brow2.y, brow3.y, brow4.y, brow5.y, brow6.y, brow7.y), _dot ); \
+                    _dot = mad( (float8)(arow.z), (float8)(brow0.z, brow1.z, brow2.z, brow3.z, brow4.z, brow5.z, brow6.z, brow7.z), _dot ); \
+                    _dot = mad( (float8)(arow.w), (float8)(brow0.w, brow1.w, brow2.w, brow3.w, brow4.w, brow5.w, brow6.w, brow7.w), _dot ); \
+                } \
+            }
 
             MM_DOT_PRODUCT(0,dot00);
             MM_DOT_PRODUCT(1,dot01);
@@ -485,53 +514,7 @@ __kernel void intelblas_gemm_buffer_NT(
             MM_DOT_PRODUCT(6,dot06);
             MM_DOT_PRODUCT(7,dot07);
 #undef MM_DOT_PRODUCT
-
-            src0_read += TILE_K;
-            slm_brow0 += TILE_K;
-            w += TILE_K;
         }
-        src1_read0 += SLM_BLOCK;
-    }
-
-    if(w < K) {
-        float4 arow;
-
-#define READ_BROW(_brow,_row) \
-        _brow = vload4(0, slm_brow0 + _row * SLM_BLOCK); \
-        _brow.x = (mad24(local_x, 4, w) < K) ? _brow.x : 0.0f; \
-        _brow.y = (mad24(local_x, 4, w + 1) < K) ? _brow.y : 0.0f; \
-        _brow.z = (mad24(local_x, 4, w + 2) < K) ? _brow.z : 0.0f; \
-        _brow.w = (mad24(local_x, 4, w + 3) < K) ? _brow.w : 0.0f;
-
-        READ_BROW(brow0,0);
-        READ_BROW(brow1,1);
-        READ_BROW(brow2,2);
-        READ_BROW(brow3,3);
-        READ_BROW(brow4,4);
-        READ_BROW(brow5,5);
-        READ_BROW(brow6,6);
-        READ_BROW(brow7,7);
-
-#define MM_DOT_PRODUCT(_row,_dot)   \
-        arow = vload4(0, src0_read + _row * ldA);  \
-        arow.x = (mad24(local_x, 4, w) < K) ? arow.x : 0.0f; \
-        arow.y = (mad24(local_x, 4, w + 1) < K) ? arow.y : 0.0f; \
-        arow.z = (mad24(local_x, 4, w + 2) < K) ? arow.z : 0.0f; \
-        arow.w = (mad24(local_x, 4, w + 3) < K) ? arow.w : 0.0f; \
-        _dot = mad( (float8)(arow.x), (float8)(brow0.x, brow1.x, brow2.x, brow3.x, brow4.x, brow5.x, brow6.x, brow7.x), _dot ); \
-        _dot = mad( (float8)(arow.y), (float8)(brow0.y, brow1.y, brow2.y, brow3.y, brow4.y, brow5.y, brow6.y, brow7.y), _dot ); \
-        _dot = mad( (float8)(arow.z), (float8)(brow0.z, brow1.z, brow2.z, brow3.z, brow4.z, brow5.z, brow6.z, brow7.z), _dot ); \
-        _dot = mad( (float8)(arow.w), (float8)(brow0.w, brow1.w, brow2.w, brow3.w, brow4.w, brow5.w, brow6.w, brow7.w), _dot );
-
-        MM_DOT_PRODUCT(0,dot00);
-        MM_DOT_PRODUCT(1,dot01);
-        MM_DOT_PRODUCT(2,dot02);
-        MM_DOT_PRODUCT(3,dot03);
-        MM_DOT_PRODUCT(4,dot04);
-        MM_DOT_PRODUCT(5,dot05);
-        MM_DOT_PRODUCT(6,dot06);
-        MM_DOT_PRODUCT(7,dot07);
-#undef MM_DOT_PRODUCT
     }
 
 #define REDUCE(_dot) \
@@ -572,21 +555,22 @@ __kernel void intelblas_gemm_buffer_NT(
     output = (local_x == 5) ? _dot.s5 : output; \
     output = (local_x == 6) ? _dot.s6 : output; \
     output = (local_x == 7) ? _dot.s7 : output; \
-    if (beta != 0.0) \
+    if (beta != 0.0f) \
         dst_write0[0] = mad(output, (float)alpha, ((float)beta * dst_write0[0])); \
     else \
         dst_write0[0] = output * (float)alpha; \
     dst_write0 += ldC;
 
-    if(global_x < N && global_y * 8 < M) {
-        OUTPUT(dot00);
-        if(mad24(global_y, 8, 1) < M) { OUTPUT(dot01); }
-        if(mad24(global_y, 8, 2) < M) { OUTPUT(dot02); }
-        if(mad24(global_y, 8, 3) < M) { OUTPUT(dot03); }
-        if(mad24(global_y, 8, 4) < M) { OUTPUT(dot04); }
-        if(mad24(global_y, 8, 5) < M) { OUTPUT(dot05); }
-        if(mad24(global_y, 8, 6) < M) { OUTPUT(dot06); }
-        if(mad24(global_y, 8, 7) < M) { OUTPUT(dot07); }
+    if (global_x < N && dst_row < M)
+    {
+        /*if (dst_row + 0 < M)*/ { OUTPUT(dot00); }
+        if (dst_row + 1 < M) { OUTPUT(dot01); }
+        if (dst_row + 2 < M) { OUTPUT(dot02); }
+        if (dst_row + 3 < M) { OUTPUT(dot03); }
+        if (dst_row + 4 < M) { OUTPUT(dot04); }
+        if (dst_row + 5 < M) { OUTPUT(dot05); }
+        if (dst_row + 6 < M) { OUTPUT(dot06); }
+        if (dst_row + 7 < M) { OUTPUT(dot07); }
     }
 #undef OUTPUT
 }