cleanup ocl4dnn spatial convolution kernels

wzw-intel · pli2-intel · commit 04edc8fe3a88 · 2017-11-24T02:19:10.000+08:00
remove unused macros and half definition macros,
also remove unused ocl::Queue

Signed-off-by: Li Peng &lt;peng.li@intel.com&gt;
diff --git a/modules/dnn/src/ocl4dnn/src/math_functions.cpp b/modules/dnn/src/ocl4dnn/src/math_functions.cpp
@@ -65,8 +65,6 @@ ocl::Image2D ocl4dnnGEMMCopyBufferToImage(UMat buffer, int offset,
                                           int padded_width, int height,
                                           int width, int ld)
 {
-    ocl::Context ctx = ocl::Context::getDefault();
-    ocl::Queue queue = ocl::Queue::getDefault();
     ocl::Image2D image;
 
     if (!is_matrix_a && transpose)
@@ -192,9 +190,6 @@ static bool ocl4dnnFastImageGEMM(const CBLAS_TRANSPOSE TransA,
     // just padding one line is enough as the sub group block read
     // will clamp to edge according to the spec.
 
-    ocl::Context ctx = ocl::Context::getDefault();
-    ocl::Queue queue = ocl::Queue::getDefault();
-
     ocl::Image2D ImA;
     ocl::Image2D ImB;
 
@@ -446,7 +441,6 @@ bool ocl4dnnGEMV<float>(const CBLAS_TRANSPOSE TransA,
                  const int32_t offx, const float beta, UMat y,
                  const int32_t offy)
 {
-    ocl::Queue queue = ocl::Queue::getDefault();
     bool ret = false;
 
     if (TransA == CblasNoTrans)
@@ -507,8 +501,6 @@ bool ocl4dnnAXPY(const int32_t N, const Dtype alpha,
                  const UMat X, const int32_t offX, UMat Y,
                  const int32_t offY)
 {
-    ocl::Context ctx = ocl::Context::getDefault();
-
     ocl::Kernel oclk_axpy(CL_KERNEL_SELECT("axpy"), cv::ocl::dnn::math_oclsrc);
     if (oclk_axpy.empty())
         return false;
diff --git a/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp b/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp
@@ -184,8 +184,6 @@ void OCL4DNNConvSpatial<Dtype>::collectCommonInformation()
     addDef("as_Dtype2", "as_float2");
     addDef("as_Dtype4", "as_float4");
     addDef("as_Dtype8", "as_float8");
-    addDef("Dtype_ID", (int)CV_32F);
-    addDef("Dtype_SIZE", (int)sizeof(Dtype));
 }
 
 typedef enum {
diff --git a/modules/dnn/src/ocl4dnn/src/ocl4dnn_lrn.cpp b/modules/dnn/src/ocl4dnn/src/ocl4dnn_lrn.cpp
@@ -92,7 +92,6 @@ bool OCL4DNNLRN<Dtype>::Forward(const UMat& bottom, UMat& top)
 template<typename Dtype>
 bool OCL4DNNLRN<Dtype>::crossChannelForward(const UMat& bottom, UMat& top)
 {
-    ocl::Queue queue = ocl::Queue::getDefault();
     CHECK_EQ(phase_test_, true) << "Only support forward inference.";
 
     cl_uint argIdx = 0;
diff --git a/modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp b/modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp
@@ -97,7 +97,6 @@ bool OCL4DNNPool<Dtype>::Forward(const UMat& bottom,
                                  UMat& top_mask)
 {
     bool ret = true;
-    ocl::Queue queue = ocl::Queue::getDefault();
     size_t global[] = { 128 * 128 };
     size_t local[] = { 128 };
     cl_uint argIdx = 0;
diff --git a/modules/dnn/src/ocl4dnn/src/ocl4dnn_softmax.cpp b/modules/dnn/src/ocl4dnn/src/ocl4dnn_softmax.cpp
@@ -83,7 +83,6 @@ template<typename Dtype>
 bool OCL4DNNSoftmax<Dtype>::Forward(const UMat& bottom, UMat& top)
 {
     bool ret = false;
-    ocl::Queue queue = ocl::Queue::getDefault();
     bool intel_subgroup = ocl::Device::getDefault().intelSubgroupsSupport();
     if (intel_subgroup && inner_num_ < 128)
     {
diff --git a/modules/dnn/src/opencl/conv_layer_spatial.cl b/modules/dnn/src/opencl/conv_layer_spatial.cl
@@ -82,7 +82,6 @@
 #define LOOP(N, VAR, STMT) CAT(LOOP, N)((VAR), (STMT))
 
 #if defined(convolve_simd) || defined(Conv_Interleaved)
-#if Dtype_SIZE == 4
 #define INT_TYPE uint
 #define INT_TYPE2 uint2
 #define INT_TYPE4 uint4
@@ -91,9 +90,6 @@
 #define SUB_GROUP_BLOCK_READ4 intel_sub_group_block_read4
 #define SUB_GROUP_BLOCK_READ8 intel_sub_group_block_read8
 #define SUB_GROUP_BLOCK_READ intel_sub_group_block_read
-#else
-#error "Unsupported type"
-#endif
 #endif
 
 #ifdef KERNEL_BASIC
@@ -176,11 +172,7 @@ __kernel void ConvolveBasic(
 
 #elif defined KERNEL_IDLF
 
-#if TYPE == TYPE_HALF
-#define VLOAD4(_v, _p) do { (_v).s0 = *(_p); (_v).s1 = *(_p + 1); (_v).s2 = *(_p + 2); (_v).s3 = *(_p + 3); } while(0)
-#else
 #define VLOAD4(_v, _p) do { _v = vload4(0, _p); } while(0)
-#endif
 
 // Each work-item computes a OUT_BLOCK_WIDTH * OUT_BLOCK_HEIGHT region of one output map.
 // Each work-group (which will be mapped to 1 SIMD16/SIMD8 EU thread) will compute 16/8 different feature maps, but each feature map is for the same region of the imput image.

Original file line number	Diff line number	Diff line change
`@@ -184,8 +184,6 @@ void OCL4DNNConvSpatial<Dtype>::collectCommonInformation()`
`184`	`184`	`addDef("as_Dtype2", "as_float2");`
`185`	`185`	`addDef("as_Dtype4", "as_float4");`
`186`	`186`	`addDef("as_Dtype8", "as_float8");`
`187`		`- addDef("Dtype_ID", (int)CV_32F);`
`188`		`- addDef("Dtype_SIZE", (int)sizeof(Dtype));`
`189`	`187`	`}`
`190`	`188`
`191`	`189`	`typedef enum {`
Original file line number	Diff line number	Diff line change
`@@ -92,7 +92,6 @@ bool OCL4DNNLRN<Dtype>::Forward(const UMat& bottom, UMat& top)`
`92`	`92`	`template<typename Dtype>`
`93`	`93`	`bool OCL4DNNLRN<Dtype>::crossChannelForward(const UMat& bottom, UMat& top)`
`94`	`94`	`{`
`95`		`- ocl::Queue queue = ocl::Queue::getDefault();`
`96`	`95`	`CHECK_EQ(phase_test_, true) << "Only support forward inference.";`
`97`	`96`
`98`	`97`	`cl_uint argIdx = 0;`
Original file line number	Diff line number	Diff line change
`@@ -97,7 +97,6 @@ bool OCL4DNNPool<Dtype>::Forward(const UMat& bottom,`
`97`	`97`	`UMat& top_mask)`
`98`	`98`	`{`
`99`	`99`	`bool ret = true;`
`100`		`- ocl::Queue queue = ocl::Queue::getDefault();`
`101`	`100`	`size_t global[] = { 128 * 128 };`
`102`	`101`	`size_t local[] = { 128 };`
`103`	`102`	`cl_uint argIdx = 0;`
Original file line number	Diff line number	Diff line change
`@@ -83,7 +83,6 @@ template<typename Dtype>`
`83`	`83`	`bool OCL4DNNSoftmax<Dtype>::Forward(const UMat& bottom, UMat& top)`
`84`	`84`	`{`
`85`	`85`	`bool ret = false;`
`86`		`- ocl::Queue queue = ocl::Queue::getDefault();`
`87`	`86`	`bool intel_subgroup = ocl::Device::getDefault().intelSubgroupsSupport();`
`88`	`87`	`if (intel_subgroup && inner_num_ < 128)`
`89`	`88`	`{`