opencv/modules/dnn/test/test_layers.cpp

2528 lines
85 KiB
C++
Raw Normal View History

/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2017, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "test_precomp.hpp"
#include <opencv2/core/ocl.hpp>
#include "npy_blob.hpp"
#include <opencv2/dnn/shape_utils.hpp>
#include <opencv2/dnn/all_layers.hpp>
#include <opencv2/dnn/layer.details.hpp> // CV_DNN_REGISTER_LAYER_CLASS
#ifdef HAVE_INF_ENGINE
#include <thread>
#endif
namespace opencv_test { namespace {
template<typename TString>
static String _tf(TString filename)
{
String basetestdir = getOpenCVExtraDir();
size_t len = basetestdir.size();
if(len > 0 && basetestdir[len-1] != '/' && basetestdir[len-1] != '\\')
return (basetestdir + "/dnn/layers") + filename;
return (basetestdir + "dnn/layers/") + filename;
}
void runLayer(Ptr<Layer> layer, std::vector<Mat> &inpBlobs, std::vector<Mat> &outBlobs)
{
size_t ninputs = inpBlobs.size();
std::vector<Mat> inp(ninputs), outp, intp;
std::vector<MatShape> inputs, outputs, internals;
for (size_t i = 0; i < ninputs; i++)
{
inp[i] = inpBlobs[i].clone();
inputs.push_back(shape(inp[i]));
}
layer->getMemoryShapes(inputs, 0, outputs, internals);
for (size_t i = 0; i < outputs.size(); i++)
{
outp.push_back(Mat(outputs[i], CV_32F));
}
for (size_t i = 0; i < internals.size(); i++)
{
intp.push_back(Mat(internals[i], CV_32F));
}
layer->finalize(inp, outp);
layer->forward(inp, outp, intp);
size_t noutputs = outp.size();
outBlobs.resize(noutputs);
for (size_t i = 0; i < noutputs; i++)
outBlobs[i] = outp[i];
}
2018-06-27 21:34:36 +08:00
class Test_Caffe_layers : public DNNTestLayer
{
2018-06-27 21:34:36 +08:00
public:
void testLayerUsingCaffeModels(const String& basename, bool useCaffeModel = false,
bool useCommonInputBlob = true, double l1 = 0.0, double lInf = 0.0,
int numInps = 1, int numOuts = 1)
2018-06-27 21:34:36 +08:00
{
CV_Assert_N(numInps >= 1, numInps <= 10, numOuts >= 1, numOuts <= 10);
2018-06-27 21:34:36 +08:00
String prototxt = _tf(basename + ".prototxt");
String caffemodel = _tf(basename + ".caffemodel");
std::vector<Mat> inps, refs, outs;
if (numInps > 1)
{
for (int i = 0; i < numInps; i++)
{
String inpfile = _tf(basename + cv::format(".input_%d.npy", i));
inps.push_back(blobFromNPY(inpfile));
}
}
else
{
String inpfile = (useCommonInputBlob) ? _tf("blob.npy") : _tf(basename + ".input.npy");
inps.push_back(blobFromNPY(inpfile));
}
if (numOuts > 1)
{
for (int i = 0; i < numOuts; i++)
{
String outfile = _tf(basename + cv::format("_%d.npy", i));
refs.push_back(blobFromNPY(outfile));
}
}
else
{
String outfile = _tf(basename + ".npy");
refs.push_back(blobFromNPY(outfile));
}
2018-06-27 21:34:36 +08:00
Net net = readNetFromCaffe(prototxt, (useCaffeModel) ? caffemodel : String());
ASSERT_FALSE(net.empty());
checkBackend(&inps[0], &refs[0]);
Merge pull request #9114 from pengli:dnn_rebase add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: Li Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: Li Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: Li Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: Li Peng <peng.li@intel.com> * set OPENCL target when it is available and disable fuseLayer for OCL target for the time being Signed-off-by: Li Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * remove trailing space Signed-off-by: Li Peng <peng.li@intel.com> * Fixed tensorflow demo bug. Root cause is that tensorflow has different algorithm with libdnn to calculate convolution output dimension. libdnn don't calculate output dimension anymore and just use one passed in by config. * split gemm ocl file split it into gemm_buffer.cl and gemm_image.cl Signed-off-by: Li Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: Li Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: Li Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: Li Peng <peng.li@intel.com> * update softmax layer on some platform subgroup extension may not work well, fallback to non subgroup ocl acceleration. Signed-off-by: Li Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: Li Peng <peng.li@intel.com> * update output message Signed-off-by: Li Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: Li Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: Li Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * DNN: Don't link OpenCL library explicitly * DNN: Make default preferableTarget to DNN_TARGET_CPU User should set it to DNN_TARGET_OPENCL explicitly if want to use OpenCL acceleration. Also don't fusion when using DNN_TARGET_OPENCL * DNN: refine coding style * Add getOpenCLErrorString * DNN: Use int32_t/uint32_t instread of alias * Use namespace ocl4dnn to include libdnn things * remove extra copyTo in softmax ocl path Signed-off-by: Li Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: Li Peng <peng.li@intel.com> * Add prefer target property for layer class It is used to indicate the target for layer forwarding, either the default CPU target or OCL target. Signed-off-by: Li Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: Li Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: wzw <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: Li Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: Li Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: Li Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: Li Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: Li Peng <peng.li@intel.com> * Refine auto-tuning mechanism. - Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory for fine-tuned kernel configuration. e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp, the cache directory will be /home/tmp/spatialkernels/ on Linux. - Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable auto-tuning. - OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling for OpenCL command queue. This fix basic kernel get wrong running time, i.e. 0ms. - If creating cache directory failed, disable auto-tuning. * Detect and create cache dir on windows Signed-off-by: Li Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: Li Peng <peng.li@intel.com> * Fix redundant swizzleWeights calling when use cached kernel config. * Fix "out of resource" bug when auto-tuning too many kernels. * replace cl_mem with UMat in ocl4dnnConvSpatial class * OCL4DNN: reduce the tuning kernel candidate. This patch could reduce 75% of the tuning candidates with less than 2% performance impact for the final result. Signed-off-by: Zhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: Li Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: Li Peng <peng.li@intel.com> * Various fixes for ocl4dnn 1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()) 2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp 3. Code comments cleanup 4. ignore check on OCL cpu device Signed-off-by: Li Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: Li Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: Li Peng <peng.li@intel.com> * check kernel build options Signed-off-by: Li Peng <peng.li@intel.com> * Handle program compilation fail properly. * Use std::numeric_limits<float>::infinity() for large float number * check ocl4dnn kernel compilation result Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: Li Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: Li Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: Li Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: Li Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: Li Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: Li Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: Li Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: Li Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: Li Peng <peng.li@intel.com> * dnn(build): don't use C++11, OPENCL_LIBRARIES fix * dnn(ocl4dnn): remove unused OpenCL kernels * dnn(ocl4dnn): extract OpenCL code into .cl files * dnn(ocl4dnn): refine auto-tuning Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING environment variable to enable it. Use a set of pre-tuned configs as default config if auto-tuning is disabled. These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet, AlexNet, ResNet-50 If default config is not suitable, use the first available kernel config from the candidates. Candidate priority from high to low is gemm like kernel, IDLF kernel, basick kernel. * dnn(ocl4dnn): pooling doesn't use OpenCL subgroups * dnn(ocl4dnn): fix perf test OpenCV has default 3sec time limit for each performance test. Warmup OpenCL backend outside of perf measurement loop. * use ocl::KernelArg as much as possible Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): Refine signature of kernel config - Use more readable string as signture of kernel config - Don't count device name and vendor in signature string - Default kernel configurations are tuned for Intel GPU with 24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model. * dnn(ocl4dnn): swap width/height in configuration * dnn(ocl4dnn): enable configs for Intel OpenCL runtime only * core: make configuration helper functions accessible from non-core modules * dnn(ocl4dnn): update kernel auto-tuning behavior Avoid unwanted creation of directories * dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash * dnn(ocl4dnn): remove redundant code * dnn(ocl4dnn): Add more clear message for simd size dismatch. * dnn(ocl4dnn): add const to const argument Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel * dnn(ocl4dnn): drop unused tuneLocalSize() * dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method * dnn(ocl4dnn): sanitize file names used for cache * dnn(perf): enable Network tests with OpenCL * dnn(ocl4dnn/conv): drop computeGlobalSize() * dnn(ocl4dnn/conv): drop unused fields * dnn(ocl4dnn/conv): simplify ctor * dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL * dnn(ocl4dnn/conv): drop unsupported double / untested half types * dnn(ocl4dnn/conv): drop unused variable * dnn(ocl4dnn/conv): alignSize/divUp * dnn(ocl4dnn/conv): use enum values * dnn(ocl4dnn): drop unused innerproduct variable Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): add an generic function to check cl option support * dnn(ocl4dnn): run softmax subgroup version kernel first Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
2018-06-27 21:34:36 +08:00
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
String inp_name = "input";
if (numInps > 1)
{
for (int i = 0; i < numInps; i++)
{
net.setInput(inps[i], inp_name + cv::format("_%d", i));
}
}
else
{
net.setInput(inps.back(), inp_name);
}
net.forward(outs);
for (int i = 0; i < refs.size(); i++)
{
normAssert(refs[i], outs[i], "", l1 ? l1 : default_l1, lInf ? lInf : default_lInf);
}
2018-06-27 21:34:36 +08:00
}
};
TEST_P(Test_Caffe_layers, Softmax)
{
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_softmax");
}
TEST_P(Test_Caffe_layers, LRN)
{
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_lrn_spatial");
testLayerUsingCaffeModels("layer_lrn_channels");
}
TEST_P(Test_Caffe_layers, Convolution)
{
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_convolution", true);
Merge pull request #9114 from pengli:dnn_rebase add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: Li Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: Li Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: Li Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: Li Peng <peng.li@intel.com> * set OPENCL target when it is available and disable fuseLayer for OCL target for the time being Signed-off-by: Li Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * remove trailing space Signed-off-by: Li Peng <peng.li@intel.com> * Fixed tensorflow demo bug. Root cause is that tensorflow has different algorithm with libdnn to calculate convolution output dimension. libdnn don't calculate output dimension anymore and just use one passed in by config. * split gemm ocl file split it into gemm_buffer.cl and gemm_image.cl Signed-off-by: Li Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: Li Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: Li Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: Li Peng <peng.li@intel.com> * update softmax layer on some platform subgroup extension may not work well, fallback to non subgroup ocl acceleration. Signed-off-by: Li Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: Li Peng <peng.li@intel.com> * update output message Signed-off-by: Li Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: Li Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: Li Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * DNN: Don't link OpenCL library explicitly * DNN: Make default preferableTarget to DNN_TARGET_CPU User should set it to DNN_TARGET_OPENCL explicitly if want to use OpenCL acceleration. Also don't fusion when using DNN_TARGET_OPENCL * DNN: refine coding style * Add getOpenCLErrorString * DNN: Use int32_t/uint32_t instread of alias * Use namespace ocl4dnn to include libdnn things * remove extra copyTo in softmax ocl path Signed-off-by: Li Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: Li Peng <peng.li@intel.com> * Add prefer target property for layer class It is used to indicate the target for layer forwarding, either the default CPU target or OCL target. Signed-off-by: Li Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: Li Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: wzw <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: Li Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: Li Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: Li Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: Li Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: Li Peng <peng.li@intel.com> * Refine auto-tuning mechanism. - Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory for fine-tuned kernel configuration. e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp, the cache directory will be /home/tmp/spatialkernels/ on Linux. - Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable auto-tuning. - OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling for OpenCL command queue. This fix basic kernel get wrong running time, i.e. 0ms. - If creating cache directory failed, disable auto-tuning. * Detect and create cache dir on windows Signed-off-by: Li Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: Li Peng <peng.li@intel.com> * Fix redundant swizzleWeights calling when use cached kernel config. * Fix "out of resource" bug when auto-tuning too many kernels. * replace cl_mem with UMat in ocl4dnnConvSpatial class * OCL4DNN: reduce the tuning kernel candidate. This patch could reduce 75% of the tuning candidates with less than 2% performance impact for the final result. Signed-off-by: Zhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: Li Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: Li Peng <peng.li@intel.com> * Various fixes for ocl4dnn 1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()) 2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp 3. Code comments cleanup 4. ignore check on OCL cpu device Signed-off-by: Li Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: Li Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: Li Peng <peng.li@intel.com> * check kernel build options Signed-off-by: Li Peng <peng.li@intel.com> * Handle program compilation fail properly. * Use std::numeric_limits<float>::infinity() for large float number * check ocl4dnn kernel compilation result Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: Li Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: Li Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: Li Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: Li Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: Li Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: Li Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: Li Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: Li Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: Li Peng <peng.li@intel.com> * dnn(build): don't use C++11, OPENCL_LIBRARIES fix * dnn(ocl4dnn): remove unused OpenCL kernels * dnn(ocl4dnn): extract OpenCL code into .cl files * dnn(ocl4dnn): refine auto-tuning Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING environment variable to enable it. Use a set of pre-tuned configs as default config if auto-tuning is disabled. These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet, AlexNet, ResNet-50 If default config is not suitable, use the first available kernel config from the candidates. Candidate priority from high to low is gemm like kernel, IDLF kernel, basick kernel. * dnn(ocl4dnn): pooling doesn't use OpenCL subgroups * dnn(ocl4dnn): fix perf test OpenCV has default 3sec time limit for each performance test. Warmup OpenCL backend outside of perf measurement loop. * use ocl::KernelArg as much as possible Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): Refine signature of kernel config - Use more readable string as signture of kernel config - Don't count device name and vendor in signature string - Default kernel configurations are tuned for Intel GPU with 24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model. * dnn(ocl4dnn): swap width/height in configuration * dnn(ocl4dnn): enable configs for Intel OpenCL runtime only * core: make configuration helper functions accessible from non-core modules * dnn(ocl4dnn): update kernel auto-tuning behavior Avoid unwanted creation of directories * dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash * dnn(ocl4dnn): remove redundant code * dnn(ocl4dnn): Add more clear message for simd size dismatch. * dnn(ocl4dnn): add const to const argument Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel * dnn(ocl4dnn): drop unused tuneLocalSize() * dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method * dnn(ocl4dnn): sanitize file names used for cache * dnn(perf): enable Network tests with OpenCL * dnn(ocl4dnn/conv): drop computeGlobalSize() * dnn(ocl4dnn/conv): drop unused fields * dnn(ocl4dnn/conv): simplify ctor * dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL * dnn(ocl4dnn/conv): drop unsupported double / untested half types * dnn(ocl4dnn/conv): drop unused variable * dnn(ocl4dnn/conv): alignSize/divUp * dnn(ocl4dnn/conv): use enum values * dnn(ocl4dnn): drop unused innerproduct variable Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): add an generic function to check cl option support * dnn(ocl4dnn): run softmax subgroup version kernel first Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
}
TEST_P(Test_Caffe_layers, DeConvolution)
Merge pull request #9114 from pengli:dnn_rebase add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: Li Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: Li Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: Li Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: Li Peng <peng.li@intel.com> * set OPENCL target when it is available and disable fuseLayer for OCL target for the time being Signed-off-by: Li Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * remove trailing space Signed-off-by: Li Peng <peng.li@intel.com> * Fixed tensorflow demo bug. Root cause is that tensorflow has different algorithm with libdnn to calculate convolution output dimension. libdnn don't calculate output dimension anymore and just use one passed in by config. * split gemm ocl file split it into gemm_buffer.cl and gemm_image.cl Signed-off-by: Li Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: Li Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: Li Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: Li Peng <peng.li@intel.com> * update softmax layer on some platform subgroup extension may not work well, fallback to non subgroup ocl acceleration. Signed-off-by: Li Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: Li Peng <peng.li@intel.com> * update output message Signed-off-by: Li Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: Li Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: Li Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * DNN: Don't link OpenCL library explicitly * DNN: Make default preferableTarget to DNN_TARGET_CPU User should set it to DNN_TARGET_OPENCL explicitly if want to use OpenCL acceleration. Also don't fusion when using DNN_TARGET_OPENCL * DNN: refine coding style * Add getOpenCLErrorString * DNN: Use int32_t/uint32_t instread of alias * Use namespace ocl4dnn to include libdnn things * remove extra copyTo in softmax ocl path Signed-off-by: Li Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: Li Peng <peng.li@intel.com> * Add prefer target property for layer class It is used to indicate the target for layer forwarding, either the default CPU target or OCL target. Signed-off-by: Li Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: Li Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: wzw <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: Li Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: Li Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: Li Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: Li Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: Li Peng <peng.li@intel.com> * Refine auto-tuning mechanism. - Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory for fine-tuned kernel configuration. e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp, the cache directory will be /home/tmp/spatialkernels/ on Linux. - Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable auto-tuning. - OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling for OpenCL command queue. This fix basic kernel get wrong running time, i.e. 0ms. - If creating cache directory failed, disable auto-tuning. * Detect and create cache dir on windows Signed-off-by: Li Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: Li Peng <peng.li@intel.com> * Fix redundant swizzleWeights calling when use cached kernel config. * Fix "out of resource" bug when auto-tuning too many kernels. * replace cl_mem with UMat in ocl4dnnConvSpatial class * OCL4DNN: reduce the tuning kernel candidate. This patch could reduce 75% of the tuning candidates with less than 2% performance impact for the final result. Signed-off-by: Zhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: Li Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: Li Peng <peng.li@intel.com> * Various fixes for ocl4dnn 1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()) 2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp 3. Code comments cleanup 4. ignore check on OCL cpu device Signed-off-by: Li Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: Li Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: Li Peng <peng.li@intel.com> * check kernel build options Signed-off-by: Li Peng <peng.li@intel.com> * Handle program compilation fail properly. * Use std::numeric_limits<float>::infinity() for large float number * check ocl4dnn kernel compilation result Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: Li Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: Li Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: Li Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: Li Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: Li Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: Li Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: Li Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: Li Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: Li Peng <peng.li@intel.com> * dnn(build): don't use C++11, OPENCL_LIBRARIES fix * dnn(ocl4dnn): remove unused OpenCL kernels * dnn(ocl4dnn): extract OpenCL code into .cl files * dnn(ocl4dnn): refine auto-tuning Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING environment variable to enable it. Use a set of pre-tuned configs as default config if auto-tuning is disabled. These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet, AlexNet, ResNet-50 If default config is not suitable, use the first available kernel config from the candidates. Candidate priority from high to low is gemm like kernel, IDLF kernel, basick kernel. * dnn(ocl4dnn): pooling doesn't use OpenCL subgroups * dnn(ocl4dnn): fix perf test OpenCV has default 3sec time limit for each performance test. Warmup OpenCL backend outside of perf measurement loop. * use ocl::KernelArg as much as possible Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): Refine signature of kernel config - Use more readable string as signture of kernel config - Don't count device name and vendor in signature string - Default kernel configurations are tuned for Intel GPU with 24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model. * dnn(ocl4dnn): swap width/height in configuration * dnn(ocl4dnn): enable configs for Intel OpenCL runtime only * core: make configuration helper functions accessible from non-core modules * dnn(ocl4dnn): update kernel auto-tuning behavior Avoid unwanted creation of directories * dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash * dnn(ocl4dnn): remove redundant code * dnn(ocl4dnn): Add more clear message for simd size dismatch. * dnn(ocl4dnn): add const to const argument Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel * dnn(ocl4dnn): drop unused tuneLocalSize() * dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method * dnn(ocl4dnn): sanitize file names used for cache * dnn(perf): enable Network tests with OpenCL * dnn(ocl4dnn/conv): drop computeGlobalSize() * dnn(ocl4dnn/conv): drop unused fields * dnn(ocl4dnn/conv): simplify ctor * dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL * dnn(ocl4dnn/conv): drop unsupported double / untested half types * dnn(ocl4dnn/conv): drop unused variable * dnn(ocl4dnn/conv): alignSize/divUp * dnn(ocl4dnn/conv): use enum values * dnn(ocl4dnn): drop unused innerproduct variable Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): add an generic function to check cl option support * dnn(ocl4dnn): run softmax subgroup version kernel first Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
{
if(target == DNN_TARGET_CUDA_FP16)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA_FP16);
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_deconvolution", true, false);
}
TEST_P(Test_Caffe_layers, InnerProduct)
{
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16)
applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL_FP16);
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_inner_product", true);
Merge pull request #9114 from pengli:dnn_rebase add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: Li Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: Li Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: Li Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: Li Peng <peng.li@intel.com> * set OPENCL target when it is available and disable fuseLayer for OCL target for the time being Signed-off-by: Li Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * remove trailing space Signed-off-by: Li Peng <peng.li@intel.com> * Fixed tensorflow demo bug. Root cause is that tensorflow has different algorithm with libdnn to calculate convolution output dimension. libdnn don't calculate output dimension anymore and just use one passed in by config. * split gemm ocl file split it into gemm_buffer.cl and gemm_image.cl Signed-off-by: Li Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: Li Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: Li Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: Li Peng <peng.li@intel.com> * update softmax layer on some platform subgroup extension may not work well, fallback to non subgroup ocl acceleration. Signed-off-by: Li Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: Li Peng <peng.li@intel.com> * update output message Signed-off-by: Li Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: Li Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: Li Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * DNN: Don't link OpenCL library explicitly * DNN: Make default preferableTarget to DNN_TARGET_CPU User should set it to DNN_TARGET_OPENCL explicitly if want to use OpenCL acceleration. Also don't fusion when using DNN_TARGET_OPENCL * DNN: refine coding style * Add getOpenCLErrorString * DNN: Use int32_t/uint32_t instread of alias * Use namespace ocl4dnn to include libdnn things * remove extra copyTo in softmax ocl path Signed-off-by: Li Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: Li Peng <peng.li@intel.com> * Add prefer target property for layer class It is used to indicate the target for layer forwarding, either the default CPU target or OCL target. Signed-off-by: Li Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: Li Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: wzw <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: Li Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: Li Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: Li Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: Li Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: Li Peng <peng.li@intel.com> * Refine auto-tuning mechanism. - Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory for fine-tuned kernel configuration. e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp, the cache directory will be /home/tmp/spatialkernels/ on Linux. - Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable auto-tuning. - OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling for OpenCL command queue. This fix basic kernel get wrong running time, i.e. 0ms. - If creating cache directory failed, disable auto-tuning. * Detect and create cache dir on windows Signed-off-by: Li Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: Li Peng <peng.li@intel.com> * Fix redundant swizzleWeights calling when use cached kernel config. * Fix "out of resource" bug when auto-tuning too many kernels. * replace cl_mem with UMat in ocl4dnnConvSpatial class * OCL4DNN: reduce the tuning kernel candidate. This patch could reduce 75% of the tuning candidates with less than 2% performance impact for the final result. Signed-off-by: Zhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: Li Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: Li Peng <peng.li@intel.com> * Various fixes for ocl4dnn 1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()) 2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp 3. Code comments cleanup 4. ignore check on OCL cpu device Signed-off-by: Li Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: Li Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: Li Peng <peng.li@intel.com> * check kernel build options Signed-off-by: Li Peng <peng.li@intel.com> * Handle program compilation fail properly. * Use std::numeric_limits<float>::infinity() for large float number * check ocl4dnn kernel compilation result Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: Li Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: Li Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: Li Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: Li Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: Li Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: Li Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: Li Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: Li Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: Li Peng <peng.li@intel.com> * dnn(build): don't use C++11, OPENCL_LIBRARIES fix * dnn(ocl4dnn): remove unused OpenCL kernels * dnn(ocl4dnn): extract OpenCL code into .cl files * dnn(ocl4dnn): refine auto-tuning Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING environment variable to enable it. Use a set of pre-tuned configs as default config if auto-tuning is disabled. These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet, AlexNet, ResNet-50 If default config is not suitable, use the first available kernel config from the candidates. Candidate priority from high to low is gemm like kernel, IDLF kernel, basick kernel. * dnn(ocl4dnn): pooling doesn't use OpenCL subgroups * dnn(ocl4dnn): fix perf test OpenCV has default 3sec time limit for each performance test. Warmup OpenCL backend outside of perf measurement loop. * use ocl::KernelArg as much as possible Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): Refine signature of kernel config - Use more readable string as signture of kernel config - Don't count device name and vendor in signature string - Default kernel configurations are tuned for Intel GPU with 24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model. * dnn(ocl4dnn): swap width/height in configuration * dnn(ocl4dnn): enable configs for Intel OpenCL runtime only * core: make configuration helper functions accessible from non-core modules * dnn(ocl4dnn): update kernel auto-tuning behavior Avoid unwanted creation of directories * dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash * dnn(ocl4dnn): remove redundant code * dnn(ocl4dnn): Add more clear message for simd size dismatch. * dnn(ocl4dnn): add const to const argument Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel * dnn(ocl4dnn): drop unused tuneLocalSize() * dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method * dnn(ocl4dnn): sanitize file names used for cache * dnn(perf): enable Network tests with OpenCL * dnn(ocl4dnn/conv): drop computeGlobalSize() * dnn(ocl4dnn/conv): drop unused fields * dnn(ocl4dnn/conv): simplify ctor * dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL * dnn(ocl4dnn/conv): drop unsupported double / untested half types * dnn(ocl4dnn/conv): drop unused variable * dnn(ocl4dnn/conv): alignSize/divUp * dnn(ocl4dnn/conv): use enum values * dnn(ocl4dnn): drop unused innerproduct variable Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): add an generic function to check cl option support * dnn(ocl4dnn): run softmax subgroup version kernel first Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
}
TEST_P(Test_Caffe_layers, Pooling_max)
Merge pull request #9114 from pengli:dnn_rebase add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: Li Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: Li Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: Li Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: Li Peng <peng.li@intel.com> * set OPENCL target when it is available and disable fuseLayer for OCL target for the time being Signed-off-by: Li Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * remove trailing space Signed-off-by: Li Peng <peng.li@intel.com> * Fixed tensorflow demo bug. Root cause is that tensorflow has different algorithm with libdnn to calculate convolution output dimension. libdnn don't calculate output dimension anymore and just use one passed in by config. * split gemm ocl file split it into gemm_buffer.cl and gemm_image.cl Signed-off-by: Li Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: Li Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: Li Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: Li Peng <peng.li@intel.com> * update softmax layer on some platform subgroup extension may not work well, fallback to non subgroup ocl acceleration. Signed-off-by: Li Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: Li Peng <peng.li@intel.com> * update output message Signed-off-by: Li Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: Li Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: Li Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * DNN: Don't link OpenCL library explicitly * DNN: Make default preferableTarget to DNN_TARGET_CPU User should set it to DNN_TARGET_OPENCL explicitly if want to use OpenCL acceleration. Also don't fusion when using DNN_TARGET_OPENCL * DNN: refine coding style * Add getOpenCLErrorString * DNN: Use int32_t/uint32_t instread of alias * Use namespace ocl4dnn to include libdnn things * remove extra copyTo in softmax ocl path Signed-off-by: Li Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: Li Peng <peng.li@intel.com> * Add prefer target property for layer class It is used to indicate the target for layer forwarding, either the default CPU target or OCL target. Signed-off-by: Li Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: Li Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: wzw <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: Li Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: Li Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: Li Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: Li Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: Li Peng <peng.li@intel.com> * Refine auto-tuning mechanism. - Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory for fine-tuned kernel configuration. e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp, the cache directory will be /home/tmp/spatialkernels/ on Linux. - Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable auto-tuning. - OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling for OpenCL command queue. This fix basic kernel get wrong running time, i.e. 0ms. - If creating cache directory failed, disable auto-tuning. * Detect and create cache dir on windows Signed-off-by: Li Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: Li Peng <peng.li@intel.com> * Fix redundant swizzleWeights calling when use cached kernel config. * Fix "out of resource" bug when auto-tuning too many kernels. * replace cl_mem with UMat in ocl4dnnConvSpatial class * OCL4DNN: reduce the tuning kernel candidate. This patch could reduce 75% of the tuning candidates with less than 2% performance impact for the final result. Signed-off-by: Zhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: Li Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: Li Peng <peng.li@intel.com> * Various fixes for ocl4dnn 1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()) 2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp 3. Code comments cleanup 4. ignore check on OCL cpu device Signed-off-by: Li Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: Li Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: Li Peng <peng.li@intel.com> * check kernel build options Signed-off-by: Li Peng <peng.li@intel.com> * Handle program compilation fail properly. * Use std::numeric_limits<float>::infinity() for large float number * check ocl4dnn kernel compilation result Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: Li Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: Li Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: Li Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: Li Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: Li Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: Li Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: Li Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: Li Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: Li Peng <peng.li@intel.com> * dnn(build): don't use C++11, OPENCL_LIBRARIES fix * dnn(ocl4dnn): remove unused OpenCL kernels * dnn(ocl4dnn): extract OpenCL code into .cl files * dnn(ocl4dnn): refine auto-tuning Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING environment variable to enable it. Use a set of pre-tuned configs as default config if auto-tuning is disabled. These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet, AlexNet, ResNet-50 If default config is not suitable, use the first available kernel config from the candidates. Candidate priority from high to low is gemm like kernel, IDLF kernel, basick kernel. * dnn(ocl4dnn): pooling doesn't use OpenCL subgroups * dnn(ocl4dnn): fix perf test OpenCV has default 3sec time limit for each performance test. Warmup OpenCL backend outside of perf measurement loop. * use ocl::KernelArg as much as possible Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): Refine signature of kernel config - Use more readable string as signture of kernel config - Don't count device name and vendor in signature string - Default kernel configurations are tuned for Intel GPU with 24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model. * dnn(ocl4dnn): swap width/height in configuration * dnn(ocl4dnn): enable configs for Intel OpenCL runtime only * core: make configuration helper functions accessible from non-core modules * dnn(ocl4dnn): update kernel auto-tuning behavior Avoid unwanted creation of directories * dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash * dnn(ocl4dnn): remove redundant code * dnn(ocl4dnn): Add more clear message for simd size dismatch. * dnn(ocl4dnn): add const to const argument Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel * dnn(ocl4dnn): drop unused tuneLocalSize() * dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method * dnn(ocl4dnn): sanitize file names used for cache * dnn(perf): enable Network tests with OpenCL * dnn(ocl4dnn/conv): drop computeGlobalSize() * dnn(ocl4dnn/conv): drop unused fields * dnn(ocl4dnn/conv): simplify ctor * dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL * dnn(ocl4dnn/conv): drop unsupported double / untested half types * dnn(ocl4dnn/conv): drop unused variable * dnn(ocl4dnn/conv): alignSize/divUp * dnn(ocl4dnn/conv): use enum values * dnn(ocl4dnn): drop unused innerproduct variable Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): add an generic function to check cl option support * dnn(ocl4dnn): run softmax subgroup version kernel first Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
{
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_pooling_max");
}
TEST_P(Test_Caffe_layers, Pooling_ave)
{
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_pooling_ave");
}
TEST_P(Test_Caffe_layers, MVN)
{
Merge pull request #14827 from YashasSamaga:cuda4dnn-csl-low CUDA backend for the DNN module * stub cuda4dnn design * minor fixes for tests and doxygen * add csl public api directory to module headers * add low-level CSL components * add high-level CSL components * integrate csl::Tensor into backbone code * switch to CPU iff unsupported; otherwise, fail on error * add fully connected layer * add softmax layer * add activation layers * support arbitary rank TensorDescriptor * pass input wrappers to `initCUDA()` * add 1d/2d/3d-convolution * add pooling layer * reorganize and refactor code * fixes for gcc, clang and doxygen; remove cxx14/17 code * add blank_layer * add LRN layer * add rounding modes for pooling layer * split tensor.hpp into tensor.hpp and tensor_ops.hpp * add concat layer * add scale layer * add batch normalization layer * split math.cu into activations.cu and math.hpp * add eltwise layer * add flatten layer * add tensor transform api * add asymmetric padding support for convolution layer * add reshape layer * fix rebase issues * add permute layer * add padding support for concat layer * refactor and reorganize code * add normalize layer * optimize bias addition in scale layer * add prior box layer * fix and optimize normalize layer * add asymmetric padding support for pooling layer * add event API * improve pooling performance for some padding scenarios * avoid over-allocation of compute resources to kernels * improve prior box performance * enable layer fusion * add const layer * add resize layer * add slice layer * add padding layer * add deconvolution layer * fix channelwise ReLU initialization * add vector traits * add vectorized versions of relu, clipped_relu, power * add vectorized concat kernels * improve concat_with_offsets performance * vectorize scale and bias kernels * add support for multi-billion element tensors * vectorize prior box kernels * fix address alignment check * improve bias addition performance of conv/deconv/fc layers * restructure code for supporting multiple targets * add DNN_TARGET_CUDA_FP64 * add DNN_TARGET_FP16 * improve vectorization * add region layer * improve tensor API, add dynamic ranks 1. use ManagedPtr instead of a Tensor in backend wrapper 2. add new methods to tensor classes - size_range: computes the combined size of for a given axis range - tensor span/view can be constructed from a raw pointer and shape 3. the tensor classes can change their rank at runtime (previously rank was fixed at compile-time) 4. remove device code from tensor classes (as they are unused) 5. enforce strict conditions on tensor class APIs to improve debugging ability * fix parametric relu activation * add squeeze/unsqueeze tensor API * add reorg layer * optimize permute and enable 2d permute * enable 1d and 2d slice * add split layer * add shuffle channel layer * allow tensors of different ranks in reshape primitive * patch SliceOp to allow Crop Layer * allow extra shape inputs in reshape layer * use `std::move_backward` instead of `std::move` for insert in resizable_static_array * improve workspace management * add spatial LRN * add nms (cpu) to region layer * add max pooling with argmax ( and a fix to limits.hpp) * add max unpooling layer * rename DNN_TARGET_CUDA_FP32 to DNN_TARGET_CUDA * update supportBackend to be more rigorous * remove stray include from preventing non-cuda build * include op_cuda.hpp outside condition #if * refactoring, fixes and many optimizations * drop DNN_TARGET_CUDA_FP64 * fix gcc errors * increase max. tensor rank limit to six * add Interp layer * drop custom layers; use BackendNode * vectorize activation kernels * fixes for gcc * remove wrong assertion * fix broken assertion in unpooling primitive * fix build errors in non-CUDA build * completely remove workspace from public API * fix permute layer * enable accuracy and perf. tests for DNN_TARGET_CUDA * add asynchronous forward * vectorize eltwise ops * vectorize fill kernel * fixes for gcc * remove CSL headers from public API * remove csl header source group from cmake * update min. cudnn version in cmake * add numerically stable FP32 log1pexp * refactor code * add FP16 specialization to cudnn based tensor addition * vectorize scale1 and bias1 + minor refactoring * fix doxygen build * fix invalid alignment assertion * clear backend wrappers before allocateLayers * ignore memory lock failures * do not allocate internal blobs * integrate NVTX * add numerically stable half precision log1pexp * fix indentation, following coding style, improve docs * remove accidental modification of IE code * Revert "add asynchronous forward" This reverts commit 1154b9da9da07e9b52f8a81bdcea48cf31c56f70. * [cmake] throw error for unsupported CC versions * fix rebase issues * add more docs, refactor code, fix bugs * minor refactoring and fixes * resolve warnings/errors from clang * remove haveCUDA() checks from supportBackend() * remove NVTX integration * changes based on review comments * avoid exception when no CUDA device is present * add color code for CUDA in Net::dump
2019-10-21 19:28:00 +08:00
if(backend == DNN_BACKEND_CUDA)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA); /* MVN is unsupported */
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_mvn");
}
void testReshape(const MatShape& inputShape, const MatShape& targetShape,
int axis = 0, int num_axes = -1,
MatShape mask = MatShape())
{
LayerParams params;
params.set("axis", axis);
params.set("num_axes", num_axes);
if (!mask.empty())
{
params.set("dim", DictValue::arrayInt<int*>(&mask[0], mask.size()));
}
Mat inp(inputShape.size(), &inputShape[0], CV_32F);
std::vector<Mat> inpVec(1, inp);
std::vector<Mat> outVec, intVec;
Ptr<Layer> rl = LayerFactory::createLayerInstance("Reshape", params);
runLayer(rl, inpVec, outVec);
Mat& out = outVec[0];
MatShape shape(out.size.p, out.size.p + out.dims);
EXPECT_EQ(shape, targetShape);
}
TEST(Layer_Test_Reshape, Accuracy)
{
{
int inp[] = {4, 3, 1, 2};
int out[] = {4, 3, 2};
testReshape(MatShape(inp, inp + 4), MatShape(out, out + 3), 2, 1);
}
{
int inp[] = {1, 128, 4, 4};
int out[] = {1, 2048};
int mask[] = {-1, 2048};
testReshape(MatShape(inp, inp + 4), MatShape(out, out + 2), 0, -1,
MatShape(mask, mask + 2));
}
2018-07-03 13:26:43 +08:00
{
int inp[] = {1, 2, 3};
int out[] = {3, 1, 2};
int mask[] = {3, 1, 2};
testReshape(MatShape(inp, inp + 3), MatShape(out, out + 3), 0, -1,
MatShape(mask, mask + 3));
}
}
2018-06-27 21:34:36 +08:00
TEST_P(Test_Caffe_layers, BatchNorm)
{
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_batch_norm", true);
testLayerUsingCaffeModels("layer_batch_norm_local_stats", true, false);
}
TEST_P(Test_Caffe_layers, ReLU)
Merge pull request #9114 from pengli:dnn_rebase add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: Li Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: Li Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: Li Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: Li Peng <peng.li@intel.com> * set OPENCL target when it is available and disable fuseLayer for OCL target for the time being Signed-off-by: Li Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * remove trailing space Signed-off-by: Li Peng <peng.li@intel.com> * Fixed tensorflow demo bug. Root cause is that tensorflow has different algorithm with libdnn to calculate convolution output dimension. libdnn don't calculate output dimension anymore and just use one passed in by config. * split gemm ocl file split it into gemm_buffer.cl and gemm_image.cl Signed-off-by: Li Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: Li Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: Li Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: Li Peng <peng.li@intel.com> * update softmax layer on some platform subgroup extension may not work well, fallback to non subgroup ocl acceleration. Signed-off-by: Li Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: Li Peng <peng.li@intel.com> * update output message Signed-off-by: Li Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: Li Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: Li Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * DNN: Don't link OpenCL library explicitly * DNN: Make default preferableTarget to DNN_TARGET_CPU User should set it to DNN_TARGET_OPENCL explicitly if want to use OpenCL acceleration. Also don't fusion when using DNN_TARGET_OPENCL * DNN: refine coding style * Add getOpenCLErrorString * DNN: Use int32_t/uint32_t instread of alias * Use namespace ocl4dnn to include libdnn things * remove extra copyTo in softmax ocl path Signed-off-by: Li Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: Li Peng <peng.li@intel.com> * Add prefer target property for layer class It is used to indicate the target for layer forwarding, either the default CPU target or OCL target. Signed-off-by: Li Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: Li Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: wzw <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: Li Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: Li Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: Li Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: Li Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: Li Peng <peng.li@intel.com> * Refine auto-tuning mechanism. - Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory for fine-tuned kernel configuration. e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp, the cache directory will be /home/tmp/spatialkernels/ on Linux. - Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable auto-tuning. - OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling for OpenCL command queue. This fix basic kernel get wrong running time, i.e. 0ms. - If creating cache directory failed, disable auto-tuning. * Detect and create cache dir on windows Signed-off-by: Li Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: Li Peng <peng.li@intel.com> * Fix redundant swizzleWeights calling when use cached kernel config. * Fix "out of resource" bug when auto-tuning too many kernels. * replace cl_mem with UMat in ocl4dnnConvSpatial class * OCL4DNN: reduce the tuning kernel candidate. This patch could reduce 75% of the tuning candidates with less than 2% performance impact for the final result. Signed-off-by: Zhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: Li Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: Li Peng <peng.li@intel.com> * Various fixes for ocl4dnn 1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()) 2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp 3. Code comments cleanup 4. ignore check on OCL cpu device Signed-off-by: Li Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: Li Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: Li Peng <peng.li@intel.com> * check kernel build options Signed-off-by: Li Peng <peng.li@intel.com> * Handle program compilation fail properly. * Use std::numeric_limits<float>::infinity() for large float number * check ocl4dnn kernel compilation result Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: Li Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: Li Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: Li Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: Li Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: Li Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: Li Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: Li Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: Li Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: Li Peng <peng.li@intel.com> * dnn(build): don't use C++11, OPENCL_LIBRARIES fix * dnn(ocl4dnn): remove unused OpenCL kernels * dnn(ocl4dnn): extract OpenCL code into .cl files * dnn(ocl4dnn): refine auto-tuning Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING environment variable to enable it. Use a set of pre-tuned configs as default config if auto-tuning is disabled. These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet, AlexNet, ResNet-50 If default config is not suitable, use the first available kernel config from the candidates. Candidate priority from high to low is gemm like kernel, IDLF kernel, basick kernel. * dnn(ocl4dnn): pooling doesn't use OpenCL subgroups * dnn(ocl4dnn): fix perf test OpenCV has default 3sec time limit for each performance test. Warmup OpenCL backend outside of perf measurement loop. * use ocl::KernelArg as much as possible Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): Refine signature of kernel config - Use more readable string as signture of kernel config - Don't count device name and vendor in signature string - Default kernel configurations are tuned for Intel GPU with 24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model. * dnn(ocl4dnn): swap width/height in configuration * dnn(ocl4dnn): enable configs for Intel OpenCL runtime only * core: make configuration helper functions accessible from non-core modules * dnn(ocl4dnn): update kernel auto-tuning behavior Avoid unwanted creation of directories * dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash * dnn(ocl4dnn): remove redundant code * dnn(ocl4dnn): Add more clear message for simd size dismatch. * dnn(ocl4dnn): add const to const argument Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel * dnn(ocl4dnn): drop unused tuneLocalSize() * dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method * dnn(ocl4dnn): sanitize file names used for cache * dnn(perf): enable Network tests with OpenCL * dnn(ocl4dnn/conv): drop computeGlobalSize() * dnn(ocl4dnn/conv): drop unused fields * dnn(ocl4dnn/conv): simplify ctor * dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL * dnn(ocl4dnn/conv): drop unsupported double / untested half types * dnn(ocl4dnn/conv): drop unused variable * dnn(ocl4dnn/conv): alignSize/divUp * dnn(ocl4dnn/conv): use enum values * dnn(ocl4dnn): drop unused innerproduct variable Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): add an generic function to check cl option support * dnn(ocl4dnn): run softmax subgroup version kernel first Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
{
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_relu");
}
2018-06-27 21:34:36 +08:00
TEST_P(Test_Caffe_layers, Dropout)
{
Merge pull request #9114 from pengli:dnn_rebase add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: Li Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: Li Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: Li Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: Li Peng <peng.li@intel.com> * set OPENCL target when it is available and disable fuseLayer for OCL target for the time being Signed-off-by: Li Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * remove trailing space Signed-off-by: Li Peng <peng.li@intel.com> * Fixed tensorflow demo bug. Root cause is that tensorflow has different algorithm with libdnn to calculate convolution output dimension. libdnn don't calculate output dimension anymore and just use one passed in by config. * split gemm ocl file split it into gemm_buffer.cl and gemm_image.cl Signed-off-by: Li Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: Li Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: Li Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: Li Peng <peng.li@intel.com> * update softmax layer on some platform subgroup extension may not work well, fallback to non subgroup ocl acceleration. Signed-off-by: Li Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: Li Peng <peng.li@intel.com> * update output message Signed-off-by: Li Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: Li Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: Li Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * DNN: Don't link OpenCL library explicitly * DNN: Make default preferableTarget to DNN_TARGET_CPU User should set it to DNN_TARGET_OPENCL explicitly if want to use OpenCL acceleration. Also don't fusion when using DNN_TARGET_OPENCL * DNN: refine coding style * Add getOpenCLErrorString * DNN: Use int32_t/uint32_t instread of alias * Use namespace ocl4dnn to include libdnn things * remove extra copyTo in softmax ocl path Signed-off-by: Li Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: Li Peng <peng.li@intel.com> * Add prefer target property for layer class It is used to indicate the target for layer forwarding, either the default CPU target or OCL target. Signed-off-by: Li Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: Li Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: wzw <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: Li Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: Li Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: Li Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: Li Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: Li Peng <peng.li@intel.com> * Refine auto-tuning mechanism. - Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory for fine-tuned kernel configuration. e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp, the cache directory will be /home/tmp/spatialkernels/ on Linux. - Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable auto-tuning. - OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling for OpenCL command queue. This fix basic kernel get wrong running time, i.e. 0ms. - If creating cache directory failed, disable auto-tuning. * Detect and create cache dir on windows Signed-off-by: Li Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: Li Peng <peng.li@intel.com> * Fix redundant swizzleWeights calling when use cached kernel config. * Fix "out of resource" bug when auto-tuning too many kernels. * replace cl_mem with UMat in ocl4dnnConvSpatial class * OCL4DNN: reduce the tuning kernel candidate. This patch could reduce 75% of the tuning candidates with less than 2% performance impact for the final result. Signed-off-by: Zhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: Li Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: Li Peng <peng.li@intel.com> * Various fixes for ocl4dnn 1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()) 2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp 3. Code comments cleanup 4. ignore check on OCL cpu device Signed-off-by: Li Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: Li Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: Li Peng <peng.li@intel.com> * check kernel build options Signed-off-by: Li Peng <peng.li@intel.com> * Handle program compilation fail properly. * Use std::numeric_limits<float>::infinity() for large float number * check ocl4dnn kernel compilation result Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: Li Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: Li Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: Li Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: Li Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: Li Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: Li Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: Li Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: Li Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: Li Peng <peng.li@intel.com> * dnn(build): don't use C++11, OPENCL_LIBRARIES fix * dnn(ocl4dnn): remove unused OpenCL kernels * dnn(ocl4dnn): extract OpenCL code into .cl files * dnn(ocl4dnn): refine auto-tuning Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING environment variable to enable it. Use a set of pre-tuned configs as default config if auto-tuning is disabled. These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet, AlexNet, ResNet-50 If default config is not suitable, use the first available kernel config from the candidates. Candidate priority from high to low is gemm like kernel, IDLF kernel, basick kernel. * dnn(ocl4dnn): pooling doesn't use OpenCL subgroups * dnn(ocl4dnn): fix perf test OpenCV has default 3sec time limit for each performance test. Warmup OpenCL backend outside of perf measurement loop. * use ocl::KernelArg as much as possible Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): Refine signature of kernel config - Use more readable string as signture of kernel config - Don't count device name and vendor in signature string - Default kernel configurations are tuned for Intel GPU with 24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model. * dnn(ocl4dnn): swap width/height in configuration * dnn(ocl4dnn): enable configs for Intel OpenCL runtime only * core: make configuration helper functions accessible from non-core modules * dnn(ocl4dnn): update kernel auto-tuning behavior Avoid unwanted creation of directories * dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash * dnn(ocl4dnn): remove redundant code * dnn(ocl4dnn): Add more clear message for simd size dismatch. * dnn(ocl4dnn): add const to const argument Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel * dnn(ocl4dnn): drop unused tuneLocalSize() * dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method * dnn(ocl4dnn): sanitize file names used for cache * dnn(perf): enable Network tests with OpenCL * dnn(ocl4dnn/conv): drop computeGlobalSize() * dnn(ocl4dnn/conv): drop unused fields * dnn(ocl4dnn/conv): simplify ctor * dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL * dnn(ocl4dnn/conv): drop unsupported double / untested half types * dnn(ocl4dnn/conv): drop unused variable * dnn(ocl4dnn/conv): alignSize/divUp * dnn(ocl4dnn/conv): use enum values * dnn(ocl4dnn): drop unused innerproduct variable Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): add an generic function to check cl option support * dnn(ocl4dnn): run softmax subgroup version kernel first Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
testLayerUsingCaffeModels("layer_dropout");
}
TEST_P(Test_Caffe_layers, Concat)
Merge pull request #9114 from pengli:dnn_rebase add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: Li Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: Li Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: Li Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: Li Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: Li Peng <peng.li@intel.com> * set OPENCL target when it is available and disable fuseLayer for OCL target for the time being Signed-off-by: Li Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * remove trailing space Signed-off-by: Li Peng <peng.li@intel.com> * Fixed tensorflow demo bug. Root cause is that tensorflow has different algorithm with libdnn to calculate convolution output dimension. libdnn don't calculate output dimension anymore and just use one passed in by config. * split gemm ocl file split it into gemm_buffer.cl and gemm_image.cl Signed-off-by: Li Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: Li Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: Li Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: Li Peng <peng.li@intel.com> * update softmax layer on some platform subgroup extension may not work well, fallback to non subgroup ocl acceleration. Signed-off-by: Li Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: Li Peng <peng.li@intel.com> * update output message Signed-off-by: Li Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: Li Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: Li Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * DNN: Don't link OpenCL library explicitly * DNN: Make default preferableTarget to DNN_TARGET_CPU User should set it to DNN_TARGET_OPENCL explicitly if want to use OpenCL acceleration. Also don't fusion when using DNN_TARGET_OPENCL * DNN: refine coding style * Add getOpenCLErrorString * DNN: Use int32_t/uint32_t instread of alias * Use namespace ocl4dnn to include libdnn things * remove extra copyTo in softmax ocl path Signed-off-by: Li Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: Li Peng <peng.li@intel.com> * Add prefer target property for layer class It is used to indicate the target for layer forwarding, either the default CPU target or OCL target. Signed-off-by: Li Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: Li Peng <peng.li@intel.com> Signed-off-by: wzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: Li Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: wzw <zhiwen.wu@intel.com> Signed-off-by: Li Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: Li Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: Li Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: Li Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: Li Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: Li Peng <peng.li@intel.com> * Refine auto-tuning mechanism. - Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory for fine-tuned kernel configuration. e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp, the cache directory will be /home/tmp/spatialkernels/ on Linux. - Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable auto-tuning. - OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling for OpenCL command queue. This fix basic kernel get wrong running time, i.e. 0ms. - If creating cache directory failed, disable auto-tuning. * Detect and create cache dir on windows Signed-off-by: Li Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: Li Peng <peng.li@intel.com> * Fix redundant swizzleWeights calling when use cached kernel config. * Fix "out of resource" bug when auto-tuning too many kernels. * replace cl_mem with UMat in ocl4dnnConvSpatial class * OCL4DNN: reduce the tuning kernel candidate. This patch could reduce 75% of the tuning candidates with less than 2% performance impact for the final result. Signed-off-by: Zhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: Li Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: Li Peng <peng.li@intel.com> * Various fixes for ocl4dnn 1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()) 2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp 3. Code comments cleanup 4. ignore check on OCL cpu device Signed-off-by: Li Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: Li Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: Li Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: Li Peng <peng.li@intel.com> * check kernel build options Signed-off-by: Li Peng <peng.li@intel.com> * Handle program compilation fail properly. * Use std::numeric_limits<float>::infinity() for large float number * check ocl4dnn kernel compilation result Signed-off-by: Li Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: Li Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: Li Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: Li Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: Li Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: Li Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: Li Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: Li Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: Li Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: Li Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: Li Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: Li Peng <peng.li@intel.com> * dnn(build): don't use C++11, OPENCL_LIBRARIES fix * dnn(ocl4dnn): remove unused OpenCL kernels * dnn(ocl4dnn): extract OpenCL code into .cl files * dnn(ocl4dnn): refine auto-tuning Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING environment variable to enable it. Use a set of pre-tuned configs as default config if auto-tuning is disabled. These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet, AlexNet, ResNet-50 If default config is not suitable, use the first available kernel config from the candidates. Candidate priority from high to low is gemm like kernel, IDLF kernel, basick kernel. * dnn(ocl4dnn): pooling doesn't use OpenCL subgroups * dnn(ocl4dnn): fix perf test OpenCV has default 3sec time limit for each performance test. Warmup OpenCL backend outside of perf measurement loop. * use ocl::KernelArg as much as possible Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): Refine signature of kernel config - Use more readable string as signture of kernel config - Don't count device name and vendor in signature string - Default kernel configurations are tuned for Intel GPU with 24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model. * dnn(ocl4dnn): swap width/height in configuration * dnn(ocl4dnn): enable configs for Intel OpenCL runtime only * core: make configuration helper functions accessible from non-core modules * dnn(ocl4dnn): update kernel auto-tuning behavior Avoid unwanted creation of directories * dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash * dnn(ocl4dnn): remove redundant code * dnn(ocl4dnn): Add more clear message for simd size dismatch. * dnn(ocl4dnn): add const to const argument Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel * dnn(ocl4dnn): drop unused tuneLocalSize() * dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method * dnn(ocl4dnn): sanitize file names used for cache * dnn(perf): enable Network tests with OpenCL * dnn(ocl4dnn/conv): drop computeGlobalSize() * dnn(ocl4dnn/conv): drop unused fields * dnn(ocl4dnn/conv): simplify ctor * dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL * dnn(ocl4dnn/conv): drop unsupported double / untested half types * dnn(ocl4dnn/conv): drop unused variable * dnn(ocl4dnn/conv): alignSize/divUp * dnn(ocl4dnn/conv): use enum values * dnn(ocl4dnn): drop unused innerproduct variable Signed-off-by: Li Peng <peng.li@intel.com> * dnn(ocl4dnn): add an generic function to check cl option support * dnn(ocl4dnn): run softmax subgroup version kernel first Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 20:38:00 +08:00
{
2019-07-27 18:30:15 +08:00
#if defined(INF_ENGINE_RELEASE)
#if INF_ENGINE_VER_MAJOR_GE(2019010000) && INF_ENGINE_VER_MAJOR_LT(2019020000)
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && target == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER, CV_TEST_TAG_DNN_SKIP_IE_VERSION);
2019-07-27 18:30:15 +08:00
#elif INF_ENGINE_VER_MAJOR_EQ(2019020000)
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 &&
(target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16))
applyTestTag(target == DNN_TARGET_OPENCL ? CV_TEST_TAG_DNN_SKIP_IE_OPENCL : CV_TEST_TAG_DNN_SKIP_IE_OPENCL_FP16,
CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER, CV_TEST_TAG_DNN_SKIP_IE_VERSION);
2019-07-27 18:30:15 +08:00
#endif
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH &&
(target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16))
applyTestTag(target == DNN_TARGET_OPENCL ? CV_TEST_TAG_DNN_SKIP_IE_OPENCL : CV_TEST_TAG_DNN_SKIP_IE_OPENCL_FP16,
CV_TEST_TAG_DNN_SKIP_IE_NGRAPH, CV_TEST_TAG_DNN_SKIP_IE_VERSION);
2019-01-31 18:38:24 +08:00
#endif
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_concat");
testLayerUsingCaffeModels("layer_concat_optim", true, false);
testLayerUsingCaffeModels("layer_concat_shared_input", true, false);
}
2018-06-27 21:34:36 +08:00
TEST_P(Test_Caffe_layers, Fused_Concat)
{
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && (target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16))
applyTestTag(target == DNN_TARGET_OPENCL ? CV_TEST_TAG_DNN_SKIP_IE_OPENCL : CV_TEST_TAG_DNN_SKIP_IE_OPENCL_FP16,
CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER, CV_TEST_TAG_DNN_SKIP_IE_VERSION);
2018-06-27 21:34:36 +08:00
checkBackend();
// Test case
// input
// |
// v
// some_layer
// | |
// v v
// concat
Net net;
int interLayer;
{
LayerParams lp;
lp.type = "AbsVal";
lp.name = "someLayer";
interLayer = net.addLayerToPrev(lp.name, lp.type, lp);
}
{
LayerParams lp;
lp.set("axis", 1);
lp.type = "Concat";
lp.name = "testConcat";
int id = net.addLayer(lp.name, lp.type, lp);
net.connect(interLayer, 0, id, 0);
net.connect(interLayer, 0, id, 1);
}
int shape[] = {1, 2, 3, 4};
Mat input(4, shape, CV_32F);
randu(input, 0.0f, 1.0f); // [0, 1] to make AbsVal an identity transformation.
net.setInput(input);
2018-06-27 21:34:36 +08:00
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
Mat out = net.forward();
2018-06-27 21:34:36 +08:00
normAssert(slice(out, Range::all(), Range(0, 2), Range::all(), Range::all()), input, "", default_l1, default_lInf);
normAssert(slice(out, Range::all(), Range(2, 4), Range::all(), Range::all()), input, "", default_l1, default_lInf);
}
TEST_P(Test_Caffe_layers, Eltwise)
{
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && target == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD);
2018-06-27 21:34:36 +08:00
testLayerUsingCaffeModels("layer_eltwise");
}
TEST_P(Test_Caffe_layers, PReLU)
{
double lInf = (target == DNN_TARGET_MYRIAD || target == DNN_TARGET_OPENCL_FP16) ? 0.021 : 0.0;
testLayerUsingCaffeModels("layer_prelu", true, true, 0.0, lInf);
2018-06-27 21:34:36 +08:00
}
// TODO: fix an unstable test case
TEST_P(Test_Caffe_layers, layer_prelu_fc)
{
if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16)
applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL_FP16);
// Reference output values are in range [-0.0001, 10.3906]
double l1 = (target == DNN_TARGET_MYRIAD) ? 0.005 : 0.0;
double lInf = (target == DNN_TARGET_MYRIAD) ? 0.021 : 0.0;
#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_VER_MAJOR_EQ(2020040000)
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH && target == DNN_TARGET_OPENCL)
{
l1 = 0.006f; lInf = 0.05f;
}
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH && target == DNN_TARGET_OPENCL_FP16)
{
l1 = 0.01f; lInf = 0.05f;
}
#endif
testLayerUsingCaffeModels("layer_prelu_fc", true, false, l1, lInf);
}
2018-06-27 21:34:36 +08:00
TEST_P(Test_Caffe_layers, Reshape_Split_Slice)
{
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
2018-06-27 21:34:36 +08:00
Net net = readNetFromCaffe(_tf("reshape_and_slice_routines.prototxt"));
ASSERT_FALSE(net.empty());
2018-06-27 21:34:36 +08:00
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
Mat input(6, 12, CV_32F);
RNG rng(0);
rng.fill(input, RNG::UNIFORM, -1, 1);
net.setInput(input, "input");
Mat output = net.forward("output");
2018-06-27 21:34:36 +08:00
normAssert(input, output, "", default_l1, default_lInf);
}
2017-08-01 21:58:34 +08:00
2018-06-27 21:34:36 +08:00
TEST_P(Test_Caffe_layers, Conv_Elu)
2017-08-01 21:58:34 +08:00
{
#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE <= 2018050000
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && target == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_VERSION);
#endif
Net net = readNetFromTensorflow(_tf("layer_elu_model.pb"));
ASSERT_FALSE(net.empty());
2017-08-01 21:58:34 +08:00
Mat inp = blobFromNPY(_tf("layer_elu_in.npy"));
Mat ref = blobFromNPY(_tf("layer_elu_out.npy"));
net.setInput(inp, "input");
2018-06-27 21:34:36 +08:00
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
2017-08-01 21:58:34 +08:00
Mat out = net.forward();
double l1 = default_l1, lInf = default_lInf;
if (target == DNN_TARGET_CUDA_FP16)
{
l1 = 0.0002;
lInf = 0.0005;
}
normAssert(ref, out, "", l1, lInf);
2017-08-01 21:58:34 +08:00
}
class Layer_LSTM_Test : public ::testing::Test
{
public:
int numInp, numOut;
Mat Wh, Wx, b;
Ptr<LSTMLayer> layer;
std::vector<Mat> inputs, outputs;
Layer_LSTM_Test() {}
2017-08-25 19:45:03 +08:00
void init(const MatShape &inpShape_, const MatShape &outShape_,
bool produceCellOutput, bool useTimestampDim)
{
numInp = total(inpShape_);
numOut = total(outShape_);
Wh = Mat::ones(4 * numOut, numOut, CV_32F);
Wx = Mat::ones(4 * numOut, numInp, CV_32F);
b = Mat::ones(4 * numOut, 1, CV_32F);
2017-08-25 19:45:03 +08:00
LayerParams lp;
lp.blobs.resize(3);
lp.blobs[0] = Wh;
lp.blobs[1] = Wx;
lp.blobs[2] = b;
lp.set<bool>("produce_cell_output", produceCellOutput);
lp.set<bool>("use_timestamp_dim", useTimestampDim);
layer = LSTMLayer::create(lp);
layer->setOutShape(outShape_);
}
};
TEST_F(Layer_LSTM_Test, get_set_test)
{
const int TN = 4;
MatShape inpShape = shape(5, 3, 2);
MatShape outShape = shape(3, 1, 2);
MatShape inpResShape = concat(shape(TN), inpShape);
MatShape outResShape = concat(shape(TN), outShape);
2017-08-25 19:45:03 +08:00
init(inpShape, outShape, true, false);
layer->setOutShape(outShape);
Mat C((int)outResShape.size(), &outResShape[0], CV_32F);
randu(C, -1., 1.);
Mat H = C.clone();
randu(H, -1., 1.);
Mat inp((int)inpResShape.size(), &inpResShape[0], CV_32F);
randu(inp, -1., 1.);
inputs.push_back(inp);
runLayer(layer, inputs, outputs);
EXPECT_EQ(2u, outputs.size());
print(outResShape, "outResShape");
print(shape(outputs[0]), "out0");
print(shape(outputs[0]), "out1");
EXPECT_EQ(outResShape, shape(outputs[0]));
EXPECT_EQ(outResShape, shape(outputs[1]));
EXPECT_EQ(0, layer->inputNameToIndex("x"));
EXPECT_EQ(0, layer->outputNameToIndex("h"));
EXPECT_EQ(1, layer->outputNameToIndex("c"));
}
TEST(Layer_LSTM_Test_Accuracy_with_, CaffeRecurrent)
{
2017-08-25 19:45:03 +08:00
LayerParams lp;
lp.blobs.resize(3);
lp.blobs[0] = blobFromNPY(_tf("lstm.prototxt.w_2.npy")); // Wh
lp.blobs[1] = blobFromNPY(_tf("lstm.prototxt.w_0.npy")); // Wx
lp.blobs[2] = blobFromNPY(_tf("lstm.prototxt.w_1.npy")); // bias
Ptr<LSTMLayer> layer = LSTMLayer::create(lp);
Mat inp = blobFromNPY(_tf("recurrent.input.npy"));
std::vector<Mat> inputs(1, inp), outputs;
runLayer(layer, inputs, outputs);
Mat h_t_reference = blobFromNPY(_tf("lstm.prototxt.h_1.npy"));
normAssert(h_t_reference, outputs[0]);
}
TEST(Layer_RNN_Test_Accuracy_with_, CaffeRecurrent)
{
Ptr<RNNLayer> layer = RNNLayer::create(LayerParams());
layer->setWeights(
blobFromNPY(_tf("rnn.prototxt.w_0.npy")),
blobFromNPY(_tf("rnn.prototxt.w_1.npy")),
blobFromNPY(_tf("rnn.prototxt.w_2.npy")),
blobFromNPY(_tf("rnn.prototxt.w_3.npy")),
blobFromNPY(_tf("rnn.prototxt.w_4.npy")) );
std::vector<Mat> output, input(1, blobFromNPY(_tf("recurrent.input.npy")));
runLayer(layer, input, output);
Mat h_ref = blobFromNPY(_tf("rnn.prototxt.h_1.npy"));
normAssert(h_ref, output[0]);
}
TEST(Layer_LSTM_Test_Accuracy_, Reverse)
{
// This handcrafted setup calculates (approximately) the prefix sum of the
// input, assuming the inputs are suitably small.
cv::Mat input(2, 1, CV_32FC1);
input.at<float>(0, 0) = 1e-5f;
input.at<float>(1, 0) = 2e-5f;
cv::Mat Wx(4, 1, CV_32FC1);
Wx.at<float>(0, 0) = 0.f; // Input gate
Wx.at<float>(1, 0) = 0.f; // Forget gate
Wx.at<float>(2, 0) = 0.f; // Output gate
Wx.at<float>(3, 0) = 1.f; // Update signal
cv::Mat Wh(4, 1, CV_32FC1);
Wh.at<float>(0, 0) = 0.f; // Input gate
Wh.at<float>(1, 0) = 0.f; // Forget gate
Wh.at<float>(2, 0) = 0.f; // Output gate
Wh.at<float>(3, 0) = 0.f; // Update signal
cv::Mat bias(4, 1, CV_32FC1);
bias.at<float>(0, 0) = 1e10f; // Input gate - always allows input to c
bias.at<float>(1, 0) = 1e10f; // Forget gate - never forget anything on c
bias.at<float>(2, 0) = 1e10f; // Output gate - always output everything
bias.at<float>(3, 0) = 0.f; // Update signal
LayerParams lp;
lp.set("reverse", true);
lp.set("use_timestamp_dim", true);
lp.blobs.clear();
lp.blobs.push_back(Wh);
lp.blobs.push_back(Wx);
lp.blobs.push_back(bias);
cv::Ptr<cv::dnn::LSTMLayer> layer = LSTMLayer::create(lp);
std::vector<cv::Mat> outputs;
std::vector<cv::Mat> inputs;
inputs.push_back(input);
runLayer(layer, inputs, outputs);
ASSERT_EQ(1, outputs.size());
cv::Mat out = outputs[0];
ASSERT_EQ(3, out.dims);
ASSERT_EQ(shape(2, 1, 1), shape(out));
float* data = reinterpret_cast<float*>(out.data);
EXPECT_NEAR(std::tanh(1e-5f) + std::tanh(2e-5f), data[0], 1e-10);
EXPECT_NEAR(std::tanh(2e-5f), data[1], 1e-10);
}
class Layer_RNN_Test : public ::testing::Test
{
public:
int nX, nH, nO, nT, nS;
Mat Whh, Wxh, bh, Who, bo;
Ptr<RNNLayer> layer;
std::vector<Mat> inputs, outputs;
Layer_RNN_Test()
{
nT = 3;
nS = 5;
nX = 31;
nH = 64;
nO = 100;
Whh = Mat::ones(nH, nH, CV_32F);
Wxh = Mat::ones(nH, nX, CV_32F);
bh = Mat::ones(nH, 1, CV_32F);
Who = Mat::ones(nO, nH, CV_32F);
bo = Mat::ones(nO, 1, CV_32F);
layer = RNNLayer::create(LayerParams());
layer->setProduceHiddenOutput(true);
layer->setWeights(Wxh, bh, Whh, Who, bo);
}
};
TEST_F(Layer_RNN_Test, get_set_test)
{
int sz[] = { nT, nS, 1, nX };
Mat inp(4, sz, CV_32F);
randu(inp, -1., 1.);
inputs.push_back(inp);
runLayer(layer, inputs, outputs);
EXPECT_EQ(outputs.size(), 2u);
EXPECT_EQ(shape(outputs[0]), shape(nT, nS, nO));
EXPECT_EQ(shape(outputs[1]), shape(nT, nS, nH));
}
TEST_P(Test_Caffe_layers, Accum)
{
if (backend == DNN_BACKEND_OPENCV && target != DNN_TARGET_CPU)
applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL, CV_TEST_TAG_DNN_SKIP_OPENCL_FP16);
testLayerUsingCaffeModels("accum", false, false, 0.0, 0.0, 2);
testLayerUsingCaffeModels("accum_ref", false, false, 0.0, 0.0, 2);
}
TEST_P(Test_Caffe_layers, FlowWarp)
{
if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16)
applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL_FP16);
testLayerUsingCaffeModels("flow_warp", false, false, 0.0, 0.0, 2);
}
TEST_P(Test_Caffe_layers, ChannelNorm)
{
if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16)
applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL_FP16);
testLayerUsingCaffeModels("channel_norm", false, false);
}
TEST_P(Test_Caffe_layers, DataAugmentation)
{
if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16)
applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL_FP16);
testLayerUsingCaffeModels("data_augmentation", true, false);
testLayerUsingCaffeModels("data_augmentation_2x1", true, false);
testLayerUsingCaffeModels("data_augmentation_8x6", true, false);
}
TEST_P(Test_Caffe_layers, Resample)
{
if (backend != DNN_BACKEND_OPENCV)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER, CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
testLayerUsingCaffeModels("nearest_2inps", false, false, 0.0, 0.0, 2);
testLayerUsingCaffeModels("nearest", false, false);
}
TEST_P(Test_Caffe_layers, Correlation)
{
if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NGRAPH, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER,
CV_TEST_TAG_DNN_SKIP_OPENCL, CV_TEST_TAG_DNN_SKIP_OPENCL_FP16);
testLayerUsingCaffeModels("correlation", false, false, 0.0, 0.0, 2);
}
TEST_P(Test_Caffe_layers, Convolution2Inputs)
{
testLayerUsingCaffeModels("conv_2_inps", true, false, 0.0, 0.0, 2);
}
TEST_P(Test_Caffe_layers, ROIPooling_Accuracy)
2017-12-04 23:45:30 +08:00
{
Net net = readNetFromCaffe(_tf("net_roi_pooling.prototxt"));
ASSERT_FALSE(net.empty());
2017-12-04 23:45:30 +08:00
Mat inp = blobFromNPY(_tf("net_roi_pooling.input.npy"));
Mat rois = blobFromNPY(_tf("net_roi_pooling.rois.npy"));
Mat ref = blobFromNPY(_tf("net_roi_pooling.npy"));
checkBackend(&inp, &ref);
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
2017-12-04 23:45:30 +08:00
net.setInput(inp, "input");
net.setInput(rois, "rois");
Mat out = net.forward();
double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 1e-3 : 1e-5;
double lInf = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 1e-3 : 1e-4;
2019-12-22 22:44:38 +08:00
if (target == DNN_TARGET_CUDA_FP16)
{
l1 = 2e-4;
lInf = 9e-4;
}
normAssert(out, ref, "", l1, lInf);
2017-12-04 23:45:30 +08:00
}
2018-04-04 19:48:29 +08:00
TEST_P(Test_Caffe_layers, FasterRCNN_Proposal)
2017-12-14 00:06:30 +08:00
{
if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16)
applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL_FP16);
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
Merge pull request #14827 from YashasSamaga:cuda4dnn-csl-low CUDA backend for the DNN module * stub cuda4dnn design * minor fixes for tests and doxygen * add csl public api directory to module headers * add low-level CSL components * add high-level CSL components * integrate csl::Tensor into backbone code * switch to CPU iff unsupported; otherwise, fail on error * add fully connected layer * add softmax layer * add activation layers * support arbitary rank TensorDescriptor * pass input wrappers to `initCUDA()` * add 1d/2d/3d-convolution * add pooling layer * reorganize and refactor code * fixes for gcc, clang and doxygen; remove cxx14/17 code * add blank_layer * add LRN layer * add rounding modes for pooling layer * split tensor.hpp into tensor.hpp and tensor_ops.hpp * add concat layer * add scale layer * add batch normalization layer * split math.cu into activations.cu and math.hpp * add eltwise layer * add flatten layer * add tensor transform api * add asymmetric padding support for convolution layer * add reshape layer * fix rebase issues * add permute layer * add padding support for concat layer * refactor and reorganize code * add normalize layer * optimize bias addition in scale layer * add prior box layer * fix and optimize normalize layer * add asymmetric padding support for pooling layer * add event API * improve pooling performance for some padding scenarios * avoid over-allocation of compute resources to kernels * improve prior box performance * enable layer fusion * add const layer * add resize layer * add slice layer * add padding layer * add deconvolution layer * fix channelwise ReLU initialization * add vector traits * add vectorized versions of relu, clipped_relu, power * add vectorized concat kernels * improve concat_with_offsets performance * vectorize scale and bias kernels * add support for multi-billion element tensors * vectorize prior box kernels * fix address alignment check * improve bias addition performance of conv/deconv/fc layers * restructure code for supporting multiple targets * add DNN_TARGET_CUDA_FP64 * add DNN_TARGET_FP16 * improve vectorization * add region layer * improve tensor API, add dynamic ranks 1. use ManagedPtr instead of a Tensor in backend wrapper 2. add new methods to tensor classes - size_range: computes the combined size of for a given axis range - tensor span/view can be constructed from a raw pointer and shape 3. the tensor classes can change their rank at runtime (previously rank was fixed at compile-time) 4. remove device code from tensor classes (as they are unused) 5. enforce strict conditions on tensor class APIs to improve debugging ability * fix parametric relu activation * add squeeze/unsqueeze tensor API * add reorg layer * optimize permute and enable 2d permute * enable 1d and 2d slice * add split layer * add shuffle channel layer * allow tensors of different ranks in reshape primitive * patch SliceOp to allow Crop Layer * allow extra shape inputs in reshape layer * use `std::move_backward` instead of `std::move` for insert in resizable_static_array * improve workspace management * add spatial LRN * add nms (cpu) to region layer * add max pooling with argmax ( and a fix to limits.hpp) * add max unpooling layer * rename DNN_TARGET_CUDA_FP32 to DNN_TARGET_CUDA * update supportBackend to be more rigorous * remove stray include from preventing non-cuda build * include op_cuda.hpp outside condition #if * refactoring, fixes and many optimizations * drop DNN_TARGET_CUDA_FP64 * fix gcc errors * increase max. tensor rank limit to six * add Interp layer * drop custom layers; use BackendNode * vectorize activation kernels * fixes for gcc * remove wrong assertion * fix broken assertion in unpooling primitive * fix build errors in non-CUDA build * completely remove workspace from public API * fix permute layer * enable accuracy and perf. tests for DNN_TARGET_CUDA * add asynchronous forward * vectorize eltwise ops * vectorize fill kernel * fixes for gcc * remove CSL headers from public API * remove csl header source group from cmake * update min. cudnn version in cmake * add numerically stable FP32 log1pexp * refactor code * add FP16 specialization to cudnn based tensor addition * vectorize scale1 and bias1 + minor refactoring * fix doxygen build * fix invalid alignment assertion * clear backend wrappers before allocateLayers * ignore memory lock failures * do not allocate internal blobs * integrate NVTX * add numerically stable half precision log1pexp * fix indentation, following coding style, improve docs * remove accidental modification of IE code * Revert "add asynchronous forward" This reverts commit 1154b9da9da07e9b52f8a81bdcea48cf31c56f70. * [cmake] throw error for unsupported CC versions * fix rebase issues * add more docs, refactor code, fix bugs * minor refactoring and fixes * resolve warnings/errors from clang * remove haveCUDA() checks from supportBackend() * remove NVTX integration * changes based on review comments * avoid exception when no CUDA device is present * add color code for CUDA in Net::dump
2019-10-21 19:28:00 +08:00
if(backend == DNN_BACKEND_CUDA)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA); /* Proposal layer is unsupported */
2017-12-14 00:06:30 +08:00
Net net = readNetFromCaffe(_tf("net_faster_rcnn_proposal.prototxt"));
Mat scores = blobFromNPY(_tf("net_faster_rcnn_proposal.scores.npy"));
Mat deltas = blobFromNPY(_tf("net_faster_rcnn_proposal.deltas.npy"));
Mat imInfo = (Mat_<float>(1, 3) << 600, 800, 1.6f);
net.setInput(scores, "rpn_cls_prob_reshape");
net.setInput(deltas, "rpn_bbox_pred");
net.setInput(imInfo, "im_info");
2018-04-04 19:48:29 +08:00
std::vector<Mat> outs;
2018-06-27 21:34:36 +08:00
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
2018-04-04 19:48:29 +08:00
net.forward(outs, "output");
2018-04-04 19:48:29 +08:00
for (int i = 0; i < 2; ++i)
{
Mat ref = blobFromNPY(_tf(i == 0 ? "net_faster_rcnn_proposal.out_rois.npy" :
"net_faster_rcnn_proposal.out_scores.npy"));
const int numDets = ref.size[0];
EXPECT_LE(numDets, outs[i].size[0]);
normAssert(outs[i].rowRange(0, numDets), ref);
if (numDets < outs[i].size[0])
2018-11-15 04:25:23 +08:00
{
2018-04-04 19:48:29 +08:00
EXPECT_EQ(countNonZero(outs[i].rowRange(numDets, outs[i].size[0])), 0);
2018-11-15 04:25:23 +08:00
}
2018-04-04 19:48:29 +08:00
}
}
typedef testing::TestWithParam<tuple<Vec4i, Vec2i, bool> > Scale_untrainable;
TEST_P(Scale_untrainable, Accuracy)
{
Vec4i inpShapeVec = get<0>(GetParam());
int axis = get<1>(GetParam())[0];
int weightsDims = get<1>(GetParam())[1];
bool testFusion = get<2>(GetParam());
const int inpShape[] = {inpShapeVec[0], inpShapeVec[1], inpShapeVec[2], inpShapeVec[3]};
// Create a network with two inputs. Scale layer multiplies a first input to
// a second one. See http://caffe.berkeleyvision.org/tutorial/layers/scale.html
Net net;
// Check that this version of Scale layer won't be fused with Convolution layer.
if (testFusion)
{
LayerParams lp;
lp.set("kernel_size", 1);
lp.set("num_output", 3);
lp.set("group", 3);
lp.set("bias_term", false);
lp.type = "Convolution";
lp.name = "testConv";
std::vector<int> weightsShape(4);
weightsShape[0] = 3; // #outChannels
weightsShape[1] = 1; // #inpChannels / group
weightsShape[2] = 1; // height
weightsShape[3] = 1; // width
Mat weights(weightsShape, CV_32F);
weights.setTo(1);
lp.blobs.push_back(weights);
net.addLayerToPrev(lp.name, lp.type, lp);
}
LayerParams lp;
lp.type = "Scale";
lp.name = "testLayer";
lp.set("axis", axis);
int id = net.addLayerToPrev(lp.name, lp.type, lp);
net.connect(0, 1, id, 1);
Mat input(4, inpShape, CV_32F);
Mat weights(weightsDims, &inpShape[axis], CV_32F);
randu(input, -1, 1);
randu(weights, -1, 1);
std::vector<String> inpNames(2);
inpNames[0] = "scale_input";
inpNames[1] = "scale_weights";
net.setInputsNames(inpNames);
net.setInput(input, inpNames[0]);
net.setInput(weights, inpNames[1]);
net.setPreferableBackend(DNN_BACKEND_OPENCV);
Mat out = net.forward();
Mat ref(input.dims, input.size, CV_32F);
float* inpData = (float*)input.data;
float* refData = (float*)ref.data;
float* weightsData = (float*)weights.data;
int spatialSize = 1;
for (int i = axis + weightsDims; i < 4; ++i)
spatialSize *= inpShape[i];
for (int i = 0; i < ref.total(); ++i)
{
float w = weightsData[(i / spatialSize) % weights.total()];
refData[i] = inpData[i] * w;
}
normAssert(out, ref);
}
INSTANTIATE_TEST_CASE_P(Layer_Test, Scale_untrainable, Combine(
/*input size*/ Values(Vec4i(2, 3, 4, 5)),
/*axis, #dims*/ Values(Vec2i(0, 1), Vec2i(0, 2), Vec2i(0, 3), Vec2i(0, 4),
Vec2i(1, 1), Vec2i(1, 2), Vec2i(1, 3),
Vec2i(2, 1), Vec2i(2, 2),
Vec2i(3, 1)),
/*conv fusion*/ testing::Bool()
));
typedef testing::TestWithParam<tuple<Vec4i, Vec4i, int, int, int> > Crop;
TEST_P(Crop, Accuracy)
{
Vec4i inpShapeVec = get<0>(GetParam());
Vec4i sizShapeVec = get<1>(GetParam());
int axis = get<2>(GetParam());
int numOffsets = get<3>(GetParam());
int offsetVal = get<4>(GetParam());
const int inpShape[] = {inpShapeVec[0], inpShapeVec[1], inpShapeVec[2], inpShapeVec[3]};
const int sizShape[] = {sizShapeVec[0], sizShapeVec[1], sizShapeVec[2], sizShapeVec[3]};
// Create a network with two inputs. Crop layer crops a first input to
// the size of a second one.
// See http://caffe.berkeleyvision.org/tutorial/layers/crop.html
Net net;
LayerParams lp;
lp.name = "testCrop";
lp.type = "Crop";
lp.set("axis", axis);
if (numOffsets > 0)
{
std::vector<int> offsets(numOffsets, offsetVal);
lp.set("offset", DictValue::arrayInt<int*>(&offsets[0], offsets.size()));
}
else
offsetVal = 0;
int id = net.addLayerToPrev(lp.name, lp.type, lp);
net.connect(0, 1, id, 1);
Mat inpImage(4, inpShape, CV_32F);
Mat sizImage(4, sizShape, CV_32F);
randu(inpImage, -1, 1);
randu(sizImage, -1, 1);
std::vector<String> inpNames(2);
inpNames[0] = "cropImage";
inpNames[1] = "sizImage";
net.setInputsNames(inpNames);
net.setInput(inpImage, inpNames[0]);
net.setInput(sizImage, inpNames[1]);
net.setPreferableBackend(DNN_BACKEND_OPENCV);
// There are a few conditions that represent invalid input to the crop
// layer, so in those cases we want to verify an exception is thrown.
bool shouldThrowException = false;
if (numOffsets > 1 && numOffsets != 4 - axis)
shouldThrowException = true;
else
for (int i = axis; i < 4; i++)
if (sizShape[i] + offsetVal > inpShape[i])
shouldThrowException = true;
Mat out;
if (shouldThrowException)
{
ASSERT_ANY_THROW(out = net.forward());
return;
}
else
out = net.forward();
// Finally, compare the cropped output blob from the DNN layer (out)
// to a reference blob (ref) that we compute here.
std::vector<Range> crop_range;
crop_range.resize(4, Range::all());
for (int i = axis; i < 4; i++)
crop_range[i] = Range(offsetVal, sizShape[i] + offsetVal);
Mat ref(sizImage.dims, sizImage.size, CV_32F);
inpImage(&crop_range[0]).copyTo(ref);
normAssert(out, ref);
}
INSTANTIATE_TEST_CASE_P(Layer_Test, Crop, Combine(
/*input blob shape*/ Values(Vec4i(1, 3, 20, 30)),
/*cropsize blob shape*/ Values(Vec4i(1, 3, 10, 12)),
/*start axis*/ Values(0, 1, 2),
/*number of offsets*/ Values(0, 1, 2, 4),
/*offset value*/ Values(3, 4)
));
// Check that by default average pooling layer should not count zero padded values
// into the normalization area.
2018-06-27 21:34:36 +08:00
TEST_P(Test_Caffe_layers, Average_pooling_kernel_area)
{
LayerParams lp;
lp.name = "testAvePool";
lp.type = "Pooling";
lp.set("kernel_size", 2);
lp.set("stride", 2);
lp.set("pool", "AVE");
Net net;
net.addLayerToPrev(lp.name, lp.type, lp);
// 1 2 | 3
// 4 5 | 6
// ----+--
// 7 8 | 9
Mat inp = (Mat_<float>(3, 3) << 1, 2, 3, 4, 5, 6, 7, 8, 9);
2018-06-27 21:34:36 +08:00
Mat ref = (Mat_<float>(2, 2) << (1 + 2 + 4 + 5) / 4.f, (3 + 6) / 2.f, (7 + 8) / 2.f, 9);
Mat tmp = blobFromImage(inp);
net.setInput(blobFromImage(inp));
2018-06-27 21:34:36 +08:00
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
Mat out = net.forward();
2018-06-27 21:34:36 +08:00
normAssert(out, blobFromImage(ref));
}
TEST_P(Test_Caffe_layers, PriorBox_repeated)
{
Net net = readNet(_tf("prior_box.prototxt"));
int inp_size[] = {1, 3, 10, 10};
int shape_size[] = {1, 2, 3, 4};
Mat inp(4, inp_size, CV_32F);
randu(inp, -1.0f, 1.0f);
Mat shape(4, shape_size, CV_32F);
randu(shape, -1.0f, 1.0f);
net.setInput(inp, "data");
net.setInput(shape, "shape");
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
Mat out = net.forward();
Mat ref = blobFromNPY(_tf("priorbox_output.npy"));
double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 1e-3 : 1e-5;
double lInf = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 1e-3 : 1e-4;
if (target == DNN_TARGET_CUDA_FP16)
{
l1 = 7e-5;
lInf = 0.0005;
}
normAssert(out, ref, "", l1, lInf);
}
// Test PriorBoxLayer in case of no aspect ratios (just squared proposals).
2018-06-27 21:34:36 +08:00
TEST_P(Test_Caffe_layers, PriorBox_squares)
{
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && target == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH && target == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
LayerParams lp;
lp.name = "testPriorBox";
lp.type = "PriorBox";
lp.set("min_size", 2);
lp.set("flip", true);
lp.set("clip", true);
float variance[] = {0.1f, 0.1f, 0.2f, 0.2f};
float aspectRatios[] = {1.0f}; // That should be ignored.
lp.set("variance", DictValue::arrayReal<float*>(&variance[0], 4));
lp.set("aspect_ratio", DictValue::arrayReal<float*>(&aspectRatios[0], 1));
Net net;
int id = net.addLayerToPrev(lp.name, lp.type, lp);
net.connect(0, 0, id, 1); // The second input is an input image. Shapes are used for boxes normalization.
Mat inp(1, 2, CV_32F);
randu(inp, -1, 1);
net.setInput(blobFromImage(inp));
2018-06-27 21:34:36 +08:00
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
Mat out = net.forward();
2018-06-27 21:34:36 +08:00
Mat ref = (Mat_<float>(4, 4) << 0.0, 0.0, 0.75, 1.0,
0.25, 0.0, 1.0, 1.0,
0.1f, 0.1f, 0.2f, 0.2f,
0.1f, 0.1f, 0.2f, 0.2f);
double l1 = 1e-5;
if (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD || target == DNN_TARGET_CUDA_FP16)
l1 = 2e-5;
2018-07-13 15:48:31 +08:00
normAssert(out.reshape(1, 4), ref, "", l1);
}
typedef TestWithParam<tuple<int, int> > Layer_Test_DWconv_Prelu;
TEST_P(Layer_Test_DWconv_Prelu, Accuracy)
{
// Test case
// input img size 3x16x16 value all 1
// |
// v
// dw_conv weight[0]=-1 weight[1]=-2 weight[2]=-3 bias={1,2,3}
// |
// v
// prelu weight={1,2,3}
// |
// v
// output out size 3x14x14 if right: out[0]=-8 out[0]=-32 out[0]=-72
// but current opencv output: out[0]=-24 out[0]=-48 out[0]=-72
const int num_input = get<0>(GetParam()); //inpChannels
const int group = 3; //outChannels=group when group>1
const int num_output = get<1>(GetParam());
const int kernel_depth = num_input/group;
CV_Assert_N(num_output >= group, num_output % group == 0, num_input % group == 0);
Net net;
//layer 1: dwconv
LayerParams lp;
lp.name = "dwconv";
lp.type = "Convolution";
lp.set("kernel_size", 3);
lp.set("num_output", num_output);
lp.set("pad", 0);
lp.set("group", group);
lp.set("stride", 1);
lp.set("engine", "CAFFE");
lp.set("bias_term", "true");
std::vector<int> weightsShape(4);
weightsShape[0] = num_output; // #outChannels
weightsShape[1] = kernel_depth; // #inpChannels / group
weightsShape[2] = 3; // height
weightsShape[3] = 3; // width
Mat weights(weightsShape, CV_32F, Scalar(1));
//assign weights
for (int i = 0; i < weightsShape[0]; ++i)
{
for (int j = 0; j < weightsShape[1]; ++j)
{
for (int k = 0; k < weightsShape[2]; ++k)
{
for (int l = 0; l < weightsShape[3]; ++l)
{
weights.ptr<float>(i, j, k)[l]=-1*(i+1);
}
}
}
}
lp.blobs.push_back(weights);
//assign bias
Mat bias(1, num_output, CV_32F, Scalar(1));
for (int i = 0; i < 1; ++i)
{
for (int j = 0; j < num_output; ++j)
{
bias.ptr<float>(i)[j]=j+1;
}
}
lp.blobs.push_back(bias);
net.addLayerToPrev(lp.name, lp.type, lp);
//layer 2: prelu
LayerParams lpr;
lpr.name = "dw_relu";
lpr.type = "PReLU";
Mat weightsp(1, num_output, CV_32F, Scalar(1));
//assign weights
for (int i = 0; i < 1; ++i)
{
for (int j = 0; j < num_output; ++j)
{
weightsp.ptr<float>(i)[j]=j+1;
}
}
lpr.blobs.push_back(weightsp);
net.addLayerToPrev(lpr.name, lpr.type, lpr);
int shape[] = {1, num_input, 16, 16};
Mat in_blob(4, &shape[0], CV_32FC1, Scalar(1));
net.setPreferableBackend(DNN_BACKEND_OPENCV);
net.setInput(in_blob);
Mat out = net.forward();
//assign target
std::vector<int> outShape(4);
outShape[0] = 1;
outShape[1] = num_output; // outChannels
outShape[2] = 14; // height
outShape[3] = 14; // width
Mat target(outShape, CV_32F, Scalar(1));
for (int i = 0; i < outShape[0]; ++i)
{
for (int j = 0; j < outShape[1]; ++j)
{
for (int k = 0; k < outShape[2]; ++k)
{
for (int l = 0; l < outShape[3]; ++l)
{
target.ptr<float>(i, j, k)[l]=(-9*kernel_depth*(j+1)+j+1)*(j+1);
}
}
}
}
normAssert(out, target);
}
INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_DWconv_Prelu, Combine(Values(3, 6), Values(3, 6)));
#ifdef HAVE_INF_ENGINE
// Using Intel's Model Optimizer generate .xml and .bin files:
// ./ModelOptimizer -w /path/to/caffemodel -d /path/to/prototxt \
// -p FP32 -i -b ${batch_size} -o /path/to/output/folder
typedef testing::TestWithParam<tuple<Backend, Target> > Layer_Test_Convolution_DLDT;
2018-12-07 17:40:34 +08:00
TEST_P(Layer_Test_Convolution_DLDT, Accuracy)
{
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
2020-07-30 23:04:22 +08:00
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_API);
else if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NGRAPH);
else
FAIL() << "Unknown backendId";
2018-12-07 17:40:34 +08:00
Net netDefault = readNet(_tf("layer_convolution.caffemodel"), _tf("layer_convolution.prototxt"));
2020-07-30 23:04:22 +08:00
Net net = readNet(_tf("layer_convolution.xml"), _tf("layer_convolution.bin"));
Mat inp = blobFromNPY(_tf("blob.npy"));
netDefault.setInput(inp);
netDefault.setPreferableBackend(DNN_BACKEND_OPENCV);
Mat outDefault = netDefault.forward();
net.setInput(inp);
net.setPreferableBackend(backendId);
2018-12-07 17:40:34 +08:00
net.setPreferableTarget(targetId);
Mat out = net.forward();
double l1 = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? 1.5e-3 : 1e-5;
double lInf = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? 1.8e-2 : 1e-4;
normAssert(outDefault, out, "", l1, lInf);
std::vector<int> outLayers = net.getUnconnectedOutLayers();
ASSERT_EQ(net.getLayer(outLayers[0])->name, "output");
2020-07-30 23:04:22 +08:00
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
ASSERT_EQ(net.getLayer(outLayers[0])->type, "Convolution");
else
ASSERT_EQ(net.getLayer(outLayers[0])->type, "Add");
}
2018-12-07 17:40:34 +08:00
TEST_P(Layer_Test_Convolution_DLDT, setInput_uint8)
{
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
2020-07-30 23:04:22 +08:00
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_API);
else if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NGRAPH);
else
FAIL() << "Unknown backendId";
int blobSize[] = {2, 6, 75, 113};
Mat inputs[] = {Mat(4, &blobSize[0], CV_8U), Mat()};
randu(inputs[0], 0, 255);
inputs[0].convertTo(inputs[1], CV_32F);
Mat outs[2];
for (int i = 0; i < 2; ++i)
{
2020-07-30 23:04:22 +08:00
Net net = readNet(_tf("layer_convolution.xml"), _tf("layer_convolution.bin"));
net.setPreferableBackend(backendId);
2018-12-07 17:40:34 +08:00
net.setPreferableTarget(targetId);
net.setInput(inputs[i]);
outs[i] = net.forward();
ASSERT_EQ(outs[i].type(), CV_32F);
}
2018-12-07 17:40:34 +08:00
if (targetId != DNN_TARGET_MYRIAD)
normAssert(outs[0], outs[1]);
}
TEST_P(Layer_Test_Convolution_DLDT, multithreading)
{
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
2020-07-30 23:04:22 +08:00
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_API);
else if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NGRAPH);
else
FAIL() << "Unknown backendId";
2020-07-30 23:04:22 +08:00
std::string xmlPath = _tf("layer_convolution.xml");
std::string binPath = _tf("layer_convolution.bin");
Net firstNet = readNet(xmlPath, binPath);
Net secondNet = readNet(xmlPath, binPath);
Mat inp = blobFromNPY(_tf("blob.npy"));
firstNet.setInput(inp);
secondNet.setInput(inp);
firstNet.setPreferableBackend(backendId);
firstNet.setPreferableTarget(targetId);
secondNet.setPreferableBackend(backendId);
secondNet.setPreferableTarget(targetId);
Mat out1, out2;
std::thread t1([&]{out1 = firstNet.forward();});
std::thread t2([&]{out2 = secondNet.forward();});
t1.join();
t2.join();
Mat ref = blobFromNPY(_tf("layer_convolution.npy"));
double l1 = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? 1.5e-3 : 1e-5;
double lInf = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? 1.8e-2 : 1e-4;
normAssert(out1, ref, "first thread", l1, lInf);
normAssert(out2, ref, "second thread", l1, lInf);
}
2018-12-07 17:40:34 +08:00
INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_Convolution_DLDT,
dnnBackendsAndTargetsIE()
);
// 1. Create a .prototxt file with the following network:
// layer {
// type: "Input" name: "data" top: "data"
// input_param { shape { dim: 1 dim: 2 dim: 3 } }
// }
// layer {
// type: "Input" name: "second_input" top: "second_input"
// input_param { shape { dim: 1 dim: 2 dim: 3 } }
// }
// layer {
// type: "Eltwise" name: "output" top: "output"
// bottom: "data" bottom: "second_input"
// eltwise_param { operation: SUM }
// }
//
// 2. Create a .caffemodel file using Caffe:
//
// import caffe
// net = caffe.Net('/path/to/prototxt', caffe.TEST)
// net.save('/path/to/caffemodel')
//
// 3. Convert using ModelOptimizer.
2019-01-11 01:29:44 +08:00
typedef testing::TestWithParam<tuple<int, int, Target, std::vector<int> > > Test_DLDT_two_inputs_3dim;
TEST_P(Test_DLDT_two_inputs_3dim, as_IR)
{
int firstInpType = get<0>(GetParam());
int secondInpType = get<1>(GetParam());
2018-12-07 17:40:34 +08:00
Target targetId = get<2>(GetParam());
2020-07-30 23:04:22 +08:00
Net net = readNet(_tf("net_two_inputs.xml"), _tf("net_two_inputs.bin"));
2019-01-11 01:29:44 +08:00
std::vector<int> inpSize = get<3>(GetParam());
Mat firstInp(3, inpSize.data(), firstInpType);
Mat secondInp(3, inpSize.data(), secondInpType);
randu(firstInp, 0, 255);
randu(secondInp, 0, 255);
net.setInput(firstInp, "data");
net.setInput(secondInp, "second_input");
2018-12-07 17:40:34 +08:00
net.setPreferableTarget(targetId);
double l1 = ((targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) &&
(firstInpType == CV_32F || secondInpType == CV_32F)) ? 0.06 : 0.0;
double lInf = ((targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) &&
(firstInpType == CV_32F || secondInpType == CV_32F)) ? 0.23 : 0.0;
Mat out = net.forward();
Mat ref;
cv::add(firstInp, secondInp, ref, Mat(), CV_32F);
normAssert(out, ref, "", l1, lInf);
}
2019-01-11 01:29:44 +08:00
std::vector< std::vector<int> > list_sizes{ {1, 2, 3}, {3, 2, 1}, {5, 5, 5}, {13, 7, 11} };
INSTANTIATE_TEST_CASE_P(/*nothing*/, Test_DLDT_two_inputs_3dim, Combine(
Values(CV_8U, CV_32F), Values(CV_8U, CV_32F),
testing::ValuesIn(getAvailableTargets(DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)),
2019-01-11 01:29:44 +08:00
testing::ValuesIn(list_sizes)
));
typedef testing::TestWithParam<tuple<int, int, tuple<Backend, Target> > > Test_DLDT_two_inputs;
TEST_P(Test_DLDT_two_inputs, as_backend)
{
static const float kScale = 0.5f;
static const float kScaleInv = 1.0f / kScale;
Backend backendId = get<0>(get<2>(GetParam()));
Target targetId = get<1>(get<2>(GetParam()));
2018-12-07 17:40:34 +08:00
Net net;
LayerParams lp;
lp.type = "Eltwise";
lp.name = "testLayer";
lp.set("operation", "sum");
int eltwiseId = net.addLayerToPrev(lp.name, lp.type, lp); // connect to a first input
net.connect(0, 1, eltwiseId, 1); // connect to a second input
2018-12-07 17:40:34 +08:00
int inpSize[] = {1, 2, 3, 4};
Mat firstInp(4, &inpSize[0], get<0>(GetParam()));
Mat secondInp(4, &inpSize[0], get<1>(GetParam()));
randu(firstInp, 0, 255);
randu(secondInp, 0, 255);
net.setInputsNames({"data", "second_input"});
net.setInput(firstInp, "data", kScale);
net.setInput(secondInp, "second_input", kScaleInv);
net.setPreferableBackend(backendId);
2018-12-07 17:40:34 +08:00
net.setPreferableTarget(targetId);
Mat out = net.forward();
Mat ref;
addWeighted(firstInp, kScale, secondInp, kScaleInv, 0, ref, CV_32F);
2018-12-07 17:40:34 +08:00
// Output values are in range [0, 637.5].
double l1 = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? 0.06 : 1e-6;
double lInf = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? 0.3 : 1e-5;
if (targetId == DNN_TARGET_CUDA_FP16)
{
l1 = 0.06;
lInf = 0.3;
}
2018-12-07 17:40:34 +08:00
normAssert(out, ref, "", l1, lInf);
}
INSTANTIATE_TEST_CASE_P(/*nothing*/, Test_DLDT_two_inputs, Combine(
2018-12-07 17:40:34 +08:00
Values(CV_8U, CV_32F), Values(CV_8U, CV_32F),
dnnBackendsAndTargets()
));
class UnsupportedLayer : public Layer
{
public:
UnsupportedLayer(const LayerParams &params) : Layer(params) {}
static Ptr<Layer> create(const LayerParams& params)
{
return Ptr<Layer>(new UnsupportedLayer(params));
}
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
return backendId == DNN_BACKEND_OPENCV;
}
virtual void forward(cv::InputArrayOfArrays inputs, cv::OutputArrayOfArrays outputs, cv::OutputArrayOfArrays internals) CV_OVERRIDE {}
};
typedef DNNTestLayer Test_DLDT_layers;
static void test_dldt_fused_output(Backend backend, Target target)
{
static const int kNumChannels = 3;
Net net;
{
LayerParams lp;
lp.set("kernel_size", 1);
lp.set("num_output", 3);
lp.set("bias_term", false);
lp.type = "Convolution";
lp.name = "testConv";
lp.blobs.push_back(Mat({kNumChannels, 1, 1, 1}, CV_32F, Scalar(1)));
net.addLayerToPrev(lp.name, lp.type, lp);
}
{
LayerParams lp;
lp.set("bias_term", false);
lp.type = "Scale";
lp.name = "testScale";
lp.blobs.push_back(Mat({kNumChannels}, CV_32F, Scalar(1)));
net.addLayerToPrev(lp.name, lp.type, lp);
}
{
LayerParams lp;
net.addLayerToPrev("unsupported_layer", "Unsupported", lp);
}
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
net.setInput(Mat({1, 1, 2, 3}, CV_32FC1, Scalar(1)));
net.forward();
}
TEST_P(Test_DLDT_layers, fused_output)
{
CV_DNN_REGISTER_LAYER_CLASS(Unsupported, UnsupportedLayer);
try
{
test_dldt_fused_output(backend, target);
}
catch (const std::exception& e)
{
ADD_FAILURE() << "Exception: " << e.what();
}
catch(...)
{
ADD_FAILURE() << "Unknown exception";
}
LayerFactory::unregisterLayer("Unsupported");
}
TEST_P(Test_DLDT_layers, multiple_networks)
{
Net nets[2];
for (int i = 0; i < 2; ++i)
{
nets[i].setInputsNames(std::vector<String>(1, format("input_%d", i)));
LayerParams lp;
lp.set("kernel_size", 1);
lp.set("num_output", 1);
lp.set("bias_term", false);
lp.type = "Convolution";
lp.name = format("testConv_%d", i);
lp.blobs.push_back(Mat({1, 1, 1, 1}, CV_32F, Scalar(1 + i)));
nets[i].addLayerToPrev(lp.name, lp.type, lp);
nets[i].setPreferableBackend(backend);
nets[i].setPreferableTarget(target);
nets[i].setInput(Mat({1, 1, 2, 3}, CV_32FC1, Scalar(1)));
}
Mat out_1 = nets[0].forward();
Mat out_2 = nets[1].forward();
// After the second model is initialized we try to receive an output from the first network again.
out_1 = nets[0].forward();
normAssert(2 * out_1, out_2);
}
INSTANTIATE_TEST_CASE_P(/*nothing*/, Test_DLDT_layers, dnnBackendsAndTargets());
#endif // HAVE_INF_ENGINE
// Test a custom layer.
2018-06-27 21:34:36 +08:00
class CustomInterpLayer CV_FINAL : public Layer
{
public:
2018-06-27 21:34:36 +08:00
CustomInterpLayer(const LayerParams &params) : Layer(params)
{
zoomFactor = params.get<int>("zoom_factor", 0);
outWidth = params.get<int>("width", 0);
outHeight = params.get<int>("height", 0);
}
2018-06-27 21:34:36 +08:00
static Ptr<Layer> create(LayerParams& params)
{
2018-06-27 21:34:36 +08:00
return Ptr<Layer>(new CustomInterpLayer(params));
}
virtual bool getMemoryShapes(const std::vector<std::vector<int> > &inputs,
const int requiredOutputs,
std::vector<std::vector<int> > &outputs,
std::vector<std::vector<int> > &internals) const CV_OVERRIDE
{
const int batchSize = inputs[0][0];
const int numChannels = inputs[0][1];
const int inpHeight = inputs[0][2];
const int inpWidth = inputs[0][3];
std::vector<int> outShape(4);
outShape[0] = batchSize;
outShape[1] = numChannels;
outShape[2] = outHeight != 0 ? outHeight : (inpHeight + (inpHeight - 1) * (zoomFactor - 1));
outShape[3] = outWidth != 0 ? outWidth : (inpWidth + (inpWidth - 1) * (zoomFactor - 1));
outputs.assign(1, outShape);
return false;
}
virtual void finalize(InputArrayOfArrays, OutputArrayOfArrays outputs_arr) CV_OVERRIDE
{
std::vector<Mat> outputs;
outputs_arr.getMatVector(outputs);
if (!outWidth && !outHeight)
{
outHeight = outputs[0].size[2];
outWidth = outputs[0].size[3];
}
}
// Implementation of this custom layer is based on https://github.com/cdmh/deeplab-public/blob/master/src/caffe/layers/interp_layer.cpp
void forward(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr, OutputArrayOfArrays internals_arr) CV_OVERRIDE
{
CV_TRACE_FUNCTION();
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
Mat& inp = inputs[0];
Mat& out = outputs[0];
const float* inpData = (float*)inp.data;
float* outData = (float*)out.data;
const int batchSize = inp.size[0];
const int numChannels = inp.size[1];
const int inpHeight = inp.size[2];
const int inpWidth = inp.size[3];
const float rheight = (outHeight > 1) ? static_cast<float>(inpHeight - 1) / (outHeight - 1) : 0.f;
const float rwidth = (outWidth > 1) ? static_cast<float>(inpWidth - 1) / (outWidth - 1) : 0.f;
for (int h2 = 0; h2 < outHeight; ++h2)
{
const float h1r = rheight * h2;
const int h1 = h1r;
const int h1p = (h1 < inpHeight - 1) ? 1 : 0;
const float h1lambda = h1r - h1;
const float h0lambda = 1.f - h1lambda;
for (int w2 = 0; w2 < outWidth; ++w2)
{
const float w1r = rwidth * w2;
const int w1 = w1r;
const int w1p = (w1 < inpWidth - 1) ? 1 : 0;
const float w1lambda = w1r - w1;
const float w0lambda = 1.f - w1lambda;
const float* pos1 = inpData + h1 * inpWidth + w1;
float* pos2 = outData + h2 * outWidth + w2;
for (int c = 0; c < batchSize * numChannels; ++c)
{
pos2[0] =
h0lambda * (w0lambda * pos1[0] + w1lambda * pos1[w1p]) +
h1lambda * (w0lambda * pos1[h1p * inpWidth] + w1lambda * pos1[h1p * inpWidth + w1p]);
pos1 += inpWidth * inpHeight;
pos2 += outWidth * outHeight;
}
}
}
}
private:
int outWidth, outHeight, zoomFactor;
};
#ifndef OPENCV_DNN_EXTERNAL_PROTOBUF
2018-06-27 21:34:36 +08:00
TEST_P(Test_Caffe_layers, Interp)
#else
TEST_P(Test_Caffe_layers, DISABLED_Interp) // requires patched protobuf (available in OpenCV source tree only)
#endif
{
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && target == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD);
// Test a custom layer.
2018-06-27 21:34:36 +08:00
CV_DNN_REGISTER_LAYER_CLASS(Interp, CustomInterpLayer);
try
{
testLayerUsingCaffeModels("layer_interp", false, false);
}
catch (...)
{
LayerFactory::unregisterLayer("Interp");
throw;
}
LayerFactory::unregisterLayer("Interp");
2018-06-27 21:34:36 +08:00
// Test an implemented layer.
testLayerUsingCaffeModels("layer_interp", false, false);
}
2018-06-27 21:34:36 +08:00
INSTANTIATE_TEST_CASE_P(/*nothing*/, Test_Caffe_layers, dnnBackendsAndTargets());
TEST(Layer_Test_PoolingIndices, Accuracy)
{
Net net;
LayerParams lp;
lp.set("pool", "max");
lp.set("kernel_w", 2);
lp.set("kernel_h", 2);
lp.set("stride_w", 2);
lp.set("stride_h", 2);
lp.set("pad_w", 0);
lp.set("pad_h", 0);
lp.name = "testLayer.name"; // This test also checks that OpenCV lets use names with dots.
lp.type = "Pooling";
net.addLayerToPrev(lp.name, lp.type, lp);
Mat inp(10, 10, CV_8U);
randu(inp, 0, 255);
Mat maxValues(5, 5, CV_32F, Scalar(-1)), indices(5, 5, CV_32F, Scalar(-1));
for (int y = 0; y < 10; ++y)
{
int dstY = y / 2;
for (int x = 0; x < 10; ++x)
{
int dstX = x / 2;
uint8_t val = inp.at<uint8_t>(y, x);
if ((float)inp.at<uint8_t>(y, x) > maxValues.at<float>(dstY, dstX))
{
maxValues.at<float>(dstY, dstX) = val;
indices.at<float>(dstY, dstX) = y * 10 + x;
}
}
}
2018-06-26 20:38:08 +08:00
net.setPreferableBackend(DNN_BACKEND_OPENCV);
net.setInput(blobFromImage(inp));
std::vector<Mat> outputs;
net.forward(outputs, lp.name);
normAssert(maxValues, outputs[0].reshape(1, 5));
normAssert(indices, outputs[1].reshape(1, 5));
}
typedef testing::TestWithParam<tuple<Vec4i, int, tuple<Backend, Target> > > Layer_Test_ShuffleChannel;
2018-06-19 19:35:07 +08:00
TEST_P(Layer_Test_ShuffleChannel, Accuracy)
{
Vec4i inpShapeVec = get<0>(GetParam());
int group = get<1>(GetParam());
ASSERT_EQ(inpShapeVec[1] % group, 0);
const int groupSize = inpShapeVec[1] / group;
int backendId = get<0>(get<2>(GetParam()));
int targetId = get<1>(get<2>(GetParam()));
2018-06-19 19:35:07 +08:00
Net net;
LayerParams lp;
lp.set("group", group);
lp.type = "ShuffleChannel";
lp.name = "testLayer";
net.addLayerToPrev(lp.name, lp.type, lp);
const int inpShape[] = {inpShapeVec[0], inpShapeVec[1], inpShapeVec[2], inpShapeVec[3]};
Mat inp(4, inpShape, CV_32F);
randu(inp, 0, 255);
net.setInput(inp);
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
2018-06-19 19:35:07 +08:00
Mat out = net.forward();
double l1 = 1e-5, lInf = 1e-4;
if (targetId == DNN_TARGET_OPENCL_FP16)
{
l1 = 5e-2;
lInf = 7e-2;
}
else if (targetId == DNN_TARGET_CUDA_FP16)
{
l1 = 0.06;
lInf = 0.07;
}
2018-06-19 19:35:07 +08:00
for (int n = 0; n < inpShapeVec[0]; ++n)
{
for (int c = 0; c < inpShapeVec[1]; ++c)
{
Mat outChannel = getPlane(out, n, c);
Mat inpChannel = getPlane(inp, n, groupSize * (c % group) + c / group);
normAssert(outChannel, inpChannel, "", l1, lInf);
2018-06-19 19:35:07 +08:00
}
}
}
INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_ShuffleChannel, Combine(
/*input shape*/ Values(Vec4i(1, 6, 5, 7), Vec4i(3, 12, 1, 4)),
/*group*/ Values(1, 2, 3, 6), dnnBackendsAndTargets(/*with IE*/ false)
2018-06-19 19:35:07 +08:00
));
2018-07-04 20:50:39 +08:00
// Check if relu is not fused to convolution if we requested it's output
TEST(Layer_Test_Convolution, relu_fusion)
{
Net net;
{
LayerParams lp;
lp.set("kernel_size", 1);
lp.set("num_output", 1);
lp.set("bias_term", false);
lp.type = "Convolution";
lp.name = "testConv";
int weightsShape[] = {1, 1, 1, 1};
Mat weights(4, &weightsShape[0], CV_32F, Scalar(1));
lp.blobs.push_back(weights);
net.addLayerToPrev(lp.name, lp.type, lp);
}
{
LayerParams lp;
lp.type = "ReLU";
lp.name = "testReLU";
net.addLayerToPrev(lp.name, lp.type, lp);
}
int sz[] = {1, 1, 2, 3};
Mat input(4, &sz[0], CV_32F);
randu(input, -1.0, -0.1);
net.setInput(input);
net.setPreferableBackend(DNN_BACKEND_OPENCV);
Mat output = net.forward("testConv");
normAssert(input, output);
}
typedef testing::TestWithParam<tuple<bool, tuple<Backend, Target> > > Layer_Test_Eltwise_unequal;
TEST_P(Layer_Test_Eltwise_unequal, accuracy_input_0_truncate)
{
bool weighted = get<0>(GetParam());
int backendId = get<0>(get<1>(GetParam()));
int targetId = get<1>(get<1>(GetParam()));
if (backendId == DNN_BACKEND_CUDA && weighted)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA);
Net net;
LayerParams lp;
lp.type = "Eltwise";
lp.name = "testLayer";
lp.set<std::string>("output_channels_mode", "input_0_truncate");
const int inpShapes[][4] = {{1, 4, 2, 2}, {1, 5, 2, 2}, {1, 3, 2, 2}};
const int out_channels = inpShapes[0][1];
std::vector<String> inpNames(3);
std::vector<Mat> inputs(3);
std::vector<float> weights(3, 1);
if (weighted)
{
for (int i = 0; i < inputs.size(); ++i)
weights[i] = -0.125f + i * 0.25f;
lp.set("coeff", DictValue::arrayReal<float*>(&weights[0], weights.size()));
}
int eltwiseId = net.addLayer(lp.name, lp.type, lp);
for (int i = 0; i < inputs.size(); ++i)
{
inputs[i].create(4, inpShapes[i], CV_32F);
size_t total = inputs[i].total();
for (size_t j = 0; j < total; j++)
inputs[i].ptr<float>()[j] = j + i * 100;
inpNames[i] = format("input_%d", i);
net.connect(0, i, eltwiseId, i);
}
Mat ref(4, inpShapes[0], CV_32F, Scalar(0));
net.setInputsNames(inpNames);
for (int i = 0; i < inputs.size(); ++i)
{
//std::cout << ref.reshape(1,1) << endl;
net.setInput(inputs[i], inpNames[i]);
for (size_t batchId = 0; batchId < ref.size[0]; batchId++)
{
int input_channels = inputs[i].size[1];
Range ranges[4] = { Range(batchId, batchId + 1), Range(0, std::min(out_channels, input_channels)), Range::all(), Range::all() };
Mat ref_slice = ref(ranges);
Mat input_slice = inputs[i](ranges);
ref_slice += weights[i] * input_slice;
}
}
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
Mat out = net.forward();
normAssert(out, ref);
if (testing::Test::HasFailure())
{
std::cout << out.reshape(1,1) << endl;
std::cout << ref.reshape(1,1) << endl;
}
}
TEST_P(Layer_Test_Eltwise_unequal, accuracy_input_0)
{
bool weighted = get<0>(GetParam());
int backendId = get<0>(get<1>(GetParam()));
int targetId = get<1>(get<1>(GetParam()));
Net net;
LayerParams lp;
lp.type = "Eltwise";
lp.name = "testLayer";
lp.set<std::string>("output_channels_mode", "input_0");
if (backendId == DNN_BACKEND_CUDA && weighted)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA);
const int inpShapes[][4] = {{1, 4, 2, 2}, {1, 2, 2, 2}, {1, 3, 2, 2}};
const int out_channels = inpShapes[0][1];
std::vector<String> inpNames(3);
std::vector<Mat> inputs(3);
std::vector<float> weights(3, 1);
if (weighted)
{
for (int i = 0; i < inputs.size(); ++i)
weights[i] = -0.125f + i * 0.25f;
lp.set("coeff", DictValue::arrayReal<float*>(&weights[0], weights.size()));
}
int eltwiseId = net.addLayer(lp.name, lp.type, lp);
for (int i = 0; i < inputs.size(); ++i)
{
inputs[i].create(4, inpShapes[i], CV_32F);
size_t total = inputs[i].total();
for (size_t j = 0; j < total; j++)
inputs[i].ptr<float>()[j] = j + i * 100;
inpNames[i] = format("input_%d", i);
net.connect(0, i, eltwiseId, i);
}
Mat ref(4, inpShapes[0], CV_32F, Scalar(0));
net.setInputsNames(inpNames);
for (int i = 0; i < inputs.size(); ++i)
{
//std::cout << ref.reshape(1,1) << endl;
net.setInput(inputs[i], inpNames[i]);
for (size_t batchId = 0; batchId < ref.size[0]; batchId++)
{
int input_channels = inputs[i].size[1];
Range ranges[4] = { Range(batchId, batchId + 1), Range(0, std::min(out_channels, input_channels)), Range::all(), Range::all() };
Mat ref_slice = ref(ranges);
Mat input_slice = inputs[i](ranges);
ref_slice += weights[i] * input_slice;
}
}
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
Mat out = net.forward();
normAssert(out, ref);
if (testing::Test::HasFailure())
{
std::cout << out.reshape(1,1) << endl;
std::cout << ref.reshape(1,1) << endl;
}
}
INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_Eltwise_unequal, Combine(
testing::Bool(),
dnnBackendsAndTargets()
));
typedef testing::TestWithParam<tuple<Backend, Target> > Layer_Test_Resize;
TEST_P(Layer_Test_Resize, change_input)
{
int backendId = get<0>(GetParam());
int targetId = get<1>(GetParam());
Net net;
LayerParams lp;
lp.type = "Resize";
lp.name = "testLayer";
lp.set("zoom_factor", 2);
lp.set("interpolation", "nearest");
net.addLayerToPrev(lp.name, lp.type, lp);
for (int i = 0; i < 2; ++i)
{
Mat inp(4 + i, 5 + i, CV_8UC3), ref;
randu(inp, 0, 255);
resize(inp, ref, Size(0, 0), 2, 2, INTER_NEAREST);
ref = blobFromImage(ref);
net.setInput(blobFromImage(inp));
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
Mat out = net.forward();
normAssert(out, ref);
}
}
INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_Resize, dnnBackendsAndTargets());
struct Layer_Test_Slice : public testing::TestWithParam<tuple<Backend, Target> >
{
template<int DIMS>
void test_slice(const int* inputShape, const int* begin, const int* end)
{
int backendId = get<0>(GetParam());
int targetId = get<1>(GetParam());
Mat input(DIMS, inputShape, CV_32FC1, Scalar::all(0));
for (int i = 0; i < (int)input.total(); ++i)
input.ptr<float>()[i] = (float)i;
std::vector<Range> range(DIMS);
for (int i = 0; i < DIMS; ++i)
range[i] = Range(begin[i], end[i]);
Net net;
LayerParams lp;
lp.type = "Slice";
lp.name = "testLayer";
lp.set("begin", DictValue::arrayInt<int*>((int*)&begin[0], DIMS));
lp.set("end", DictValue::arrayInt<int*>((int*)&end[0], DIMS));
net.addLayerToPrev(lp.name, lp.type, lp);
{
net.setInput(input);
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
Mat out = net.forward();
EXPECT_GT(cv::norm(out, NORM_INF), 0);
normAssert(out, input(range));
#if 0
cout << input(range).clone().reshape(1, 1) << endl;
cout << out.reshape(1, 1) << endl;
#endif
}
}
};
TEST_P(Layer_Test_Slice, slice_channels_17762)
{
const int inputShape[4] = {1, 16, 6, 8};
const int begin[] = {0, 4, 0, 0};
const int end[] = {1, 8, 6, 8};
test_slice<4>(inputShape, begin, end);
}
TEST_P(Layer_Test_Slice, slice_channels_with_batch_17762)
{
const int inputShape[4] = {4, 4, 3, 4};
const int begin[] = {0, 1, 0, 0};
const int end[] = {4, 3, 3, 4};
test_slice<4>(inputShape, begin, end);
}
TEST_P(Layer_Test_Slice, slice_channels_and_batch_17762)
{
int backend = get<0>(GetParam());
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER, CV_TEST_TAG_DNN_SKIP_IE_VERSION);
const int inputShape[4] = {4, 4, 3, 4};
const int begin[] = {2, 1, 0, 0};
const int end[] = {4, 3, 3, 4};
test_slice<4>(inputShape, begin, end);
}
TEST_P(Layer_Test_Slice, slice_rows)
{
const int inputShape[4] = {1, 2, 6, 4};
const int begin[] = {0, 0, 4, 0};
const int end[] = {1, 2, 6, 4};
test_slice<4>(inputShape, begin, end);
}
TEST_P(Layer_Test_Slice, slice_cols)
{
const int inputShape[4] = {1, 2, 3, 8};
const int begin[] = {0, 0, 0, 4};
const int end[] = {1, 2, 3, 8};
test_slice<4>(inputShape, begin, end);
}
TEST_P(Layer_Test_Slice, slice_complex_1_unaligned)
{
const int inputShape[4] = {1, 4, 2, 3};
const int begin[] = {0, 2, 1, 0};
const int end[] = {1, 3, 2, 2};
test_slice<4>(inputShape, begin, end);
}
TEST_P(Layer_Test_Slice, slice_complex_2_x4)
{
const int inputShape[4] = {1, 3, 2, 4};
const int begin[] = {0, 2, 1, 0};
const int end[] = {1, 3, 2, 2};
test_slice<4>(inputShape, begin, end);
}
TEST_P(Layer_Test_Slice, slice_complex_3)
{
const int inputShape[4] = {1, 6, 4, 8};
const int begin[] = {0, 2, 1, 4};
const int end[] = {1, 4, 3, 8};
test_slice<4>(inputShape, begin, end);
}
2020-05-05 18:07:26 +08:00
TEST_P(Layer_Test_Slice, variable_input_shape)
{
int backendId = get<0>(GetParam());
int targetId = get<1>(GetParam());
int begin[] = {0, 0, 0, 0};
int end[] = {-1, -1, -1, -1};
Net net;
LayerParams lp;
lp.type = "Slice";
lp.name = "testLayer";
lp.set("begin", DictValue::arrayInt<int*>(&begin[0], 4));
lp.set("end", DictValue::arrayInt<int*>(&end[0], 4));
net.addLayerToPrev(lp.name, lp.type, lp);
for (int i = 0; i < 2; ++i)
{
Mat inp(4 + i, 5 + i, CV_8UC1);
randu(inp, 0, 255);
inp = blobFromImage(inp);
net.setInput(inp);
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
Mat out = net.forward();
normAssert(out, inp);
}
}
INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_Slice, dnnBackendsAndTargets());
typedef testing::TestWithParam<tuple<Backend, Target> > Layer_Test_BatchNorm;
TEST_P(Layer_Test_BatchNorm, fusion)
{
// This tests reinitializes network by forwarding different batch size input.
// We check BatchNorm layer weights restoring after fusion.
int backendId = get<0>(GetParam());
int targetId = get<1>(GetParam());
const int ch = 4;
Mat mean(1, ch, CV_32F), var(1, ch, CV_32F), weights(1, ch, CV_32F);
randu(mean, 0, 1);
randu(var, 0, 1);
randu(weights, 0, 1);
Net net;
{
LayerParams lp;
lp.type = "BatchNorm";
lp.name = "bn";
lp.set("has_weight", false);
lp.set("has_bias", false);
lp.blobs.push_back(mean);
lp.blobs.push_back(var);
net.addLayerToPrev(lp.name, lp.type, lp);
}
{
LayerParams lp;
lp.type = "Scale";
lp.name = "scale";
lp.set("has_bias", false);
lp.blobs.push_back(weights);
net.addLayerToPrev(lp.name, lp.type, lp);
}
Mat inp(4, 5, CV_32FC(ch));
randu(inp, 0, 1);
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
net.setInput(blobFromImage(inp));
Mat ref = net.forward();
net.setInput(blobFromImages(std::vector<Mat>(2, inp)));
Mat out = net.forward();
for (int i = 0; i < 2; ++i)
{
std::vector<Range> ranges(4, Range::all());
ranges[0].start = i;
ranges[0].end = i + 1;
normAssert(out(ranges), ref);
}
}
INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_BatchNorm, dnnBackendsAndTargets());
class TestLayerFusion : public DNNTestLayer {
public:
static void makeDefaultTestConvolutionLayer(LayerParams& convParams, int in_channels, int num_filters, bool bias_term)
{
const int kernel_h = 3, kernel_w = 3;
const int pad_h = kernel_h / 2, pad_w = kernel_w / 2;
convParams.set("kernel_h", kernel_h);
convParams.set("kernel_w", kernel_w);
convParams.set("pad_h", pad_h);
convParams.set("pad_w", pad_w);
convParams.set("num_output", num_filters);
convParams.set("bias_term", bias_term);
convParams.type = "Convolution";
convParams.name = "convolution";
float conv_init_magnitude = 1.0f / in_channels / kernel_h / kernel_w;
int weightsShape[] = {num_filters, in_channels, kernel_h, kernel_w};
Mat weights(4, &weightsShape[0], CV_32F);
randu(weights, -conv_init_magnitude, conv_init_magnitude);
convParams.blobs.push_back(weights);
if (bias_term)
{
Mat bias(1, num_filters, CV_32F);
randu(bias, -1.0f, 1.0f);
convParams.blobs.push_back(bias);
}
}
static void makeDefaultTestActivationLayer(LayerParams& activationParams, const std::string& type, int in_channels)
{
activationParams.type = type;
activationParams.name = "activation";
if (activationParams.type == "ReLU")
activationParams.set("negative_slope", 0.1f);
else if (activationParams.type == "Power")
{
activationParams.set("power", 2.0f);
activationParams.set("scale", 0.5f);
activationParams.set("shift", 0.3f);
}
else if (activationParams.type == "ReLU6")
{
activationParams.set("min_value", -1.0f);
activationParams.set("max_value", 1.0f);
}
else if (activationParams.type == "ChannelsPReLU")
{
Mat scales(1, in_channels, CV_32F);
randu(scales, -1.0f, 1.0f);
activationParams.blobs.push_back(scales);
}
}
static void makeDefaultTestEltwiseLayer(LayerParams& eltwiseParams, const std::string& op, bool withCoefficients)
{
eltwiseParams.type = "Eltwise";
eltwiseParams.name = "eltwise";
eltwiseParams.set("operation", op);
if (withCoefficients)
{
float coeff[] = {0.3f, 0.5f};
eltwiseParams.set("coeff", DictValue::arrayReal<float*>(coeff, 2));
}
}
static void test(Mat& input, Net& net, Backend backendId, Target targetId, std::vector<int> expectedFusedLayers = std::vector<int>(), double l1 = 0.0, double lInf = 0.0)
{
DNNTestLayer::checkBackend(backendId, targetId);
net.enableFusion(false);
net.setPreferableBackend(DNN_BACKEND_OPENCV);
net.setPreferableTarget(DNN_TARGET_CPU);
net.setInput(input);
Mat outputReference = net.forward().clone();
std::vector<double> refTimings;
net.getPerfProfile(refTimings);
for (int i = 0; i < refTimings.size(); i++)
{
CV_Assert(refTimings[i] != 0.0);
}
net.enableFusion(true);
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
net.setInput(input);
Mat outputTest = net.forward().clone();
std::vector<double> testTimings;
net.getPerfProfile(testTimings);
for (int i = 0; i < testTimings.size(); i++)
{
if(std::find(expectedFusedLayers.begin(), expectedFusedLayers.end(), i + 1) != expectedFusedLayers.end())
{
EXPECT_EQ(testTimings[i], 0.0);
}
else
{
EXPECT_NE(testTimings[i], 0.0);
}
}
// double ref_max_value, ref_min_value;
// minMaxLoc(outputReference.reshape(1, 1), &ref_min_value, &ref_max_value);
// std::cout << "reference range: " << ref_min_value << ' ' << ref_max_value << std::endl;
double default_l1, default_lInf;
DNNTestLayer::getDefaultThresholds(backendId, targetId, &default_l1, &default_lInf);
if (l1 == 0.0)
l1 = default_l1;
if (lInf == 0.0)
lInf = default_lInf;
normAssert(outputReference, outputTest, "", l1, lInf);
}
static testing::internal::ParamGenerator<std::string> eltwiseOpList()
{
// TODO: automate list generation
return Values("sum", "max", "prod", "div");
}
static testing::internal::ParamGenerator<std::string> activationLayersList()
{
// TODO: automate list generation
return Values("ReLU", "ReLU6", "ChannelsPReLU", "TanH", "Swish", "Mish", "Sigmoid", "ELU", "AbsVal", "BNLL", "Power");
}
static testing::internal::ParamGenerator<tuple<Backend, Target> > dnnBackendsAndTargetsForFusionTests()
{
return dnnBackendsAndTargets(false, false, true, false, true, false); // OCV OpenCL + OCV CPU + CUDA
}
};
typedef TestWithParam<tuple<bool, std::string, tuple<Backend, Target> > > ConvolutionActivationFusion;
TEST_P(ConvolutionActivationFusion, Accuracy)
{
// input
// |
// -----------------------
// | convolution |
// -----------------------
// |
// -----------------------
// | activation |
// -----------------------
// |
// output
const int batch_size = 2, in_channels = 16;
const int in_height = 16, in_width = 16;
int inputShape[] = {batch_size, in_channels, in_height, in_width};
Mat input(4, &inputShape[0], CV_32F);
randu(input, 1.0f, 2.0f);
bool bias_term = get<0>(GetParam());
LayerParams convParams;
TestLayerFusion::makeDefaultTestConvolutionLayer(convParams, in_channels, in_channels, bias_term);
std::string actType = get<1>(GetParam());
LayerParams activationParams;
TestLayerFusion::makeDefaultTestActivationLayer(activationParams, actType, in_channels);
Backend backendId = get<0>(get<2>(GetParam()));
Target targetId = get<1>(get<2>(GetParam()));
Net net;
int convId = net.addLayer(convParams.name, convParams.type, convParams);
int activId = net.addLayerToPrev(activationParams.name, activationParams.type, activationParams);
net.connect(0, 0, convId, 0);
std::vector<int> expectedFusedLayers;
if (backendId == DNN_BACKEND_OPENCV)
{
if (targetId == DNN_TARGET_CPU)
expectedFusedLayers.push_back(activId); // all activations are fused
else if (targetId == DNN_TARGET_OPENCL || targetId == DNN_TARGET_OPENCL_FP16)
{
if (actType == "ReLU" || actType == "ChannelsPReLU" || actType == "ReLU6" || actType == "TanH" /*|| actType == "Power"*/)
expectedFusedLayers.push_back(activId);
}
}
else if (backendId == DNN_BACKEND_CUDA)
{
if (actType == "ReLU" || actType == "ReLU6" || actType == "TanH" || actType == "Swish" ||
actType == "Mish" || actType == "Sigmoid" || actType == "Power")
expectedFusedLayers.push_back(activId);
}
TestLayerFusion::test(input, net, backendId, targetId, expectedFusedLayers);
}
INSTANTIATE_TEST_CASE_P(TestLayerFusion, ConvolutionActivationFusion, Combine(
/* bias */ testing::Bool(),
/* activation */ TestLayerFusion::activationLayersList(),
TestLayerFusion::dnnBackendsAndTargetsForFusionTests()
));
typedef TestWithParam<tuple<bool, std::string, bool, tuple<Backend, Target> > > ConvolutionEltwiseFusion;
TEST_P(ConvolutionEltwiseFusion, Accuracy)
{
// input
// |
// -------------------------------
// | |
// | ---------------
// | | convolution |
// | ---------------
// | |
// | ---------------- |
// --------| eltwise op |-------
// ----------------
// |
// output
const int batch_size = 2, in_channels = 16;
const int in_height = 16, in_width = 16;
int inputShape[] = {batch_size, in_channels, in_height, in_width};
Mat input(4, &inputShape[0], CV_32F);
randu(input, 1.0f, 2.0f); // avoid small values to test eltwise div
bool bias_term = get<0>(GetParam());
LayerParams convParams;
TestLayerFusion::makeDefaultTestConvolutionLayer(convParams, in_channels, in_channels, bias_term);
std::string eltwiseOp = get<1>(GetParam());
bool weightedEltwise = get<2>(GetParam());
if (eltwiseOp != "sum" && weightedEltwise)
throw SkipTestException("weighted eltwise not supported");
LayerParams eltwiseParams;
TestLayerFusion::makeDefaultTestEltwiseLayer(eltwiseParams, eltwiseOp, weightedEltwise);
Net net;
int convId = net.addLayer(convParams.name, convParams.type, convParams);
int eltwiseId = net.addLayer(eltwiseParams.name, eltwiseParams.type, eltwiseParams);
net.connect(0, 0, convId, 0);
net.connect(convId, 0, eltwiseId, 0);
net.connect(0, 0, eltwiseId, 1);
Backend backendId = get<0>(get<3>(GetParam()));
Target targetId = get<1>(get<3>(GetParam()));
std::vector<int> expectedFusedLayers;
if (backendId == DNN_BACKEND_CUDA && eltwiseOp == "sum" && !weightedEltwise)
expectedFusedLayers.push_back(eltwiseId);
TestLayerFusion::test(input, net, backendId, targetId, expectedFusedLayers);
}
INSTANTIATE_TEST_CASE_P(TestLayerFusion, ConvolutionEltwiseFusion, Combine(
/* bias */ testing::Bool(),
/* eltwise op */ TestLayerFusion::eltwiseOpList(),
/* eltwise weighted */ testing::Bool(),
TestLayerFusion::dnnBackendsAndTargetsForFusionTests()
));
typedef TestWithParam<tuple<bool, std::string, bool, std::string, tuple<Backend, Target> > > ConvolutionEltwiseActivationFusion;
TEST_P(ConvolutionEltwiseActivationFusion, Accuracy)
{
// input
// |
// -------------------------------
// | |
// | ---------------
// | | convolution |
// | ---------------
// | |
// | ---------------- |
// --------| eltwise op |-------
// ----------------
// |
// ----------------
// | activation |
// ----------------
// |
// output
const int batch_size = 2, in_channels = 16;
const int in_height = 16, in_width = 16;
int inputShape[] = {batch_size, in_channels, in_height, in_width};
Mat input(4, &inputShape[0], CV_32F);
randu(input, 1.0f, 2.0f); // avoid small values to test eltwise div
bool bias_term = get<0>(GetParam());
LayerParams convParams;
TestLayerFusion::makeDefaultTestConvolutionLayer(convParams, in_channels, in_channels, bias_term);
std::string eltwiseOp = get<1>(GetParam());
bool weightedEltwise = get<2>(GetParam());
if (eltwiseOp != "sum" && weightedEltwise)
throw SkipTestException("weighted eltwise not supported");
LayerParams eltwiseParams;
2020-09-02 16:55:36 +08:00
TestLayerFusion::makeDefaultTestEltwiseLayer(eltwiseParams, eltwiseOp, weightedEltwise);
std::string actType = get<3>(GetParam());
LayerParams activationParams;
TestLayerFusion::makeDefaultTestActivationLayer(activationParams, actType, in_channels);
Backend backendId = get<0>(get<4>(GetParam()));
Target targetId = get<1>(get<4>(GetParam()));
Net net;
int convId = net.addLayer(convParams.name, convParams.type, convParams);
int eltwiseId = net.addLayer(eltwiseParams.name, eltwiseParams.type, eltwiseParams);
int activId = net.addLayer(activationParams.name, activationParams.type, activationParams);
net.connect(0, 0, convId, 0);
net.connect(convId, 0, eltwiseId, 0);
net.connect(0, 0, eltwiseId, 1);
net.connect(eltwiseId, 0, activId, 0);
std::vector<int> expectedFusedLayers;
if (backendId == DNN_BACKEND_OPENCV)
{
if (targetId == DNN_TARGET_CPU)
expectedFusedLayers.push_back(activId); // activation is fused with eltwise layer
else if (targetId == DNN_TARGET_OPENCL || targetId == DNN_TARGET_OPENCL_FP16)
{
if (eltwiseOp == "sum" && !weightedEltwise &&
(actType == "ReLU" || actType == "ChannelsPReLU" /*|| actType == "Power"*/)
)
{
expectedFusedLayers.push_back(eltwiseId);
expectedFusedLayers.push_back(activId);
}
}
}
else if(backendId == DNN_BACKEND_CUDA)
{
if (eltwiseOp == "sum" && !weightedEltwise)
{
expectedFusedLayers.push_back(eltwiseId);
if (actType == "ReLU" || actType == "ReLU6" || actType == "TanH" || actType == "Swish" ||
actType == "Mish" || actType == "Sigmoid" || actType == "Power")
expectedFusedLayers.push_back(activId);
}
}
TestLayerFusion::test(input, net, backendId, targetId, expectedFusedLayers);
}
INSTANTIATE_TEST_CASE_P(TestLayerFusion, ConvolutionEltwiseActivationFusion, Combine(
/* bias */ testing::Bool(),
/* eltwise op */ TestLayerFusion::eltwiseOpList(),
/* eltwise weighted */ testing::Bool(),
/* activation */ TestLayerFusion::activationLayersList(),
TestLayerFusion::dnnBackendsAndTargetsForFusionTests()
));
typedef TestWithParam<tuple<bool, std::string, std::string, bool, tuple<Backend, Target> > > ConvolutionActivationEltwiseFusion;
TEST_P(ConvolutionActivationEltwiseFusion, Accuracy)
{
// input
// |
// -------------------------------
// | |
// | ----------------
// | | convolution |
// | ----------------
// | |
// | ----------------
// | | activation |
// | ----------------
// | |
// | ---------------- |
// --------| eltwise sum |-------
// ----------------
// |
const int batch_size = 2, in_channels = 16;
const int in_height = 16, in_width = 16;
int inputShape[] = {batch_size, in_channels, in_height, in_width};
Mat input(4, &inputShape[0], CV_32F);
randu(input, 1.0f, 2.0f); // avoid small values to test eltwise div
bool bias_term = get<0>(GetParam());
LayerParams convParams;
TestLayerFusion::makeDefaultTestConvolutionLayer(convParams, in_channels, in_channels, bias_term);
std::string actType = get<1>(GetParam());
LayerParams activationParams;
TestLayerFusion::makeDefaultTestActivationLayer(activationParams, actType, in_channels);
std::string eltwiseOp = get<2>(GetParam());
bool weightedEltwise = get<3>(GetParam());
if (eltwiseOp != "sum" && weightedEltwise)
throw SkipTestException("weighted eltwise not supported");
LayerParams eltwiseParams;
2020-09-02 16:55:36 +08:00
TestLayerFusion::makeDefaultTestEltwiseLayer(eltwiseParams, eltwiseOp, weightedEltwise);
Backend backendId = get<0>(get<4>(GetParam()));
Target targetId = get<1>(get<4>(GetParam()));
Net net;
int convId = net.addLayer(convParams.name, convParams.type, convParams);
int activId = net.addLayer(activationParams.name, activationParams.type, activationParams);
int eltwiseId = net.addLayer(eltwiseParams.name, eltwiseParams.type, eltwiseParams);
net.connect(0, 0, convId, 0);
net.connect(convId, 0, activId, 0);
net.connect(activId, 0, eltwiseId, 0);
net.connect(0, 0, eltwiseId, 1);
std::vector<int> expectedFusedLayers;
if (backendId == DNN_BACKEND_OPENCV)
{
if (targetId == DNN_TARGET_CPU)
expectedFusedLayers.push_back(activId); // activation fused with convolution
else if (targetId == DNN_TARGET_OPENCL || targetId == DNN_TARGET_OPENCL_FP16)
{
if (actType == "ReLU" || actType == "ChannelsPReLU" || actType == "ReLU6" || actType == "TanH" /*|| actType == "Power"*/)
expectedFusedLayers.push_back(activId); // activation fused with convolution
}
}
else if(backendId == DNN_BACKEND_CUDA)
{
if (actType == "ReLU" || actType == "ReLU6" || actType == "TanH" || actType == "Swish" ||
actType == "Mish" || actType == "Sigmoid" || actType == "Power")
{
expectedFusedLayers.push_back(activId);
if (eltwiseOp == "sum" && !weightedEltwise)
expectedFusedLayers.push_back(eltwiseId);
}
}
TestLayerFusion::test(input, net, backendId, targetId, expectedFusedLayers);
}
INSTANTIATE_TEST_CASE_P(TestLayerFusion, ConvolutionActivationEltwiseFusion, Combine(
/* bias */ testing::Bool(),
/* activation */ TestLayerFusion::activationLayersList(),
/* eltwise op */ TestLayerFusion::eltwiseOpList(),
/* eltwise weighted */ testing::Bool(),
TestLayerFusion::dnnBackendsAndTargetsForFusionTests()
));
}} // namespace