diff --git a/modules/dnn/src/ocl4dnn/src/math_functions.cpp b/modules/dnn/src/ocl4dnn/src/math_functions.cpp index 42b35572aa..5fe52ac1ba 100644 --- 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(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 index af58cad833..65d6b4db7e 100644 --- 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::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 index 6cc65b7189..476d05287f 100644 --- a/modules/dnn/src/ocl4dnn/src/ocl4dnn_lrn.cpp +++ b/modules/dnn/src/ocl4dnn/src/ocl4dnn_lrn.cpp @@ -92,7 +92,6 @@ bool OCL4DNNLRN::Forward(const UMat& bottom, UMat& top) template bool OCL4DNNLRN::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 index e0bdf71e67..fe8b84b394 100644 --- a/modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp +++ b/modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp @@ -97,7 +97,6 @@ bool OCL4DNNPool::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 index 9ac5ddc8cc..f2452ff654 100644 --- a/modules/dnn/src/ocl4dnn/src/ocl4dnn_softmax.cpp +++ b/modules/dnn/src/ocl4dnn/src/ocl4dnn_softmax.cpp @@ -83,7 +83,6 @@ template bool OCL4DNNSoftmax::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 index 8b892b5e7c..de827d70ee 100644 --- 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.