From c08f29c8038f4a9bb927130a8babd0570aa19f8d Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Sun, 27 Sep 2020 23:42:30 +0000 Subject: [PATCH] dnn(opencl): fix convolution kernel w/o bias with activation --- modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp | 3 +++ modules/dnn/src/opencl/conv_layer_spatial.cl | 2 +- modules/dnn/test/test_layers.cpp | 7 ------- 3 files changed, 4 insertions(+), 8 deletions(-) diff --git a/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp b/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp index 3cca52d5bd..3707a31846 100644 --- a/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp +++ b/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp @@ -607,6 +607,7 @@ void OCL4DNNConvSpatial::calculateBenchmark(const UMat &bottom, UMat &ver { options_.str(""); options_.clear(); // clear contents and state flags createBasicKernel(1, 1, 1); + CV_Assert(!kernelQueue.empty()); // basic kernel must be available kernel_index_ = kernelQueue.size() - 1; convolve(bottom, verifyTop, weight, bias, numImages, kernelQueue[kernel_index_]); CV_Assert(phash.find(kernelQueue[kernel_index_]->kernelName) != phash.end()); @@ -1713,6 +1714,7 @@ void OCL4DNNConvSpatial::useFirstAvailable(const UMat &bottom, tunerItems[i]->blockHeight, tunerItems[i]->blockDepth)) { + CV_Assert(!kernelQueue.empty()); // basic kernel must be available int kernelIdx = kernelQueue.size() - 1; kernelConfig* config = kernelQueue[kernelIdx].get(); bool failed = false; @@ -1883,6 +1885,7 @@ void OCL4DNNConvSpatial::setupConvolution(const UMat &bottom, CV_LOG_INFO(NULL, "fallback to basic kernel"); options_.str(""); options_.clear(); // clear contents and state flags createBasicKernel(1, 1, 1); + CV_Assert(!kernelQueue.empty()); // basic kernel must be available kernel_index_ = kernelQueue.size() - 1; } this->bestKernelConfig = kernelQueue[kernel_index_]; diff --git a/modules/dnn/src/opencl/conv_layer_spatial.cl b/modules/dnn/src/opencl/conv_layer_spatial.cl index 5d4d6f3add..236e8d029a 100644 --- a/modules/dnn/src/opencl/conv_layer_spatial.cl +++ b/modules/dnn/src/opencl/conv_layer_spatial.cl @@ -205,7 +205,7 @@ __kernel void ConvolveBasic( #if APPLY_BIAS ACTIVATION_FUNCTION(convolved_image, offset, sum[kern] + bias[biasIndex + kern], biasIndex + kern); #else - ACTIVATION_FUNCTION(convolved_image, offset, sum[kern], biasIndex + kern); + ACTIVATION_FUNCTION(convolved_image, offset, sum[kern], kernelNum + kern); #endif } } diff --git a/modules/dnn/test/test_layers.cpp b/modules/dnn/test/test_layers.cpp index f9b0f62379..085e5a51b8 100644 --- a/modules/dnn/test/test_layers.cpp +++ b/modules/dnn/test/test_layers.cpp @@ -2223,13 +2223,6 @@ TEST_P(ConvolutionActivationFusion, Accuracy) if (actType == "Power" && backendId == DNN_BACKEND_OPENCV && (targetId == DNN_TARGET_OPENCL || targetId == DNN_TARGET_OPENCL_FP16)) applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL); - // bug: https://github.com/opencv/opencv/issues/17953 - if (actType == "ChannelsPReLU" && bias_term == false && - backendId == DNN_BACKEND_OPENCV && (targetId == DNN_TARGET_OPENCL || targetId == DNN_TARGET_OPENCL_FP16)) - { - applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL); - } - Net net; int convId = net.addLayer(convParams.name, convParams.type, convParams); int activId = net.addLayerToPrev(activationParams.name, activationParams.type, activationParams);