{"action":"created","comment":{"url":"https://api.github.com/repos/opencv/opencv/pulls/comments/137175671","pull_request_review_id":60807321,"id":137175671,"diff_hunk":"@@ -0,0 +1,2000 @@\n+/*M///////////////////////////////////////////////////////////////////////////////////////\n+//\n+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\n+//\n+//  By downloading, copying, installing or using the software you agree to this license.\n+//  If you do not agree to this license, do not download, install,\n+//  copy or use the software.\n+//\n+//\n+//                           License Agreement\n+//                For Open Source Computer Vision Library\n+//\n+// Copyright (C) 2017, Intel Corporation, all rights reserved.\n+// Copyright (c) 2016-2017 Fabian David Tschopp, all rights reserved.\n+// Third party copyrights are property of their respective owners.\n+//\n+// Redistribution and use in source and binary forms, with or without modification,\n+// are permitted provided that the following conditions are met:\n+//\n+//   * Redistribution's of source code must retain the above copyright notice,\n+//     this list of conditions and the following disclaimer.\n+//\n+//   * Redistribution's in binary form must reproduce the above copyright notice,\n+//     this list of conditions and the following disclaimer in the documentation\n+//     and/or other materials provided with the distribution.\n+//\n+//   * The name of the copyright holders may not be used to endorse or promote products\n+//     derived from this software without specific prior written permission.\n+//\n+// This software is provided by the copyright holders and contributors \"as is\" and\n+// any express or implied warranties, including, but not limited to, the implied\n+// warranties of merchantability and fitness for a particular purpose are disclaimed.\n+// In no event shall the Intel Corporation or contributors be liable for any direct,\n+// indirect, incidental, special, exemplary, or consequential damages\n+// (including, but not limited to, procurement of substitute goods or services;\n+// loss of use, data, or profits; or business interruption) however caused\n+// and on any theory of liability, whether in contract, strict liability,\n+// or tort (including negligence or otherwise) arising in any way out of\n+// the use of this software, even if advised of the possibility of such damage.\n+//\n+//M*/\n+\n+#include \"../../precomp.hpp\"\n+#include <string>\n+#include <vector>\n+#include <fstream>\n+#include <sys/stat.h>\n+#include <assert.h>\n+#include \"common.hpp\"\n+#include \"ocl4dnn.hpp\"\n+#include \"opencl_kernels_dnn.hpp\"\n+#include \"math_functions.hpp\"\n+\n+#if defined WIN32 || defined _WIN32\n+#include <windows.h>\n+#include <direct.h>\n+#endif\n+\n+#ifdef HAVE_OPENCL\n+namespace cv { namespace dnn { namespace ocl4dnn {\n+template<typename Dtype>\n+OCL4DNNConvSpatial<Dtype>::OCL4DNNConvSpatial(OCL4DNNConvConfig config)\n+{\n+    bias_term_ = config.bias_term;\n+    int32_t dims = config.in_shape.size();\n+    int32_t spatial_dims = config.kernel.size();\n+\n+    channels_   = config.in_shape[dims - spatial_dims - 1];\n+    num_output_ = config.out_shape[dims - spatial_dims - 1];\n+    group_ = config.group;\n+\n+    std::vector<int32_t> pad_;\n+    std::vector<int32_t> stride_;\n+    std::vector<int32_t> dilation_;\n+    std::vector<int32_t> kernel_shape_;\n+    std::vector<int32_t> im_in_shape_;\n+    std::vector<int32_t> im_out_shape_;\n+\n+    for (int32_t i = 0; i < spatial_dims; ++i) {\n+        kernel_shape_.push_back(config.kernel[i]);\n+        pad_.push_back(config.pad[i]);\n+        stride_.push_back(config.stride[i]);\n+        dilation_.push_back(config.dilation[i]);\n+        im_in_shape_.push_back(config.in_shape[dims - spatial_dims + i]);\n+        im_out_shape_.push_back(config.out_shape[dims - spatial_dims + i]);\n+    }\n+\n+    bestKernelConfig = NULL;\n+    prev_kernel_type_ = -1;\n+    bias_ = NULL;\n+    tuned_ = false;\n+    kernel_dim_ = channels_ / group_;\n+    out_spatial_dim_ = 1;\n+\n+    int32_t in_spatial_dim_ = 1;\n+    for (int32_t i = 0; i < spatial_dims; ++i) {\n+        kernel_dim_ *= config.kernel[i];\n+        in_spatial_dim_ *= config.in_shape[dims - spatial_dims + i];\n+        out_spatial_dim_ *= config.out_shape[dims - spatial_dims + i];\n+    }\n+\n+    // assumption: spatial dimension is 2.\n+    kernel_h_ = kernel_shape_[0];\n+    kernel_w_ = kernel_shape_[1];\n+    pad_h_ = pad_[0];\n+    pad_w_ = pad_[1];\n+    stride_h_ = stride_[0];\n+    stride_w_ = stride_[1];\n+    dilation_h_ = dilation_[0];\n+    dilation_w_ = dilation_[1];\n+    M_ = num_output_ / group_;\n+    K_ = channels_ * kernel_h_ * kernel_w_ / group_;\n+    height_ = im_in_shape_[0];\n+    width_ = im_in_shape_[1];\n+    output_h_ = im_out_shape_[0];\n+    output_w_ = im_out_shape_[1];\n+    bottom_dim_ = channels_ * in_spatial_dim_;\n+    top_dim_ = num_output_ * out_spatial_dim_;\n+    auto_tuning_ = false;\n+\n+    if (std::getenv(\"OPENCV_OCL4DNN_KERNEL_CONFIG_PATH\")) {\n+        cache_path_ << std::getenv(\"OPENCV_OCL4DNN_KERNEL_CONFIG_PATH\");\n+    }\n+\n+    bool hasCacheDir = false;\n+#if defined WIN32 || defined _WIN32\n+    if (std::getenv(\"HOME\"))\n+        cache_path_ << std::getenv(\"USERPROFILE\");\n+\n+    struct _stat file_stat;\n+    cache_path_ << \"\\\\spatialkernels\\\\\";\n+    hasCacheDir = _stat(cache_path_.str().c_str(), &file_stat) == 0 &&\n+                  ((_S_IFDIR & file_stat.st_mode) != 0);\n+#else\n+    if (std::getenv(\"HOME\"))\n+        cache_path_ << std::getenv(\"HOME\");\n+\n+    struct stat file_stat;\n+    cache_path_ << \"/spatialkernels/\";\n+    hasCacheDir = stat(cache_path_.str().c_str(), &file_stat) == 0 &&\n+                  S_ISDIR(file_stat.st_mode);\n+#endif\n+\n+    if (!hasCacheDir) {\n+#if defined WIN32 || defined _WIN32\n+        int result = _mkdir(cache_path_.str().c_str());\n+#else\n+        int result = mkdir(cache_path_.str().c_str(), 0755);\n+#endif\n+        hasCacheDir = (result == -1) ? false : true;\n+    }\n+    if (!hasCacheDir) {\n+        std::cout << \"Failed to create cache directory: \"\n+                  << cache_path_.str()\n+                  << \", disable auto-tuning.\" << std::endl;\n+    }\n+    auto_tuning_ = hasCacheDir;\n+}\n+\n+template<typename Dtype>\n+OCL4DNNConvSpatial<Dtype>::~OCL4DNNConvSpatial()\n+{\n+    swizzled_weights_.release();\n+\n+    if (bestKernelConfig) {\n+        delete bestKernelConfig;\n+    }\n+}\n+\n+template<typename Dtype>\n+std::string OCL4DNNConvSpatial<Dtype>::generateHeader()\n+{\n+    std::stringstream ss;\n+\n+    if (std::is_same<Dtype, double>::value) {\n+        // Test/enable KHR 64 bit (double)\n+        ss << \"#if defined(cl_khr_fp64)\" << std::endl;\n+        ss << \"#pragma OPENCL EXTENSION cl_khr_fp64 : enable\" << std::endl;\n+        ss << \"#define DOUBLE_SUPPORT_AVAILABLE\" << std::endl;\n+\n+        // Test/enable AMD 64 bit (double)\n+        ss << \"#elif defined(cl_amd_fp64)\" << std::endl;\n+        ss << \"#pragma OPENCL EXTENSION cl_amd_fp64 : enable\" << std::endl;\n+        ss << \"#define DOUBLE_SUPPORT_AVAILABLE\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+    }\n+\n+    // Test/enable 32 bit atomics\n+    ss << \"#if defined(cl_khr_int32_base_atomics)\" << std::endl;\n+    ss << \"#pragma OPENCL EXTENSION cl_khr_int32_base_atomics : enable\"\n+       << std::endl;\n+    ss << \"#define ATOMICS_32_AVAILABLE\" << std::endl;\n+    ss << \"#endif\" << std::endl;\n+    ss << \"#if defined(cl_khr_global_int32_base_atomics)\" << std::endl;\n+    ss << \"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\"\n+       << std::endl;\n+    ss << \"#define ATOMICS_32_AVAILABLE\" << std::endl;\n+    ss << \"#endif\" << std::endl;\n+\n+    // 64 bit integers\n+    if (sizeof(int32_t) == 8 || std::is_same<Dtype, double>::value) {\n+        // Test/enable 64 bit atomics\n+        ss << \"#if defined(cl_khr_int64_base_atomics)\" << std::endl;\n+        ss << \"#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable\"\n+           << std::endl;\n+        ss << \"#define ATOMICS_64_AVAILABLE\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+    }\n+\n+    if (std::is_same<Dtype, double>::value) {\n+        ss << \"#define Dtype double\" << std::endl;\n+        ss << \"#define Dtype1 double\" << std::endl;\n+        // double2, double4, double8, double16\n+        for (int32_t i = 2; i <= 16; i *= 2) {\n+            ss << \"#define Dtype\" << i << \" double\" << i << std::endl;\n+        }\n+    } else {\n+        ss << \"#define Dtype float\" << std::endl;\n+        ss << \"#define Dtype1 float\" << std::endl;\n+        // float2, float4, float8, float16\n+        for (int32_t i = 2; i <= 16; i *= 2) {\n+            ss << \"#define Dtype\" << i << \" float\" << i << std::endl;\n+        }\n+    }\n+\n+    if (sizeof(int32_t) == 8) {\n+        ss << \"#define int32_t long\" << std::endl;\n+        ss << \"#define uint32_t unsigned long\" << std::endl;\n+        ss << \"#define int32_tc long\" << std::endl;\n+        ss << \"#define uint32_tc unsigned long\" << std::endl;\n+    } else {\n+        ss << \"#define int32_t int\" << std::endl;\n+        ss << \"#define uint32_t unsigned int\" << std::endl;\n+        ss << \"#define int32_tc int\" << std::endl;\n+        ss << \"#define uint32_tc unsigned int\" << std::endl;\n+    }\n+\n+    return ss.str();\n+}\n+\n+template<typename Dtype>\n+std::string OCL4DNNConvSpatial<Dtype>::generateDefs()\n+{\n+    std::stringstream ss;\n+\n+    ss << \"#define __CAT(x, y) x##y\" << std::endl;\n+    ss << \"#define CAT(x, y) __CAT(x, y)\" << std::endl;\n+    ss << \"#define LOOP0(VAR, STMT)\" << std::endl;\n+    ss << \"#define LOOP1(VAR, STMT) (STMT); (VAR)++;\" << std::endl;\n+    ss << \"#define LOOP2(VAR, STMT) LOOP1(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP3(VAR, STMT) LOOP2(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP4(VAR, STMT) LOOP3(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP5(VAR, STMT) LOOP4(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP6(VAR, STMT) LOOP5(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP7(VAR, STMT) LOOP6(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP8(VAR, STMT) LOOP7(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP9(VAR, STMT) LOOP8(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP10(VAR, STMT) LOOP9(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP11(VAR, STMT) LOOP10(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP12(VAR, STMT) LOOP11(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP13(VAR, STMT) LOOP12(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP14(VAR, STMT) LOOP13(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP15(VAR, STMT) LOOP14(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP16(VAR, STMT) LOOP15(VAR, STMT); (STMT); (VAR)++;\"\n+       << std::endl;\n+    ss << \"#define LOOP(N, VAR, STMT) CAT(LOOP, N)((VAR), (STMT))\"\n+       << std::endl;\n+\n+    addDef(ss, \"KERNEL_WIDTH\", kernel_w_);\n+    addDef(ss, \"KERNEL_HEIGHT\" , kernel_h_);\n+    addDef(ss, \"STRIDE_X\", stride_w_);\n+    addDef(ss, \"STRIDE_Y\", stride_h_);\n+    addDef(ss, \"DILATION_X\", dilation_w_);\n+    addDef(ss, \"DILATION_Y\", dilation_h_);\n+    addDef(ss, \"INPUT_PAD_W\", pad_w_);\n+    addDef(ss, \"INPUT_PAD_H\", pad_h_);\n+\n+    return ss.str();\n+}\n+\n+typedef enum {\n+    KERNEL_TYPE_INTEL_IDLF = 2,\n+    KERNEL_TYPE_BASIC = 4,\n+    KERNEL_TYPE_GEMM_LIKE = 5\n+} ocl4dnnConvSpatialKernelType_t;\n+\n+template<typename Dtype>\n+std::string OCL4DNNConvSpatial<Dtype>::generateKernels(int32_t kernelType,\n+                                                      int32_t blockM,\n+                                                      int32_t blockK,\n+                                                      int32_t blockN)\n+{\n+    std::stringstream ss;\n+    std::stringstream opts;\n+    std::string kernelUKey;\n+    int32_t simd_size;\n+\n+    if (kernelType == KERNEL_TYPE_INTEL_IDLF) {\n+        simd_size = blockN;\n+        kernelUKey = generateSpecificKey(2, blockM, blockK, 1);\n+\n+        // kernel name\n+        kernel_name_ = \"IDLF_\";\n+        kernel_name_ += kernelUKey.c_str();\n+        if (simd_size == 16)\n+            kernel_name_ += \"_SIMD16\";\n+        else\n+            kernel_name_ += \"_SIMD8\";\n+\n+        // options\n+        opts << \"-cl-fast-relaxed-math -D convolve_simd=\" << kernel_name_;\n+        if (isBeignet())\n+            opts << \" -D__BEIGNET__ \";\n+        else\n+            opts << \" -cl-no-subgroup-ifp \";\n+        options_ = opts.str();\n+\n+        // defs\n+        int32_t output_width = output_w_;\n+        int32_t output_height = output_h_;\n+        int32_t output_block_width = blockM;\n+        int32_t output_block_height = blockK;\n+        const int32_t last_block_width = (output_width % output_block_width == 0) ?\n+                                        output_block_width : output_width % output_block_width;\n+        const int32_t last_block_height = (output_height % output_block_height == 0) ?\n+                                         output_block_height : output_height % output_block_height;\n+        int tile_x = (((output_block_width - 1) * stride_w_ + kernel_w_ * dilation_w_) + 3) & ~3;\n+        int tile_y = (output_block_height -1) * stride_h_ + kernel_h_ * dilation_h_;\n+        int tile_y_stride = (4 * simd_size) / tile_x;\n+        int invec_size = (tile_y + tile_y_stride - 1) / tile_y_stride;\n+\n+        addDef(ss, \"SIMD_SIZE\", simd_size);\n+        addDef(ss, \"filter_qualifier\", \"__global\");\n+        addDef(ss, \"OUT_BLOCK_WIDTH\", output_block_width);\n+        addDef(ss, \"OUT_BLOCK_HEIGHT\", output_block_height);\n+        addDef(ss, \"LAST_BLOCK_WIDTH\", last_block_width);\n+        addDef(ss, \"LAST_BLOCK_HEIGHT\", last_block_height);\n+        addDef(ss, \"INPUT_DEPTH\", channels_ / group_);\n+        addDef(ss, \"TOTAL_INPUT_DEPTH_SIZE\", channels_);\n+        addDef(ss, \"TOTAL_OUTPUT_DEPTH\", num_output_);\n+        addDef(ss, \"INPUT_START_X\", 0);\n+        addDef(ss, \"INPUT_START_Y\", 0);\n+        addDef(ss, \"INPUT_START_Z\", 0);\n+        addDef(ss, \"NUM_FILTERS\", M_);\n+        addDef(ss, \"OUT_BUFF_OFFSET\", 0);\n+        addDef(ss, \"TILE_X\", tile_x);\n+        addDef(ss, \"TILE_Y\", tile_y);\n+        addDef(ss, \"TILE_Y_STRIDE\", tile_y_stride);\n+        addDef(ss, \"INVEC_SIZE\", invec_size);\n+        addDef(ss, \"ALIGNED_NUM_FILTERS\", ALIGN(M_, simd_size));\n+        addDef(ss, \"OUT_BLOCK_SIZE\", (output_block_width*output_block_height));\n+\n+        // kernel source\n+        // Each work-item computes\n+        // a OUT_BLOCK_WIDTH * OUT_BLOCK_HEIGHT region of one output map.\n+        // Each work-group (which will be mapped to 1 SIMD16/SIMD8 EU thread)\n+        // will compute 16/8 different feature maps,\n+        // but each feature map is for the same region of the imput image.\n+        // NDRange:  (output_width+pad)/ OUT_BLOCK_WIDTH,\n+        //           (output_height+pad)/OUT_BLOCK_HEIGHT,\n+        //           NUM_FILTERS/OUT_BLOCK_DEPTH\n+        // NOTE: for beignet\n+        // this reqd_work_group_size does not guarantee that\n+        // SIMD16/8 mode will be used,\n+        // the compiler could choose to use two SIMD8 threads,\n+        // and if that happens the code will break.\n+        ss << \"#define activation_function(x) (x)\" << std::endl;\n+        ss << \"__attribute__((reqd_work_group_size(1, 1, SIMD_SIZE)))\" << std::endl;\n+        ss << \"kernel void\" << std::endl;\n+        ss << \"convolve_simd(\" << std::endl;\n+        ss << \"__global float* inputs_base,\" << std::endl;\n+        ss << \"filter_qualifier float* weights_base,\" << std::endl;\n+        ss << \"__global float* biases_base,\" << std::endl;\n+        ss << \"__global float* outputs_base,\" << std::endl;\n+        ss << \"const ushort input_width,\" << std::endl;\n+        ss << \"const ushort input_height,\" << std::endl;\n+        ss << \"const ushort output_width,\" << std::endl;\n+        ss << \"const ushort output_height)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"__global float* outputs = outputs_base;\" << std::endl;\n+        ss << \"__global float* inputs = inputs_base;\" << std::endl;\n+        ss << \"filter_qualifier float* weights = weights_base;\" << std::endl;\n+        ss << \"__global float* biases = biases_base;\" << std::endl;\n+        // oc = Output Column\n+        ss << \"uint32_t oc = get_global_id(0) * OUT_BLOCK_WIDTH;\" << std::endl;\n+        // or = Output Row\n+        ss << \"uint32_t or = get_global_id(1) * OUT_BLOCK_HEIGHT;\" << std::endl;\n+        // fm = Feature Map = od = Output Depth\n+        ss << \"uint32_t fm = get_global_id(2);\" << std::endl;\n+        ss << \"uint32_t fmg = get_group_id(2);\" << std::endl;\n+        ss << \"uint32_t lid = get_local_id(2);\" << std::endl;\n+        ss << \"float out[OUT_BLOCK_SIZE];\" << std::endl;\n+        ss << \"int32_t in_addr;\" << std::endl;\n+        // find weights adress of given neuron (lid is index)\n+        ss << \"uint32_t weight_addr = (fmg % (ALIGNED_NUM_FILTERS/SIMD_SIZE)) * \"\n+           << \"INPUT_DEPTH * KERNEL_WIDTH * KERNEL_HEIGHT * SIMD_SIZE + lid;\"\n+           << std::endl;\n+        ss << \"for(int32_t i=0;i<OUT_BLOCK_SIZE;i++) {\" << std::endl;\n+        ss << \"out[i]=0.0f;\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"uint32_t num_in_batch = ( fm ) / ALIGNED_NUM_FILTERS;\" << std::endl;\n+        ss << \"uint32_t input_batch_offset = \"\n+           << \"num_in_batch * input_height * input_width * TOTAL_INPUT_DEPTH_SIZE;\"\n+           << std::endl;\n+        ss << \"int curr_y = or * STRIDE_Y + INPUT_START_Y + (lid / (TILE_X/4));\"\n+           << std::endl;\n+        ss << \"int curr_x = oc * STRIDE_X + INPUT_START_X + (lid % (TILE_X/4)) * 4;\"\n+           << std::endl;\n+        ss << \"#if INPUT_PAD_W != 0 || INPUT_PAD_H != 0\" << std::endl;\n+        ss << \"int saved_y = curr_y;\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"in_addr = \"\n+           << \"input_batch_offset + INPUT_START_Z * input_height * input_width\"\n+            // y tile offset\n+           << \"+  (curr_y - INPUT_PAD_H) * input_width\"\n+            // x tile offset\n+           << \"+   curr_x - INPUT_PAD_W;\"\n+           << std::endl;\n+        ss << \"union {\" << std::endl;\n+        ss << \"float4 in_vec[INVEC_SIZE];\" << std::endl;\n+        ss << \"float in_array[INVEC_SIZE * 4];\" << std::endl;\n+        ss << \"} in_buf;\" << std::endl;\n+        ss << \"for(int32_t kd = 0; kd < INPUT_DEPTH; kd++)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"int32_t in_offset = in_addr;\" << std::endl;\n+        ss << \"int32_t reg = 0;\" << std::endl;\n+        ss << \"LOOP(INVEC_SIZE, reg,\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"#if INPUT_PAD_W != 0 || INPUT_PAD_H != 0\" << std::endl;\n+        ss << \"if (curr_y >= INPUT_PAD_H && \"\n+           << \"curr_y < input_height + INPUT_PAD_H && \"\n+           << \"curr_x + 3 >= INPUT_PAD_W && \"\n+           << \"curr_x < input_width + INPUT_PAD_W) {\" << std::endl;\n+        ss << \"if (curr_x < INPUT_PAD_W) {\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s0 = 0;\" << std::endl;\n+        ss << \"if (curr_x + 1 >= INPUT_PAD_W)\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s1 = *(inputs + in_offset + 1);\" << std::endl;\n+        ss << \"else\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s1 = 0;\" << std::endl;\n+        ss << \"if (curr_x + 2 >= INPUT_PAD_W)\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s2 = *(inputs + in_offset + 2);\" << std::endl;\n+        ss << \"else\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s2 = 0;\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s3 = *(inputs + in_offset + 3);\" << std::endl;\n+        ss << \"} else {\" << std::endl;\n+        // read SIMD_SIZE elements\n+        ss << \"in_buf.in_vec[reg] = *(global float4*)(inputs + in_offset);\"\n+           << std::endl;\n+        ss << \"if (curr_x + 1 >= input_width + INPUT_PAD_W)\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s1 = 0;\" << std::endl;\n+        ss << \"if (curr_x + 2 >= input_width + INPUT_PAD_W)\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s2 = 0;\" << std::endl;\n+        ss << \"if (curr_x + 3 >= input_width + INPUT_PAD_W)\" << std::endl;\n+        ss << \"in_buf.in_vec[reg].s3 = 0;\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"} else {\" << std::endl;\n+        ss << \"in_buf.in_vec[reg] = 0;\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"curr_y += TILE_Y_STRIDE;\" << std::endl;\n+        ss << \"#else\" << std::endl;\n+        // read SIMD_SIZE elements\n+        ss << \"in_buf.in_vec[reg] = *(global float4*)(inputs + in_offset);\"\n+           << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"in_offset += input_width * TILE_Y_STRIDE;\" << std::endl;\n+        ss << \"});\" << std::endl;\n+        ss << \"in_addr += input_height * input_width;\" << std::endl;\n+        ss << \"#if INPUT_PAD_W != 0 || INPUT_PAD_H != 0\" << std::endl;\n+        ss << \"curr_y = saved_y;\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"#if KERNEL_WIDTH * KERNEL_HEIGHT != 1\" << std::endl;\n+        ss << \"#define WEIGHT_PREF 8\" << std::endl;\n+        ss << \"#else\" << std::endl;\n+        ss << \"#define WEIGHT_PREF 1\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"union {\" << std::endl;\n+        ss << \"float w[WEIGHT_PREF];\" << std::endl;\n+        ss << \"#if KERNEL_WIDTH * KERNEL_HEIGHT != 1\" << std::endl;\n+        ss << \"uint8 ui8;\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"} weight_buf;\" << std::endl;\n+        ss << \"int32_t w_idx=0;\" << std::endl;\n+        ss << \"uint32_t orig_weight_addr = weight_addr;\" << std::endl;\n+        ss << \"#if KERNEL_WIDTH * KERNEL_HEIGHT != 1\" << std::endl;\n+        ss << \"weight_buf.ui8 = \"\n+           << \"intel_sub_group_block_read8((__global uint *)&weights[weight_addr]);\"\n+           << std::endl;\n+        ss << \"weight_addr += SIMD_SIZE * WEIGHT_PREF;\" << std::endl;\n+        ss << \"#else\" << std::endl;\n+        ss << \"weight_buf.w[0] = as_float(\"\n+           << \"intel_sub_group_block_read((__global uint *)&weights[weight_addr]));\"\n+           << std::endl;\n+        ss << \"weight_addr += SIMD_SIZE * 1;\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"#define BLOCK_IN(n) \"\n+           << \"sub_group_broadcast(\"\n+           << \"in_buf.in_array[((n)%4) + ((n) / (TILE_Y_STRIDE * TILE_X)) * 4], \"\n+           << \"(((n) % (TILE_Y_STRIDE * TILE_X))/4))\" << std::endl;\n+        // kr = Kernel Row\n+        ss << \"int32_t kr = 0;\" << std::endl;\n+        ss << \"LOOP(KERNEL_HEIGHT, kr,\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        // kc = Kernel Column\n+        ss << \"int32_t kc = 0;\" << std::endl;\n+        ss << \"LOOP(KERNEL_WIDTH, kc,\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"for(int32_t br=0; br < OUT_BLOCK_HEIGHT; br++) {\" << std::endl;\n+        ss << \"for(int32_t bc=0; bc < OUT_BLOCK_WIDTH; bc++) {\" << std::endl;\n+        ss << \"float input = BLOCK_IN((br * STRIDE_Y + kr * DILATION_Y) * \"\n+           << \"TILE_X + bc * STRIDE_X + kc * DILATION_X);\" << std::endl;\n+        ss << \"out[br * OUT_BLOCK_WIDTH + bc] = \"\n+           << \"mad(weight_buf.w[w_idx % WEIGHT_PREF], \"\n+           << \"input, out[br * OUT_BLOCK_WIDTH + bc]);\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"#if KERNEL_WIDTH * KERNEL_HEIGHT > WEIGHT_PREF\" << std::endl;\n+        // We assume KERNEL_WIDTH is equal to KERNEL_HEIGHT here.\n+        ss << \"if ((w_idx + 1) % WEIGHT_PREF == 0\" << std::endl;\n+        ss << \"#if KERNEL_WIDTH * KERNEL_HEIGHT % 8 != 0\" << std::endl;\n+        ss << \"&& ((w_idx + 1) <= (KERNEL_WIDTH * KERNEL_HEIGHT - WEIGHT_PREF))\"\n+           << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \") {\" << std::endl;\n+        ss << \"weight_buf.ui8 = \"\n+           << \"intel_sub_group_block_read8((__global uint *)&weights[weight_addr]);\"\n+           << std::endl;\n+        // weights must be stored in just the right SIMD swizzled format\n+        // for this to work, see host code for details.\n+        ss << \"weight_addr += SIMD_SIZE * WEIGHT_PREF;\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"#if KERNEL_WIDTH*KERNEL_HEIGHT % 8 == 0\" << std::endl;\n+        // need to do nothing\n+        ss << \"#else\" << std::endl;\n+        ss << \"else if ((w_idx + 1) %  WEIGHT_PREF == 0 && \"\n+           << \"((w_idx + 1) > (KERNEL_WIDTH * KERNEL_HEIGHT - WEIGHT_PREF)))\"\n+           << std::endl;\n+        ss << \"#if KERNEL_WIDTH * KERNEL_HEIGHT % 8 == 1\" << std::endl;\n+        ss << \"weight_buf.w[0] = weights[weight_addr];\" << std::endl;\n+        ss << \"#elif KERNEL_WIDTH * KERNEL_HEIGHT % 8 == 2\" << std::endl;\n+        ss << \"weight_buf.ui8.s01 = \"\n+           << \"intel_sub_group_block_read2((__global uint *)&weights[weight_addr]);\"\n+           << std::endl;\n+        ss << \"#elif KERNEL_WIDTH * KERNEL_HEIGHT % 8 <= 4\" << std::endl;\n+        ss << \"weight_buf.ui8.s0123 = \"\n+           << \"intel_sub_group_block_read4((__global uint *)&weights[weight_addr]);\"\n+           << std::endl;\n+        ss << \"#else\" << std::endl;\n+        ss << \"weight_buf.ui8 = \"\n+           << \"intel_sub_group_block_read8((__global uint *)&weights[weight_addr]);\"\n+           << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"#endif\" << std::endl;\n+        ss << \"++w_idx;\" << std::endl;\n+        ss << \"});\" << std::endl;\n+        ss << \"});\" << std::endl;\n+        ss << \"weight_addr = \"\n+           << \"orig_weight_addr + KERNEL_WIDTH * KERNEL_HEIGHT * SIMD_SIZE;\"\n+           << std::endl;\n+        ss << \"}\" << std::endl;\n+        // dead code to work around possible compiler bug.\n+        ss << \"if (ALIGNED_NUM_FILTERS != NUM_FILTERS && fm > 0xfffffffeul) {\"\n+           << std::endl;\n+        ss << \"outputs[0] = BLOCK_IN(fm % SIMD_SIZE);\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"fm = fm % ALIGNED_NUM_FILTERS;\" << std::endl;\n+        ss << \"if ((ALIGNED_NUM_FILTERS == NUM_FILTERS || fm < NUM_FILTERS)) {\"\n+           << std::endl;\n+        ss << \"uint32_t out_addr = \"\n+           << \"OUT_BUFF_OFFSET + \"\n+           << \"( num_in_batch * TOTAL_OUTPUT_DEPTH + fm ) * \"\n+           << \"output_width * output_height;\"\n+           << std::endl;\n+        ss << \"out_addr += or * output_width + oc;\" << std::endl;\n+        ss << \"float bias = biases[fm];\" << std::endl;\n+        ss << \"for(uint32_t r = 0; r < OUT_BLOCK_HEIGHT; r++) {\" << std::endl;\n+        ss << \"if (r + or >= output_height) break;\" << std::endl;\n+        ss << \"for(uint32_t c = 0; c < OUT_BLOCK_WIDTH; c++) {\" << std::endl;\n+        ss << \"if (c + oc >= output_width) break;\" << std::endl;\n+        ss << \"outputs[out_addr + r * output_width + c] = activation_function(bias + out[r * OUT_BLOCK_WIDTH + c]);\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+    } else if (kernelType == KERNEL_TYPE_BASIC) {\n+        kernelUKey = generateSpecificKey(4, blockM, blockK, blockN);\n+        kernel_name_ = \"BASIC_\";\n+        kernel_name_ += kernelUKey.c_str();\n+\n+        // opts\n+        opts << \" -cl-fast-relaxed-math -D CFMultiNoPadding=\" << kernel_name_;\n+        if (isBeignet())\n+            opts << \" -D__BEIGNET__ \";\n+        else\n+            opts << \" -cl-no-subgroup-ifp \";\n+        options_ = opts.str();\n+\n+        // defs\n+        addDef(ss, \"CHANNELS\", channels_ / group_);\n+        addDef(ss, \"APPLY_BIAS\", bias_term_);\n+        addDef(ss, \"OUTPUT_Z\", M_);\n+        addDef(ss, \"ZPAR\", 1);\n+\n+        // kernel\n+        ss << \"#define ACTIVATION_FUNCTION(_dst_, _offset_, _data_) \"\n+           << \"do { (_dst_)[(_offset_)] = (_data_);} while(0)\" << std::endl;\n+        ss << \"__kernel void CFMultiNoPadding(\" << std::endl;\n+        ss << \"__global Dtype* image_data,\" << std::endl;\n+        ss << \"int32_t image_offset,\" << std::endl;\n+        ss << \"__global Dtype* kernel_data, \" << std::endl;\n+        ss << \"int32_t kernel_offset,\" << std::endl;\n+        ss << \"__global Dtype* bias,\" << std::endl;\n+        ss << \"const int32_t bias_offset,\" << std::endl;\n+        ss << \"__global Dtype* convolved_image, \" << std::endl;\n+        ss << \"const int32_t convolved_image_offset,\" << std::endl;\n+        ss << \"const ushort input_width,\" << std::endl;\n+        ss << \"const ushort input_height,\" << std::endl;\n+        ss << \"const ushort output_width,\" << std::endl;\n+        ss << \"const ushort output_height,\" << std::endl;\n+        ss << \"const ushort pad_w,\" << std::endl;\n+        ss << \"const ushort pad_h) {\" << std::endl;\n+        ss << \"const int32_t outputX = get_global_id(0);\" << std::endl;\n+        ss << \"const int32_t outputY = get_global_id(1);\" << std::endl;\n+        ss << \"const int32_t kernelNum = get_global_id(2)*ZPAR;\" << std::endl;\n+        ss << \"if(outputX < output_width && outputY < output_height)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"Dtype sum[ZPAR];\" << std::endl;\n+        ss << \"for(int32_t kern =0; kern < ZPAR; kern++)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"sum[kern] = 0.0f;\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"const int32_t org_y = outputY * STRIDE_Y - pad_h;\" << std::endl;\n+        ss << \"const int32_t org_x = outputX * STRIDE_X - pad_w;\" << std::endl;\n+        ss << \"const int32_t currentKernelOffset = \"\n+           << \"kernel_offset + kernelNum*KERNEL_HEIGHT*KERNEL_WIDTH*CHANNELS;\"\n+           << std::endl;\n+        ss << \"const int32_t biasIndex=bias_offset + kernelNum;\" << std::endl;\n+        ss << \"const int32_t local_image_offset = org_y*input_width + org_x;\"\n+           << std::endl;\n+        ss << \"const int32_t imageSize = input_width*input_height;\" << std::endl;\n+        ss << \"__global Dtype* image_dataPtrFloat = \"\n+           << \"(image_data + (image_offset + local_image_offset));\" << std::endl;\n+        ss << \"__global Dtype* kernel_dataPtrFloat = \"\n+           << \"(kernel_data + (currentKernelOffset));\" << std::endl;\n+        ss << \"for(int32_t c = 0; c < CHANNELS; c++)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"for(int32_t y = 0; y < KERNEL_HEIGHT; y++)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"for(int32_t x = 0; x < KERNEL_WIDTH; x++)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"if(!(org_y + y * DILATION_Y >= 0 && \"\n+           << \"org_y + y * DILATION_Y < input_height && \"\n+           << \"org_x + x * DILATION_X >= 0 && \"\n+           << \"org_x + x * DILATION_X < input_width))\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"continue;\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"for(int32_t kern =0; kern < ZPAR; kern++)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"sum[kern] += image_dataPtrFloat[x * DILATION_X] * \"\n+           << \"kernel_dataPtrFloat[kern*KERNEL_HEIGHT*KERNEL_WIDTH*CHANNELS + x];\"\n+           << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"image_dataPtrFloat += input_width * DILATION_Y;\" << std::endl;\n+        ss << \"kernel_dataPtrFloat += KERNEL_WIDTH;\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"image_dataPtrFloat += \"\n+           << \"imageSize - input_width*KERNEL_HEIGHT*DILATION_Y;\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"if(APPLY_BIAS == 1)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"for(int32_t kern = 0; kern < ZPAR; kern++)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"if(kernelNum+kern < OUTPUT_Z)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"int32_t offset = convolved_image_offset + \"\n+           << \"(kernelNum+kern)*output_height*output_width + \"\n+           << \"outputY*output_width + outputX;\" << std::endl;\n+        ss << \"ACTIVATION_FUNCTION(convolved_image, offset, sum[kern] + \"\n+           << \"bias[biasIndex +kern]);\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"else\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"for(int32_t kern = 0; kern < ZPAR; kern++)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"if(kernelNum+kern < OUTPUT_Z)\" << std::endl;\n+        ss << \"{\" << std::endl;\n+        ss << \"int32_t offset = convolved_image_offset + \"\n+           << \"(kernelNum+kern)*output_height*output_width + \"\n+           << \"outputY*output_width + outputX;\" << std::endl;\n+        ss << \"ACTIVATION_FUNCTION(convolved_image, offset, sum[kern]);\"\n+           << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+        ss << \"}\" << std::endl;\n+    }\n+    return ss.str();\n+}\n+\n+template<typename Dtype>\n+void OCL4DNNConvSpatial<Dtype>::generateKernelSrc()\n+{\n+    std::stringstream ss;\n+\n+    ss << generateHeader();\n+    ss << generateDefs();\n+    ss << generateKernels(kernelType_, blockM_, blockK_, blockN_);\n+\n+    kernel_ = ss.str();\n+}\n+\n+template<typename Dtype>\n+bool OCL4DNNConvSpatial<Dtype>::Forward(const UMat& bottom,\n+                                        const UMat& weight,\n+                                        const UMat& bias, UMat& top,\n+                                        int32_t batch_size)\n+{\n+    cl_int ret;\n+    bottom_data_ = (Dtype*)bottom.handle(ACCESS_READ);\n+    top_data_ = (Dtype*)top.handle(ACCESS_WRITE);\n+    weight_ = (Dtype*)weight.handle(ACCESS_READ);\n+    bias_ = (Dtype*)bias.handle(ACCESS_READ);\n+    num_ = batch_size;\n+\n+    prepareKernel();\n+    ret = convolve(bottom_data_, top_data_, 0, num_, bestKernelConfig);\n+\n+    return ret == CL_SUCCESS ? true : false;\n+}\n+\n+template<typename Dtype>\n+void OCL4DNNConvSpatial<Dtype>::tune(Dtype* top_data,\n+                                    const Dtype* weight,\n+                                    const Dtype* bias,\n+                                    const Dtype* bottom_data,\n+                                    int32_t batch_size)\n+{\n+    UMat verify_mat;\n+    Dtype *verify_data;\n+\n+    verify_mat.create(1, batch_size * num_output_ * out_spatial_dim_, CV_32FC1);\n+    verify_data = reinterpret_cast<Dtype*>(verify_mat.handle(ACCESS_WRITE));\n+\n+    calculateBenchmark(bottom_data, weight, bias, verify_data);\n+    setupConvolution(bottom_data, top_data, verify_data);\n+    verify_mat.release();\n+    CHECK_EQ(tuned_, true) << \"Spatial convolution auto-tuning failed.\";\n+}\n+\n+template<typename Dtype>\n+void OCL4DNNConvSpatial<Dtype>::calculateBenchmark(const Dtype* bottom,\n+                                                  const Dtype* w,\n+                                                  const Dtype* bias,\n+                                                  Dtype* verify_data)\n+{\n+    createBasicKernel(1, 1, 1);\n+    kernel_index_ = kernelQueue.size() - 1;\n+    convolve(bottom, verify_data, 0, num_, kernelQueue[kernel_index_]);\n+    ocl::Queue &queue = ocl::Queue::getDefault();\n+    clEnqueueCopyBuffer((cl_command_queue)queue.ptr(),\n+                        (cl_mem)top_data_,\n+                        (cl_mem)verify_data, 0, 0,\n+                        sizeof(float) * num_ * top_dim_, 0, NULL, NULL);\n+    CV_Assert(phash.find(kernelQueue[kernel_index_]->kernelName) != phash.end());\n+    unloadProgram(phash.find(kernelQueue[kernel_index_]->kernelName)->second);\n+    phash.erase(kernelQueue[kernel_index_]->kernelName);\n+    kernelQueue.pop_back();\n+    return;\n+}\n+\n+#define dbg\n+#ifdef dbg\n+#define dbgPrint(x) (x)\n+#else\n+#define dbgPrint(x)\n+#endif\n+\n+// For large enough input size, we do not need to tune kernels for different\n+// size. The reason is with large input size, there will be enough work items\n+// to feed al the EUs.\n+// FIXME for the gemm like convolution, switch back to eaxct image size.\n+\n+#define TUNING_SIZE(x) ((x) > 256 ? 256 : (ALIGN(x, 16)))\n+\n+// Computes 64-bit \"cyclic redundancy check\" sum, as specified in ECMA-182\n+template<typename Dtype>\n+uint64 OCL4DNNConvSpatial<Dtype>::crc64(const uchar* data, size_t size, uint64 crc0)\n+{\n+    static uint64 table[256];\n+    static bool initialized = false;\n+\n+    if( !initialized )\n+    {\n+        for( int i = 0; i < 256; i++ )\n+        {\n+            uint64 c = i;\n+            for( int j = 0; j < 8; j++ )\n+                c = ((c & 1) ? CV_BIG_UINT(0xc96c5795d7870f42) : 0) ^ (c >> 1);\n+            table[i] = c;\n+        }\n+        initialized = true;\n+    }\n+\n+    uint64 crc = ~crc0;\n+    for( size_t idx = 0; idx < size; idx++ )\n+        crc = table[(uchar)crc ^ data[idx]] ^ (crc >> 8);\n+\n+    return ~crc;\n+}\n+\n+template<typename Dtype>\n+void OCL4DNNConvSpatial<Dtype>::generateKey()\n+{\n+    std::stringstream keyBuilder;\n+    // FIXME: to support fuse?\n+    keyBuilder << kernel_w_ << \"_\"\n+               << kernel_h_ << \"_\"\n+               << channels_ << \"_\"\n+               << group_ << \"_\"\n+               << stride_h_ << \"_\"\n+               << stride_w_ << \"_\"\n+               << dilation_h_ << \"_\"\n+               << dilation_w_ << \"_\"\n+               << bias_term_ << \"_\"\n+               << TUNING_SIZE(width_) << \"_\"\n+               << TUNING_SIZE(height_) << \"_\"\n+               << pad_w_ << \"_\"\n+               << pad_h_ << \"_\"\n+               << num_ << \"_\"\n+               << M_;\n+\n+    std::string prefix = ocl::Device::getDefault().name() +\n+                         ocl::Device::getDefault().vendorName() +\n+                         ocl::Device::getDefault().driverVersion() +\n+                         std::to_string(ocl::Device::getDefault().maxComputeUnits());\n+    prefix = prefix + keyBuilder.str();\n+    key_ = std::to_string(crc64((uchar*)prefix.c_str(), prefix.size()));\n+    short_key_ = keyBuilder.str();\n+}\n+\n+template<typename Dtype>\n+std::string OCL4DNNConvSpatial<Dtype>::generateSpecificKey(int32_t type, int32_t blockWidth,\n+                                                            int32_t blockHeight, int32_t blockDepth)\n+{\n+    std::stringstream keyBuilder;\n+    keyBuilder << short_key_\n+               << \"_\" << type\n+               << \"_\" << blockWidth\n+               << \"_\" << blockHeight\n+               << \"_\" << blockDepth;\n+    return keyBuilder.str();\n+}\n+\n+template<typename Dtype>\n+void interleaveMatrix(Dtype* mem_dst, const Dtype *mem,\n+                      int r, int c, int interleavedRows, int nonInterleavedRows,\n+                      int blockWidth, int rowAlignment )\n+{\n+    CHECK_EQ(interleavedRows % 2, 0) <<\n+             \"interleaveMatrix only supports even values for interleavedRows.\";\n+\n+    size_t memSize = r * c * sizeof(float);\n+    size_t dstSize = memSize *\n+                     (interleavedRows + nonInterleavedRows * 2) /\n+                     (interleavedRows + nonInterleavedRows);\n+    memset(mem_dst, 0, dstSize);    // NOLINT\n+\n+    const int xStride = blockWidth;\n+    const int yStride = c * 2;\n+    const Dtype *pSrc = mem;\n+    Dtype* pDst = mem_dst;\n+    for (int y = 0; y < r;) {\n+        for (int rows = 0; rows < interleavedRows; rows += 2) {\n+            if ( y >= r ) break;\n+            if ((c % xStride) == 0) {\n+                for (int x = 0; x < c / xStride; x++) {\n+                    memcpy(pDst + x * xStride * 2,                         // NOLINT\n+                           pSrc + x * xStride,     xStride * sizeof(Dtype));\n+                    memcpy(pDst + x * xStride * 2 + xStride,               // NOLINT\n+                           pSrc + x * xStride + c, xStride * sizeof(Dtype));\n+                }\n+            } else {\n+                const int count = c / xStride;\n+                int x = 0;\n+                for (; x < count - 1; x++) {\n+                    memcpy(pDst + x * xStride * 2,                          // NOLINT\n+                           pSrc + x * xStride, xStride * sizeof(Dtype));\n+                    memcpy(pDst + x * xStride * 2 + xStride,                // NOLINT\n+                           pSrc + x * xStride + c, xStride * sizeof(Dtype));\n+                }\n+                memcpy(pDst + x * xStride * 2,                            // NOLINT\n+                       pSrc + x * xStride, xStride * sizeof(Dtype));\n+            }\n+            pSrc += yStride;\n+            pDst += yStride;\n+            y += 2;\n+        }\n+\n+        for (int rows = 0; rows < nonInterleavedRows; rows++) {\n+            if (y >= r) break;\n+            const int stride = rowAlignment;\n+            int remaining = c;\n+            for (int x = 0; x < c; x += stride) {\n+                if (remaining >= stride) {\n+                    memcpy(pDst + x * 2, pSrc + x, stride * sizeof(Dtype));    // NOLINT\n+                    remaining -=stride;\n+                } else {\n+                    memcpy(pDst + x * 2, pSrc + x, remaining * sizeof(Dtype));  // NOLINT\n+                }\n+            }\n+            pSrc += yStride / 2;\n+            pDst += yStride;\n+            y++;\n+        }\n+    }\n+}\n+\n+template<typename Dtype>\n+void OCL4DNNConvSpatial<Dtype>::swizzleWeights(const Dtype *bottom,\n+                                              const Dtype *top,\n+                                              int32_t swizzled_factor,\n+                                              bool interleave)\n+{\n+    // Simply skip the weight swizzle if we already got a swizzled_weights_\n+    // in test phase and not in auto tuning\n+    // This requires we always call convolve again with the winner configuration\n+    // during the auto tuning stage.\n+    if (tuned_ && !swizzled_weights_.empty())\n+        return;\n+\n+    if (swizzled_weights_.empty()) {\n+        swizzled_weights_.create(1, ((num_output_ + 15) & ~15) * channels_ *\n+                                 kernel_h_ * ((kernel_w_ + 1) & ~1), CV_32FC1);\n+    }\n+\n+    ocl::Queue queue = ocl::Queue::getDefault();\n+\n+    if (!interleave) {\n+        ocl::Kernel oclk_copy_weight(CL_KERNEL_SELECT(\"copyWeightsSwizzled\"), cv::ocl::dnn::conv_spatial_helper_oclsrc);\n+        cl_uint argIdx = 0;\n+\n+        int32_t channels = channels_ / group_;\n+        oclk_copy_weight.set(argIdx++, (cl_mem) weight_);\n+        oclk_copy_weight.set(argIdx++, (cl_mem) swizzled_weights_.handle(ACCESS_WRITE));\n+        oclk_copy_weight.set(argIdx++, kernel_w_);\n+        oclk_copy_weight.set(argIdx++, kernel_h_);\n+        oclk_copy_weight.set(argIdx++, channels);\n+        oclk_copy_weight.set(argIdx++, this->num_output_);\n+        oclk_copy_weight.set(argIdx++, swizzled_factor);\n+        const size_t global_work_size_Copy[3] = {\n+            (size_t) (ALIGN(this->num_output_, swizzled_factor) * channels * kernel_w_ * kernel_h_), 1, 1 };\n+\n+        OCL_CHECK(clEnqueueNDRangeKernel((cl_command_queue)queue.ptr(),\n+                                         (cl_kernel)oclk_copy_weight.ptr(), 3, NULL,\n+                                         global_work_size_Copy, NULL, 0, NULL,\n+                                         NULL));\n+    } else {\n+        Dtype* cpu_weight =\n+            reinterpret_cast<Dtype*>(clEnqueueMapBuffer((cl_command_queue)queue.ptr(),\n+                                                        (cl_mem)weight_, true, CL_MAP_READ, 0,\n+                                                        sizeof(Dtype) * num_output_ * kernel_dim_ * group_,\n+                                                        0, NULL, NULL, NULL));\n+\n+        // assumption: kernel dimesion is 2\n+        Dtype* cpu_swizzled_weight =\n+            reinterpret_cast<Dtype*>(clEnqueueMapBuffer((cl_command_queue)queue.ptr(),\n+                                                        (cl_mem)swizzled_weights_.handle(ACCESS_READ),\n+                                                        true, CL_MAP_WRITE, 0,\n+                                                        sizeof(Dtype) *\n+                                                        ((num_output_ + 15) & ~15) *\n+                                                        channels_ * kernel_h_ * ((kernel_w_ + 1) & ~1),\n+                                                        0, NULL, NULL, NULL));\n+\n+        int interleavedRows = (kernel_w_ / 2) * 2;\n+        int nonInterleavedRows = kernel_w_ % 2;\n+        int blockWidth = swizzled_factor;  // should equal to simd size.\n+        int rowAlignment = 32;\n+        size_t interleaved_filter_size = M_ * kernel_w_ * kernel_h_ * channels_ * sizeof(Dtype);\n+        Dtype * tmpSwizzledWeight = reinterpret_cast<Dtype*>(malloc(interleaved_filter_size));\n+        CHECK_EQ(tmpSwizzledWeight != NULL, true) << \"Failed to allocate temporary swizzled weight\";\n+        for (int od = 0; od < M_; od++)\n+            for (int id = 0; id < channels_; id++)\n+                for (int r = 0; r < kernel_h_; r++)\n+                    for (int c = 0; c < kernel_w_; c++)\n+                        tmpSwizzledWeight[((id * kernel_h_ + r)* kernel_w_ + c) * M_ + od] =\n+                            cpu_weight[((od * channels_ + id) * kernel_h_ + r)*kernel_w_+c];\n+        interleaveMatrix(cpu_swizzled_weight,\n+                         tmpSwizzledWeight,\n+                         kernel_w_ * kernel_h_ * channels_, M_,\n+                         interleavedRows,\n+                         nonInterleavedRows,\n+                         blockWidth,\n+                         rowAlignment);\n+\n+        clEnqueueUnmapMemObject((cl_command_queue)queue.ptr(),\n+                                (cl_mem)weight_,\n+                                cpu_weight, 0, NULL,\n+                                NULL);\n+        clEnqueueUnmapMemObject((cl_command_queue)queue.ptr(),\n+                                (cl_mem)swizzled_weights_.handle(ACCESS_READ),\n+                                cpu_swizzled_weight, 0, NULL,\n+                                NULL);\n+        free(tmpSwizzledWeight);\n+    }\n+}\n+\n+template<>\n+void OCL4DNNConvSpatial<float>::computeGlobalSize(int32_t batch,\n+                                                 int32_t* wio,    // work item output size\n+                                                 size_t* lSize,  // local size\n+                                                 size_t* gSize)  // global size\n+{\n+    gSize[0] = ceil((fmax(static_cast<float>(output_w_) / wio[0], 1.0)) / lSize[0]) * lSize[0];\n+    gSize[1] = ceil((fmax(static_cast<float>(output_h_) / wio[1], 1.0)) / lSize[1]) * lSize[1];\n+    gSize[2] = ceil(static_cast<float>((ceil(static_cast<float>(M_) * batch / wio[2]))) / lSize[2]) * lSize[2];\n+}\n+\n+template<>\n+bool OCL4DNNConvSpatial<float>::createBasicKernel(int32_t blockWidth,\n+                                                   int32_t blockHeight, int32_t blockDepth)\n+{\n+    int32_t workItemOutput[3];\n+    workItemOutput[0] = 1;\n+    workItemOutput[1] = 1;\n+    workItemOutput[2] = 1;\n+\n+    kernelType_ = 4;\n+    blockM_ = blockWidth;\n+    blockK_ = blockHeight;\n+    blockN_ = blockDepth;\n+    generateKernelSrc();\n+    compileKernel();\n+\n+    size_t localSize[3] = { 1, 1, 1 };\n+    size_t globalSize[3];\n+\n+    computeGlobalSize(1, workItemOutput, localSize, globalSize);\n+    kernelQueue.push_back(new kernelConfig(kernel_name_, globalSize, localSize, workItemOutput,\n+                                           false, false, true, 4));\n+\n+    return true;\n+}\n+\n+template<typename Dtype>\n+void OCL4DNNConvSpatial<Dtype>::setBufferKernelArg(const Dtype *bottom, const Dtype *top,\n+                                                  ocl::Kernel *kernel,\n+                                                  const cl_uint &argIdx,\n+                                                  ocl::Context *ctx,\n+                                                  cl_mem buffer, size_t offset,\n+                                                  size_t size, bool readOnly,\n+                                                  bool preserved)\n+{\n+\n+    if (offset == 0) {\n+        kernel->set(argIdx, (cl_mem)buffer);\n+        return;\n+    }\n+\n+    if (preserved &&\n+        subBufferMap.find(std::make_tuple(buffer, offset, size)) != subBufferMap.end()) {\n+        kernel->set(argIdx, (cl_mem)(subBufferMap.find(std::make_tuple(buffer, offset, size))->second));\n+        return;\n+    }\n+    cl_buffer_region region;\n+    region.origin = offset * sizeof(Dtype);\n+    region.size = size * sizeof(Dtype);\n+    cl_mem_flags memFlags = readOnly ? CL_MEM_READ_ONLY : CL_MEM_READ_WRITE;\n+    cl_int error;\n+    cl_mem sub_buffer = clCreateSubBuffer(buffer, memFlags,\n+                                          CL_BUFFER_CREATE_TYPE_REGION,\n+                                          &region, &error);\n+    CHECK_EQ(error, CL_SUCCESS) << \"Failed to create sub buffer.\" << std::endl;\n+    if (error != CL_SUCCESS) {\n+        dbgPrint(std::cout << \"Failed to create sub buffer (\" << error << \").\" << std::endl);\n+        throw(error);\n+    }\n+    kernel->set(argIdx, (cl_mem)sub_buffer);\n+    if (preserved)\n+        subBufferMap.insert(std::make_pair(std::make_tuple(buffer, offset, size),\n+                            sub_buffer));\n+    else\n+        tmpSubBuffers.push_back(sub_buffer);\n+}\n+\n+template<typename Dtype>\n+void OCL4DNNConvSpatial<Dtype>::cleanTmpSubBuffers(const Dtype *bottom, const Dtype *top)\n+{\n+    for (auto &buffer : tmpSubBuffers)\n+        clReleaseMemObject(buffer);\n+    tmpSubBuffers.clear();\n+}\n+\n+template<>\n+cl_int OCL4DNNConvSpatial<float>::convolve(const float *bottom, const float *top,\n+                                          int32_t index,\n+                                          int32_t numImages, kernelConfig* config)\n+{\n+    ocl::Context ctx = ocl::Context::getDefault();\n+    ocl::Program program;\n+    phash_t::iterator it = phash.find(config->kernelName);\n+    if (it != phash.end())\n+        program = it->second;\n+    else\n+        return CL_INVALID_PROGRAM;\n+    ocl::Kernel kernel(config->kernelName.c_str(), program);\n+    cl_int err = CL_SUCCESS;\n+\n+    int32_t bias_offset;\n+    if (config->kernelType == KERNEL_TYPE_INTEL_IDLF) {\n+        swizzleWeights(bottom, top, config->workItem_output[2], false);\n+        size_t total_bottom_size = bottom_dim_ * numImages;\n+        size_t total_kernel_size = kernel_h_ * kernel_w_ * channels_ * M_;\n+        size_t total_bias_size = M_ * group_;\n+        size_t total_top_size = top_dim_ * numImages;\n+        for (int32_t g = 0; g < group_; ++g) {\n+            bias_offset = M_ * g;\n+            int32_t image_offset = width_ * height_ * (channels_ / group_) * g;\n+            int32_t output_image_offset = output_w_ * output_h_ * M_ * g;\n+\n+            int32_t kernel_offset = kernel_h_ * kernel_w_ * (channels_ / group_) * M_ * g;\n+            cl_uint argIdx = 0;\n+\n+            try {\n+                setBufferKernelArg(bottom, top, &kernel, argIdx++, &ctx,\n+                                   (cl_mem) bottom_data_,\n+                                   image_offset,\n+                                   total_bottom_size - image_offset,\n+                                   true, false);\n+                setBufferKernelArg(bottom, top, &kernel, argIdx++, &ctx,\n+                                   (cl_mem) swizzled_weights_.handle(ACCESS_READ),\n+                                   kernel_offset,\n+                                   total_kernel_size - kernel_offset,\n+                                   true, true);\n+                setBufferKernelArg(bottom, top, &kernel, argIdx++, &ctx,\n+                                   (cl_mem) bias_,\n+                                   bias_offset,\n+                                   total_bias_size - bias_offset,\n+                                   true, true);\n+                setBufferKernelArg(bottom, top, &kernel, argIdx++, &ctx,\n+                                   (cl_mem) top_data_,\n+                                   output_image_offset,\n+                                   total_top_size - output_image_offset,\n+                                   false, false);\n+            } catch (int e) {\n+                err = e;\n+            }\n+\n+            if (err == CL_SUCCESS) {\n+                kernel.set(argIdx++, (uint16_t)width_);\n+                kernel.set(argIdx++, (uint16_t)height_);\n+                kernel.set(argIdx++, (uint16_t)output_w_);\n+                kernel.set(argIdx++, (uint16_t)output_h_);\n+                ocl::Queue queue = ocl::Queue::getDefault();\n+                err = clEnqueueNDRangeKernel((cl_command_queue)queue.ptr(),\n+                                             (cl_kernel)kernel.ptr(), 3,\n+                                             NULL,\n+                                             config->global_work_size,\n+                                             config->local_work_size, 0, NULL,\n+                                             NULL);\n+            }\n+            if (err != CL_SUCCESS)\n+                break;\n+        }\n+\n+        if (group_ > 1) {\n+            cleanTmpSubBuffers(bottom, top);\n+        }\n+        if (err != CL_SUCCESS)\n+            return err;\n+    } else if (config->kernelType == KERNEL_TYPE_GEMM_LIKE) {\n+        swizzleWeights(bottom, top, config->workItem_output[1], true);\n+        size_t total_bottom_size = bottom_dim_ * numImages;\n+        size_t total_kernel_size = kernel_h_ * kernel_w_ * channels_ * M_;\n+        size_t total_bias_size = M_ * group_;\n+        size_t total_top_size = top_dim_ * numImages;\n+        for (int32_t g = 0; g < group_; ++g) {\n+            bias_offset = M_ * g;\n+            int32_t image_offset = width_ * height_ * (channels_ / group_) * g;\n+            int32_t output_image_offset = output_w_ * output_h_ * M_ * g;\n+\n+            cl_uint argIdx = 0;\n+            int32_t kernel_offset = kernel_h_ * kernel_w_ * (channels_ / group_) * M_ * g;\n+            try {\n+                setBufferKernelArg(bottom, top, &kernel, argIdx++, &ctx,\n+                                   (cl_mem) bottom_data_,\n+                                   image_offset,\n+                                   total_bottom_size - image_offset,\n+                                   true, false);\n+                setBufferKernelArg(bottom, top, &kernel, argIdx++, &ctx,\n+                                   (cl_mem) swizzled_weights_.handle(ACCESS_READ),\n+                                   kernel_offset,\n+                                   total_kernel_size - kernel_offset,\n+                                   true, true);\n+                setBufferKernelArg(bottom, top, &kernel, argIdx++, &ctx,\n+                                   (cl_mem) bias_,\n+                                   bias_offset,\n+                                   total_bias_size - bias_offset,\n+                                   true, true);\n+                setBufferKernelArg(bottom, top, &kernel, argIdx++, &ctx,\n+                                   (cl_mem) top_data_,\n+                                   output_image_offset,\n+                                   total_top_size - output_image_offset,\n+                                   false, false);\n+            } catch (int e) {\n+                err = e;\n+            }\n+\n+            if (err == CL_SUCCESS) {\n+                kernel.set(argIdx++, (uint16_t)width_);\n+                kernel.set(argIdx++, (uint16_t)height_);\n+                kernel.set(argIdx++, (uint16_t)output_w_);\n+                kernel.set(argIdx++, (uint16_t)output_h_);\n+                int out_pitch_y = output_w_ * output_h_;\n+                int out_pitch_z = out_pitch_y * M_;\n+                int aligned_input_size = height_ * width_ * channels_ / group_;\n+                int slice_pitch = width_ * height_;\n+                kernel.set(argIdx++, (uint32_t)out_pitch_y);\n+                kernel.set(argIdx++, (uint32_t)out_pitch_z);\n+                kernel.set(argIdx++, (uint32_t)aligned_input_size);\n+                kernel.set(argIdx++, (uint32_t)slice_pitch);\n+\n+                int blockM = config->workItem_output[0];\n+                int blockK = config->workItem_output[1];\n+                int blockN = config->workItem_output[2];\n+                int alignedFilterWidth = ALIGN(M_, blockN);\n+                int alignedExpandHeight = ALIGN(output_w_ * output_h_, blockM);\n+                int globalWorkSizeDX = blockN;\n+                int globalWorkSizeDY = blockM;\n+                size_t sgemm_m = alignedExpandHeight;\n+                size_t sgemm_n = alignedFilterWidth;\n+                size_t gx = (size_t) ceil( (float) sgemm_n /\n+                        (float) globalWorkSizeDX );\n+                size_t gy = (size_t) ceil( (float) sgemm_m /\n+                        (float) globalWorkSizeDY );\n+                gy = ALIGN(gy, blockK);\n+                size_t global_size[3] = { gx, gy, config->global_work_size[2] };\n+\n+                ocl::Queue queue = ocl::Queue::getDefault();\n+                err = clEnqueueNDRangeKernel((cl_command_queue)queue.ptr(),\n+                                             (cl_kernel)kernel.ptr(), 3,\n+                                             NULL,\n+                                             global_size,\n+                                             config->local_work_size, 0, NULL,\n+                                             NULL);\n+                OCL_CHECK(err);\n+            }\n+            if (err != CL_SUCCESS)\n+                break;\n+        }\n+\n+        if (group_ > 1) {\n+            cleanTmpSubBuffers(bottom, top);\n+        }\n+        if (err != CL_SUCCESS)\n+            return err;\n+    } else {\n+        for (int32_t n = 0; n < numImages; ++n) {\n+            for (int32_t g = 0; g < group_; ++g) {\n+                bias_offset = M_ * g;\n+                int32_t image_offset = n * bottom_dim_\n+                    + width_ * height_ * (channels_ / group_) * g;\n+                int32_t output_image_offset = n * top_dim_\n+                    + output_w_ * output_h_ * M_ * g;\n+\n+                cl_uint argIdx = 0;\n+                int32_t kernel_offset = kernel_h_ * kernel_w_ * (channels_ / group_) * M_ * g;\n+\n+                kernel.set(argIdx++, (cl_mem) bottom_data_);\n+                kernel.set(argIdx++, image_offset);\n+                kernel.set(argIdx++, (cl_mem) weight_);\n+                kernel.set(argIdx++, kernel_offset);\n+                kernel.set(argIdx++, (cl_mem) bias_);\n+                kernel.set(argIdx++, bias_offset);\n+                kernel.set(argIdx++, (cl_mem) top_data_);\n+                kernel.set(argIdx++, output_image_offset);\n+                kernel.set(argIdx++, (uint16_t)width_);\n+                kernel.set(argIdx++, (uint16_t)height_);\n+                kernel.set(argIdx++, (uint16_t)output_w_);\n+                kernel.set(argIdx++, (uint16_t)output_h_);\n+                kernel.set(argIdx++, (uint16_t)pad_w_);\n+                kernel.set(argIdx++, (uint16_t)pad_h_);\n+                ocl::Queue queue = ocl::Queue::getDefault();\n+                if (config->use_null_local) {\n+                    err = clEnqueueNDRangeKernel((cl_command_queue)queue.ptr(),\n+                                                 (cl_kernel)kernel.ptr(), 3,\n+                                                 NULL,\n+                                                 config->global_work_size, NULL, 0, NULL,\n+                                                 NULL);\n+                } else {\n+                    err = clEnqueueNDRangeKernel((cl_command_queue)queue.ptr(),\n+                                                 (cl_kernel)kernel.ptr(), 3,\n+                                                 NULL,\n+                                                 config->global_work_size,\n+                                                 config->local_work_size, 0, NULL,\n+                                                 NULL);\n+                }\n+\n+                if (err != CL_SUCCESS)\n+                    return err;\n+            }\n+        }\n+    }\n+\n+    return err;\n+}\n+\n+template<>\n+float OCL4DNNConvSpatial<float>::timedConvolve(const float *bottom,\n+                                              const float *top,\n+                                              int32_t index,\n+                                              int32_t numImages,\n+                                              kernelConfig* config)\n+{\n+    // warm up.\n+    bool saved_tuned = tuned_;\n+    tuned_ = false;\n+    convolve(bottom, top, index, num_, config);\n+    cv::ocl::Timer timer;\n+    timer.start();\n+    cl_int err;\n+    dbgPrint(std::cout << \"Benchmarking kernel: \" << config->kernelName << std::endl);\n+    tuned_ = true;\n+    int loop_cnt = 4;\n+    for (int i = 0; i < loop_cnt; i++) {\n+        err = convolve(bottom, top, index, num_, config);\n+        if (err != CL_SUCCESS)\n+            break;\n+    }\n+    tuned_ = saved_tuned;\n+    timer.stop();\n+    if (err != CL_SUCCESS) {\n+        config->tested = true;\n+        config->verified = false;\n+        dbgPrint(std::cout << \"convolution failed with error code \" << err << std::endl);\n+        return 1e5;\n+    }\n+\n+    float elapsedTime = timer.milliSeconds() / loop_cnt;\n+    #ifdef dbg\n+    double out_w = output_w_;\n+    double out_h = output_h_;\n+    double out_z = M_;\n+    double k_w = kernel_w_;\n+    double k_h = kernel_h_;\n+    double k_z = channels_;\n+    double totalFlops = ((k_w*k_h*k_z -1)*2)*(out_w*out_h*out_z)*num_;\n+    std::cout << \"\\tEstimated Gflops:\" << ((totalFlops/1000)/1000)/1000\n+              << std::endl;\n+    std::cout << \"\\tEstimated GFLOPS/S: \" << (((totalFlops/1000)/1000)/1000)*(1000.0/elapsedTime)\n+              << std::endl;\n+    #if 0\n+    std::cout << \"Estimated utilization: \" <<\n+        ((((totalFlops/1000)/1000)/1000)*(1000.0/elapsedTime))/880.0\n+        << std::endl;\n+    #endif\n+    #endif\n+    return elapsedTime;\n+}\n+\n+template<>\n+bool OCL4DNNConvSpatial<float>::verifyResult(const float *bottom, float *top,\n+                                             int32_t index,\n+                                             int32_t numImages, const float *verify_blob, kernelConfig* config)\n+{\n+\n+    uint32_t verificationFail = 0;\n+\n+    if (config->verified)\n+        return true;\n+    else if (config->tested)\n+        return false;\n+\n+    ocl4dnnSet(0, numImages * top_dim_, (float) 0, (cl_mem) top, 0);\n+    config->executionTime = timedConvolve(bottom, top, index, numImages, config);\n+    const float *verify_data;\n+    float *data;\n+    float *tmp_verify_data;\n+    ocl::Queue queue = ocl::Queue::getDefault();\n+    data = reinterpret_cast<float *>(clEnqueueMapBuffer((cl_command_queue)queue.ptr(),\n+                                                        (cl_mem)top, true, CL_MAP_READ,\n+                                                        0, sizeof(float) * numImages * top_dim_,\n+                                                        0, NULL, NULL, NULL));\n+    tmp_verify_data =\n+        reinterpret_cast<float *>(clEnqueueMapBuffer((cl_command_queue)queue.ptr(),\n+                                                     (cl_mem)verify_blob, true, CL_MAP_READ,\n+                                                     0, sizeof(float) * numImages * top_dim_,\n+                                                     0, NULL, NULL, NULL));\n+    verify_data = tmp_verify_data;\n+\n+    for (int32_t n = 0; n < numImages; ++n) {\n+        for (int32_t g = 0; g < group_; ++g) {\n+            int32_t output_image_offset = n * top_dim_ + output_w_ * output_h_ * M_ * g;\n+            for (int out_ch = 0; out_ch < M_ && !verificationFail; out_ch++)\n+                for (int h = 0; h < output_h_ && !verificationFail; h++)\n+                    for (int w = 0; w < output_w_; w++) {\n+                        size_t offset = output_image_offset + out_ch * output_w_ * output_h_ + h * output_w_ + w;\n+                        if (fabs(data[offset] - verify_data[offset]) > 0.1 * fabs(verify_data[offset]) &&\n+                            !(fabs(verify_data[offset]) < 1.e-3 &&\n+                            fabs(data[offset] - verify_data[offset]) < 1.e-4))\n+                        {\n+                            dbgPrint(printf(\"test verification failed @ image %d group %d\"\n+                                            \"out_ch %d h %d w %d got %G expected %G\\n\",\n+                                            n, g, out_ch, h, w, data[offset], verify_data[offset]));\n+                            verificationFail = 1;\n+                            goto out;\n+                        }\n+                    }\n+        }\n+    }\n+\n+out:\n+    clEnqueueUnmapMemObject((cl_command_queue)queue.ptr(),\n+                            (cl_mem)top, data, 0, NULL, NULL);\n+    clEnqueueUnmapMemObject((cl_command_queue)queue.ptr(),\n+                            (cl_mem)verify_blob, tmp_verify_data, 0, NULL, NULL);\n+    if (verificationFail == 1)\n+        return false;\n+    else\n+        return true;\n+}\n+\n+template<typename Dtype>\n+void OCL4DNNConvSpatial<Dtype>::unloadProgram(ocl::Program& prog)\n+{\n+    ocl::Context ctx = ocl::Context::getDefault();\n+    ctx.unloadProg(prog);\n+}\n+\n+template<typename Dtype>\n+ocl::Program OCL4DNNConvSpatial<Dtype>::compileKernel()\n+{\n+    String errmsg;\n+    ocl::Context ctx = ocl::Context::getDefault();\n+    ocl::ProgramSource src(kernel_.c_str());\n+    ocl::Program program = ctx.getProg(src, options_, errmsg);\n+    CV_Assert(!kernel_name_.empty());","path":"modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp","position":null,"original_position":1455,"commit_id":"d50b01714f6fdfd9c69a01a815414dc13e510584","original_commit_id":"44b4764647fd3ac4c918798c168dbc399af6bdb1","user":{"login":"pengli","id":985286,"avatar_url":"https://avatars1.githubusercontent.com/u/985286?v=4","gravatar_id":"","url":"https://api.github.com/users/pengli","html_url":"https://github.com/pengli","followers_url":"https://api.github.com/users/pengli/followers","following_url":"https://api.github.com/users/pengli/following{/other_user}","gists_url":"https://api.github.com/users/pengli/gists{/gist_id}","starred_url":"https://api.github.com/users/pengli/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/pengli/subscriptions","organizations_url":"https://api.github.com/users/pengli/orgs","repos_url":"https://api.github.com/users/pengli/repos","events_url":"https://api.github.com/users/pengli/events{/privacy}","received_events_url":"https://api.github.com/users/pengli/received_events","type":"User","site_admin":false},"body":"yes, we have checked this build option based on opencl version and carefully check the program compilation result if it fails.","created_at":"2017-09-06T05:48:00Z","updated_at":"2017-09-06T05:48:00Z","html_url":"https://github.com/opencv/opencv/pull/9114#discussion_r137175671","pull_request_url":"https://api.github.com/repos/opencv/opencv/pulls/9114","author_association":"MEMBER","_links":{"self":{"href":"https://api.github.com/repos/opencv/opencv/pulls/comments/137175671"},"html":{"href":"https://github.com/opencv/opencv/pull/9114#discussion_r137175671"},"pull_request":{"href":"https://api.github.com/repos/opencv/opencv/pulls/9114"}},"in_reply_to_id":136838366},"pull_request":{"url":"https://api.github.com/repos/opencv/opencv/pulls/9114","id":129342157,"html_url":"https://github.com/opencv/opencv/pull/9114","diff_url":"https://github.com/opencv/opencv/pull/9114.diff","patch_url":"https://github.com/opencv/opencv/pull/9114.patch","issue_url":"https://api.github.com/repos/opencv/opencv/issues/9114","number":9114,"state":"open","locked":false,"title":"add libdnn acceleration to dnn module ","user":{"login":"pengli","id":985286,"avatar_url":"https://avatars1.githubusercontent.com/u/985286?v=4","gravatar_id":"","url":"https://api.github.com/users/pengli","html_url":"https://github.com/pengli","followers_url":"https://api.github.com/users/pengli/followers","following_url":"https://api.github.com/users/pengli/following{/other_user}","gists_url":"https://api.github.com/users/pengli/gists{/gist_id}","starred_url":"https://api.github.com/users/pengli/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/pengli/subscriptions","organizations_url":"https://api.github.com/users/pengli/orgs","repos_url":"https://api.github.com/users/pengli/repos","events_url":"https://api.github.com/users/pengli/events{/privacy}","received_events_url":"https://api.github.com/users/pengli/received_events","type":"User","site_admin":false},"body":"libdnn provide ocl acceleration to current dnn module.  Currently it provide ocl kernels for 5 layers including convolution, softmax, LRN, fully connect and pooling.  In convolution it implements an auto-tuning mechanism to enumerate all possible convolution kernels and find the best performance one at first run time. So application can use the tuned kernel to achieve best performance. ","created_at":"2017-07-07T07:07:20Z","updated_at":"2017-09-06T05:48:00Z","closed_at":null,"merged_at":null,"merge_commit_sha":"348914dace76f7e5971f463af46cee383a0442c8","assignee":null,"assignees":[],"requested_reviewers":[],"milestone":null,"commits_url":"https://api.github.com/repos/opencv/opencv/pulls/9114/commits","review_comments_url":"https://api.github.com/repos/opencv/opencv/pulls/9114/comments","review_comment_url":"https://api.github.com/repos/opencv/opencv/pulls/comments{/number}","comments_url":"https://api.github.com/repos/opencv/opencv/issues/9114/comments","statuses_url":"https://api.github.com/repos/opencv/opencv/statuses/d50b01714f6fdfd9c69a01a815414dc13e510584","head":{"label":"pengli:dnn_rebase","ref":"dnn_rebase","sha":"d50b01714f6fdfd9c69a01a815414dc13e510584","user":{"login":"pengli","id":985286,"avatar_url":"https://avatars1.githubusercontent.com/u/985286?v=4","gravatar_id":"","url":"https://api.github.com/users/pengli","html_url":"https://github.com/pengli","followers_url":"https://api.github.com/users/pengli/followers","following_url":"https://api.github.com/users/pengli/following{/other_user}","gists_url":"https://api.github.com/users/pengli/gists{/gist_id}","starred_url":"https://api.github.com/users/pengli/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/pengli/subscriptions","organizations_url":"https://api.github.com/users/pengli/orgs","repos_url":"https://api.github.com/users/pengli/repos","events_url":"https://api.github.com/users/pengli/events{/privacy}","received_events_url":"https://api.github.com/users/pengli/received_events","type":"User","site_admin":false},"repo":{"id":71089972,"name":"opencv","full_name":"pengli/opencv","owner":{"login":"pengli","id":985286,"avatar_url":"https://avatars1.githubusercontent.com/u/985286?v=4","gravatar_id":"","url":"https://api.github.com/users/pengli","html_url":"https://github.com/pengli","followers_url":"https://api.github.com/users/pengli/followers","following_url":"https://api.github.com/users/pengli/following{/other_user}","gists_url":"https://api.github.com/users/pengli/gists{/gist_id}","starred_url":"https://api.github.com/users/pengli/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/pengli/subscriptions","organizations_url":"https://api.github.com/users/pengli/orgs","repos_url":"https://api.github.com/users/pengli/repos","events_url":"https://api.github.com/users/pengli/events{/privacy}","received_events_url":"https://api.github.com/users/pengli/received_events","type":"User","site_admin":false},"private":false,"html_url":"https://github.com/pengli/opencv","description":"Open Source Computer Vision Library","fork":true,"url":"https://api.github.com/repos/pengli/opencv","forks_url":"https://api.github.com/repos/pengli/opencv/forks","keys_url":"https://api.github.com/repos/pengli/opencv/keys{/key_id}","collaborators_url":"https://api.github.com/repos/pengli/opencv/collaborators{/collaborator}","teams_url":"https://api.github.com/repos/pengli/opencv/teams","hooks_url":"https://api.github.com/repos/pengli/opencv/hooks","issue_events_url":"https://api.github.com/repos/pengli/opencv/issues/events{/number}","events_url":"https://api.github.com/repos/pengli/opencv/events","assignees_url":"https://api.github.com/repos/pengli/opencv/assignees{/user}","branches_url":"https://api.github.com/repos/pengli/opencv/branches{/branch}","tags_url":"https://api.github.com/repos/pengli/opencv/tags","blobs_url":"https://api.github.com/repos/pengli/opencv/git/blobs{/sha}","git_tags_url":"https://api.github.com/repos/pengli/opencv/git/tags{/sha}","git_refs_url":"https://api.github.com/repos/pengli/opencv/git/refs{/sha}","trees_url":"https://api.github.com/repos/pengli/opencv/git/trees{/sha}","statuses_url":"https://api.github.com/repos/pengli/opencv/statuses/{sha}","languages_url":"https://api.github.com/repos/pengli/opencv/languages","stargazers_url":"https://api.github.com/repos/pengli/opencv/stargazers","contributors_url":"https://api.github.com/repos/pengli/opencv/contributors","subscribers_url":"https://api.github.com/repos/pengli/opencv/subscribers","subscription_url":"https://api.github.com/repos/pengli/opencv/subscription","commits_url":"https://api.github.com/repos/pengli/opencv/commits{/sha}","git_commits_url":"https://api.github.com/repos/pengli/opencv/git/commits{/sha}","comments_url":"https://api.github.com/repos/pengli/opencv/comments{/number}","issue_comment_url":"https://api.github.com/repos/pengli/opencv/issues/comments{/number}","contents_url":"https://api.github.com/repos/pengli/opencv/contents/{+path}","compare_url":"https://api.github.com/repos/pengli/opencv/compare/{base}...{head}","merges_url":"https://api.github.com/repos/pengli/opencv/merges","archive_url":"https://api.github.com/repos/pengli/opencv/{archive_format}{/ref}","downloads_url":"https://api.github.com/repos/pengli/opencv/downloads","issues_url":"https://api.github.com/repos/pengli/opencv/issues{/number}","pulls_url":"https://api.github.com/repos/pengli/opencv/pulls{/number}","milestones_url":"https://api.github.com/repos/pengli/opencv/milestones{/number}","notifications_url":"https://api.github.com/repos/pengli/opencv/notifications{?since,all,participating}","labels_url":"https://api.github.com/repos/pengli/opencv/labels{/name}","releases_url":"https://api.github.com/repos/pengli/opencv/releases{/id}","deployments_url":"https://api.github.com/repos/pengli/opencv/deployments","created_at":"2016-10-17T01:43:13Z","updated_at":"2016-10-17T01:43:26Z","pushed_at":"2017-09-06T05:18:31Z","git_url":"git://github.com/pengli/opencv.git","ssh_url":"git@github.com:pengli/opencv.git","clone_url":"https://github.com/pengli/opencv.git","svn_url":"https://github.com/pengli/opencv","homepage":"opencv.org","size":440522,"stargazers_count":0,"watchers_count":0,"language":"C++","has_issues":false,"has_projects":true,"has_downloads":true,"has_wiki":true,"has_pages":false,"forks_count":0,"mirror_url":null,"open_issues_count":0,"forks":0,"open_issues":0,"watchers":0,"default_branch":"master"}},"base":{"label":"opencv:master","ref":"master","sha":"1caca2112b1054b4e265b4bcf930ff10b5626545","user":{"login":"opencv","id":5009934,"avatar_url":"https://avatars3.githubusercontent.com/u/5009934?v=4","gravatar_id":"","url":"https://api.github.com/users/opencv","html_url":"https://github.com/opencv","followers_url":"https://api.github.com/users/opencv/followers","following_url":"https://api.github.com/users/opencv/following{/other_user}","gists_url":"https://api.github.com/users/opencv/gists{/gist_id}","starred_url":"https://api.github.com/users/opencv/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/opencv/subscriptions","organizations_url":"https://api.github.com/users/opencv/orgs","repos_url":"https://api.github.com/users/opencv/repos","events_url":"https://api.github.com/users/opencv/events{/privacy}","received_events_url":"https://api.github.com/users/opencv/received_events","type":"Organization","site_admin":false},"repo":{"id":5108051,"name":"opencv","full_name":"opencv/opencv","owner":{"login":"opencv","id":5009934,"avatar_url":"https://avatars3.githubusercontent.com/u/5009934?v=4","gravatar_id":"","url":"https://api.github.com/users/opencv","html_url":"https://github.com/opencv","followers_url":"https://api.github.com/users/opencv/followers","following_url":"https://api.github.com/users/opencv/following{/other_user}","gists_url":"https://api.github.com/users/opencv/gists{/gist_id}","starred_url":"https://api.github.com/users/opencv/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/opencv/subscriptions","organizations_url":"https://api.github.com/users/opencv/orgs","repos_url":"https://api.github.com/users/opencv/repos","events_url":"https://api.github.com/users/opencv/events{/privacy}","received_events_url":"https://api.github.com/users/opencv/received_events","type":"Organization","site_admin":false},"private":false,"html_url":"https://github.com/opencv/opencv","description":"Open Source Computer Vision Library","fork":false,"url":"https://api.github.com/repos/opencv/opencv","forks_url":"https://api.github.com/repos/opencv/opencv/forks","keys_url":"https://api.github.com/repos/opencv/opencv/keys{/key_id}","collaborators_url":"https://api.github.com/repos/opencv/opencv/collaborators{/collaborator}","teams_url":"https://api.github.com/repos/opencv/opencv/teams","hooks_url":"https://api.github.com/repos/opencv/opencv/hooks","issue_events_url":"https://api.github.com/repos/opencv/opencv/issues/events{/number}","events_url":"https://api.github.com/repos/opencv/opencv/events","assignees_url":"https://api.github.com/repos/opencv/opencv/assignees{/user}","branches_url":"https://api.github.com/repos/opencv/opencv/branches{/branch}","tags_url":"https://api.github.com/repos/opencv/opencv/tags","blobs_url":"https://api.github.com/repos/opencv/opencv/git/blobs{/sha}","git_tags_url":"https://api.github.com/repos/opencv/opencv/git/tags{/sha}","git_refs_url":"https://api.github.com/repos/opencv/opencv/git/refs{/sha}","trees_url":"https://api.github.com/repos/opencv/opencv/git/trees{/sha}","statuses_url":"https://api.github.com/repos/opencv/opencv/statuses/{sha}","languages_url":"https://api.github.com/repos/opencv/opencv/languages","stargazers_url":"https://api.github.com/repos/opencv/opencv/stargazers","contributors_url":"https://api.github.com/repos/opencv/opencv/contributors","subscribers_url":"https://api.github.com/repos/opencv/opencv/subscribers","subscription_url":"https://api.github.com/repos/opencv/opencv/subscription","commits_url":"https://api.github.com/repos/opencv/opencv/commits{/sha}","git_commits_url":"https://api.github.com/repos/opencv/opencv/git/commits{/sha}","comments_url":"https://api.github.com/repos/opencv/opencv/comments{/number}","issue_comment_url":"https://api.github.com/repos/opencv/opencv/issues/comments{/number}","contents_url":"https://api.github.com/repos/opencv/opencv/contents/{+path}","compare_url":"https://api.github.com/repos/opencv/opencv/compare/{base}...{head}","merges_url":"https://api.github.com/repos/opencv/opencv/merges","archive_url":"https://api.github.com/repos/opencv/opencv/{archive_format}{/ref}","downloads_url":"https://api.github.com/repos/opencv/opencv/downloads","issues_url":"https://api.github.com/repos/opencv/opencv/issues{/number}","pulls_url":"https://api.github.com/repos/opencv/opencv/pulls{/number}","milestones_url":"https://api.github.com/repos/opencv/opencv/milestones{/number}","notifications_url":"https://api.github.com/repos/opencv/opencv/notifications{?since,all,participating}","labels_url":"https://api.github.com/repos/opencv/opencv/labels{/name}","releases_url":"https://api.github.com/repos/opencv/opencv/releases{/id}","deployments_url":"https://api.github.com/repos/opencv/opencv/deployments","created_at":"2012-07-19T09:40:17Z","updated_at":"2017-09-06T03:28:10Z","pushed_at":"2017-09-06T05:29:09Z","git_url":"git://github.com/opencv/opencv.git","ssh_url":"git@github.com:opencv/opencv.git","clone_url":"https://github.com/opencv/opencv.git","svn_url":"https://github.com/opencv/opencv","homepage":"opencv.org","size":440824,"stargazers_count":17986,"watchers_count":17986,"language":"C++","has_issues":true,"has_projects":false,"has_downloads":true,"has_wiki":true,"has_pages":false,"forks_count":13435,"mirror_url":null,"open_issues_count":1274,"forks":13435,"open_issues":1274,"watchers":17986,"default_branch":"master"}},"_links":{"self":{"href":"https://api.github.com/repos/opencv/opencv/pulls/9114"},"html":{"href":"https://github.com/opencv/opencv/pull/9114"},"issue":{"href":"https://api.github.com/repos/opencv/opencv/issues/9114"},"comments":{"href":"https://api.github.com/repos/opencv/opencv/issues/9114/comments"},"review_comments":{"href":"https://api.github.com/repos/opencv/opencv/pulls/9114/comments"},"review_comment":{"href":"https://api.github.com/repos/opencv/opencv/pulls/comments{/number}"},"commits":{"href":"https://api.github.com/repos/opencv/opencv/pulls/9114/commits"},"statuses":{"href":"https://api.github.com/repos/opencv/opencv/statuses/d50b01714f6fdfd9c69a01a815414dc13e510584"}},"author_association":"MEMBER"}}
{"actor":{"display_login":"pengli"}}