mirror of
https://github.com/opencv/opencv.git
synced 2025-01-19 15:04:01 +08:00
speedup compilation of row_filter.cu and column_filter.cu
split them into several small files
This commit is contained in:
Vladislav Vinogradov
parent
29f89e8930
commit
810829f32e
53
modules/gpu/src/cuda/column_filter.0.cu
Normal file
53
modules/gpu/src/cuda/column_filter.0.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "column_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearColumn<float, uchar>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/column_filter.1.cu
Normal file
53
modules/gpu/src/cuda/column_filter.1.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "column_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearColumn<float3, uchar3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/column_filter.2.cu
Normal file
53
modules/gpu/src/cuda/column_filter.2.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "column_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearColumn<float4, uchar4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/column_filter.3.cu
Normal file
53
modules/gpu/src/cuda/column_filter.3.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "column_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearColumn<float3, short3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/column_filter.4.cu
Normal file
53
modules/gpu/src/cuda/column_filter.4.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "column_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearColumn<float, int>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/column_filter.5.cu
Normal file
53
modules/gpu/src/cuda/column_filter.5.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "column_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearColumn<float, float>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/column_filter.6.cu
Normal file
53
modules/gpu/src/cuda/column_filter.6.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "column_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearColumn<float3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/column_filter.7.cu
Normal file
53
modules/gpu/src/cuda/column_filter.7.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "column_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearColumn<float4, float4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
@ -1,391 +0,0 @@
|
|||||||
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
|
||||||
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
|
||||||
// Copyright (C) 1993-2011, NVIDIA 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*/
|
|
||||||
|
|
||||||
#if !defined CUDA_DISABLER
|
|
||||||
|
|
||||||
#include "internal_shared.hpp"
|
|
||||||
#include "opencv2/gpu/device/saturate_cast.hpp"
|
|
||||||
#include "opencv2/gpu/device/vec_math.hpp"
|
|
||||||
#include "opencv2/gpu/device/limits.hpp"
|
|
||||||
#include "opencv2/gpu/device/border_interpolate.hpp"
|
|
||||||
#include "opencv2/gpu/device/static_check.hpp"
|
|
||||||
|
|
||||||
namespace cv { namespace gpu { namespace device
|
|
||||||
{
|
|
||||||
namespace column_filter
|
|
||||||
{
|
|
||||||
#define MAX_KERNEL_SIZE 32
|
|
||||||
|
|
||||||
__constant__ float c_kernel[MAX_KERNEL_SIZE];
|
|
||||||
|
|
||||||
void loadKernel(const float* kernel, int ksize, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
if (stream == 0)
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) );
|
|
||||||
else
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbolAsync(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );
|
|
||||||
}
|
|
||||||
|
|
||||||
template <int KSIZE, typename T, typename D, typename B>
|
|
||||||
__global__ void linearColumnFilter(const PtrStepSz<T> src, PtrStep<D> dst, const int anchor, const B brd)
|
|
||||||
{
|
|
||||||
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 200)
|
|
||||||
const int BLOCK_DIM_X = 16;
|
|
||||||
const int BLOCK_DIM_Y = 16;
|
|
||||||
const int PATCH_PER_BLOCK = 4;
|
|
||||||
const int HALO_SIZE = KSIZE <= 16 ? 1 : 2;
|
|
||||||
#else
|
|
||||||
const int BLOCK_DIM_X = 16;
|
|
||||||
const int BLOCK_DIM_Y = 8;
|
|
||||||
const int PATCH_PER_BLOCK = 2;
|
|
||||||
const int HALO_SIZE = 2;
|
|
||||||
#endif
|
|
||||||
|
|
||||||
typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type sum_t;
|
|
||||||
|
|
||||||
__shared__ sum_t smem[(PATCH_PER_BLOCK + 2 * HALO_SIZE) * BLOCK_DIM_Y][BLOCK_DIM_X];
|
|
||||||
|
|
||||||
const int x = blockIdx.x * BLOCK_DIM_X + threadIdx.x;
|
|
||||||
|
|
||||||
if (x >= src.cols)
|
|
||||||
return;
|
|
||||||
|
|
||||||
const T* src_col = src.ptr() + x;
|
|
||||||
|
|
||||||
const int yStart = blockIdx.y * (BLOCK_DIM_Y * PATCH_PER_BLOCK) + threadIdx.y;
|
|
||||||
|
|
||||||
if (blockIdx.y > 0)
|
|
||||||
{
|
|
||||||
//Upper halo
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < HALO_SIZE; ++j)
|
|
||||||
smem[threadIdx.y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart - (HALO_SIZE - j) * BLOCK_DIM_Y, x));
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
//Upper halo
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < HALO_SIZE; ++j)
|
|
||||||
smem[threadIdx.y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_low(yStart - (HALO_SIZE - j) * BLOCK_DIM_Y, src_col, src.step));
|
|
||||||
}
|
|
||||||
|
|
||||||
if (blockIdx.y + 2 < gridDim.y)
|
|
||||||
{
|
|
||||||
//Main data
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
|
||||||
smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart + j * BLOCK_DIM_Y, x));
|
|
||||||
|
|
||||||
//Lower halo
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < HALO_SIZE; ++j)
|
|
||||||
smem[threadIdx.y + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_Y, x));
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
//Main data
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
|
||||||
smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_high(yStart + j * BLOCK_DIM_Y, src_col, src.step));
|
|
||||||
|
|
||||||
//Lower halo
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < HALO_SIZE; ++j)
|
|
||||||
smem[threadIdx.y + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_high(yStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_Y, src_col, src.step));
|
|
||||||
}
|
|
||||||
|
|
||||||
__syncthreads();
|
|
||||||
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
|
||||||
{
|
|
||||||
const int y = yStart + j * BLOCK_DIM_Y;
|
|
||||||
|
|
||||||
if (y < src.rows)
|
|
||||||
{
|
|
||||||
sum_t sum = VecTraits<sum_t>::all(0);
|
|
||||||
|
|
||||||
#pragma unroll
|
|
||||||
for (int k = 0; k < KSIZE; ++k)
|
|
||||||
sum = sum + smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y - anchor + k][threadIdx.x] * c_kernel[k];
|
|
||||||
|
|
||||||
dst(y, x) = saturate_cast<D>(sum);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
template <int KSIZE, typename T, typename D, template<typename> class B>
|
|
||||||
void linearColumnFilter_caller(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
int BLOCK_DIM_X;
|
|
||||||
int BLOCK_DIM_Y;
|
|
||||||
int PATCH_PER_BLOCK;
|
|
||||||
|
|
||||||
if (cc >= 20)
|
|
||||||
{
|
|
||||||
BLOCK_DIM_X = 16;
|
|
||||||
BLOCK_DIM_Y = 16;
|
|
||||||
PATCH_PER_BLOCK = 4;
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
BLOCK_DIM_X = 16;
|
|
||||||
BLOCK_DIM_Y = 8;
|
|
||||||
PATCH_PER_BLOCK = 2;
|
|
||||||
}
|
|
||||||
|
|
||||||
const dim3 block(BLOCK_DIM_X, BLOCK_DIM_Y);
|
|
||||||
const dim3 grid(divUp(src.cols, BLOCK_DIM_X), divUp(src.rows, BLOCK_DIM_Y * PATCH_PER_BLOCK));
|
|
||||||
|
|
||||||
B<T> brd(src.rows);
|
|
||||||
|
|
||||||
linearColumnFilter<KSIZE, T, D><<<grid, block, 0, stream>>>(src, dst, anchor, brd);
|
|
||||||
|
|
||||||
cudaSafeCall( cudaGetLastError() );
|
|
||||||
|
|
||||||
if (stream == 0)
|
|
||||||
cudaSafeCall( cudaDeviceSynchronize() );
|
|
||||||
}
|
|
||||||
|
|
||||||
template <typename T, typename D>
|
|
||||||
void linearColumnFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
typedef void (*caller_t)(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream);
|
|
||||||
|
|
||||||
static const caller_t callers[5][33] =
|
|
||||||
{
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearColumnFilter_caller< 1, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller< 2, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller< 3, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller< 4, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller< 5, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller< 6, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller< 7, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller< 8, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller< 9, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<10, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<11, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<12, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<13, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<14, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<15, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<16, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<17, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<18, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<19, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<20, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<21, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<22, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<23, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<24, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<25, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<26, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<27, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<28, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<29, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<30, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<31, T, D, BrdColReflect101>,
|
|
||||||
linearColumnFilter_caller<32, T, D, BrdColReflect101>
|
|
||||||
},
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearColumnFilter_caller< 1, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller< 2, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller< 3, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller< 4, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller< 5, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller< 6, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller< 7, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller< 8, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller< 9, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<10, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<11, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<12, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<13, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<14, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<15, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<16, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<17, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<18, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<19, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<20, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<21, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<22, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<23, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<24, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<25, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<26, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<27, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<28, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<29, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<30, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<31, T, D, BrdColReplicate>,
|
|
||||||
linearColumnFilter_caller<32, T, D, BrdColReplicate>
|
|
||||||
},
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearColumnFilter_caller< 1, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller< 2, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller< 3, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller< 4, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller< 5, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller< 6, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller< 7, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller< 8, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller< 9, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<10, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<11, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<12, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<13, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<14, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<15, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<16, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<17, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<18, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<19, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<20, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<21, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<22, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<23, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<24, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<25, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<26, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<27, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<28, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<29, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<30, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<31, T, D, BrdColConstant>,
|
|
||||||
linearColumnFilter_caller<32, T, D, BrdColConstant>
|
|
||||||
},
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearColumnFilter_caller< 1, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller< 2, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller< 3, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller< 4, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller< 5, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller< 6, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller< 7, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller< 8, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller< 9, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<10, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<11, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<12, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<13, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<14, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<15, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<16, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<17, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<18, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<19, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<20, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<21, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<22, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<23, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<24, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<25, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<26, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<27, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<28, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<29, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<30, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<31, T, D, BrdColReflect>,
|
|
||||||
linearColumnFilter_caller<32, T, D, BrdColReflect>
|
|
||||||
},
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearColumnFilter_caller< 1, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller< 2, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller< 3, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller< 4, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller< 5, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller< 6, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller< 7, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller< 8, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller< 9, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<10, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<11, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<12, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<13, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<14, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<15, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<16, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<17, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<18, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<19, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<20, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<21, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<22, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<23, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<24, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<25, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<26, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<27, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<28, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<29, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<30, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<31, T, D, BrdColWrap>,
|
|
||||||
linearColumnFilter_caller<32, T, D, BrdColWrap>
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
loadKernel(kernel, ksize, stream);
|
|
||||||
|
|
||||||
callers[brd_type][ksize]((PtrStepSz<T>)src, (PtrStepSz<D>)dst, anchor, cc, stream);
|
|
||||||
}
|
|
||||||
|
|
||||||
template void linearColumnFilter_gpu<float , uchar >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearColumnFilter_gpu<float3, uchar3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearColumnFilter_gpu<float4, uchar4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearColumnFilter_gpu<float3, short3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearColumnFilter_gpu<float , int >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearColumnFilter_gpu<float , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearColumnFilter_gpu<float3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearColumnFilter_gpu<float4, float4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
} // namespace column_filter
|
|
||||||
}}} // namespace cv { namespace gpu { namespace device
|
|
||||||
|
|
||||||
|
|
||||||
#endif /* CUDA_DISABLER */
|
|
||||||
378
modules/gpu/src/cuda/column_filter.h
Normal file
378
modules/gpu/src/cuda/column_filter.h
Normal file
@ -0,0 +1,378 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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 "opencv2/gpu/device/common.hpp"
|
||||||
|
#include "opencv2/gpu/device/saturate_cast.hpp"
|
||||||
|
#include "opencv2/gpu/device/vec_math.hpp"
|
||||||
|
#include "opencv2/gpu/device/border_interpolate.hpp"
|
||||||
|
|
||||||
|
using namespace cv::gpu;
|
||||||
|
using namespace cv::gpu::device;
|
||||||
|
|
||||||
|
namespace
|
||||||
|
{
|
||||||
|
#define MAX_KERNEL_SIZE 32
|
||||||
|
|
||||||
|
__constant__ float c_kernel[MAX_KERNEL_SIZE];
|
||||||
|
|
||||||
|
void loadKernel(const float* kernel, int ksize, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
if (stream == 0)
|
||||||
|
cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) );
|
||||||
|
else
|
||||||
|
cudaSafeCall( cudaMemcpyToSymbolAsync(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );
|
||||||
|
}
|
||||||
|
|
||||||
|
template <int KSIZE, typename T, typename D, typename B>
|
||||||
|
__global__ void linearColumnFilter(const PtrStepSz<T> src, PtrStep<D> dst, const int anchor, const B brd)
|
||||||
|
{
|
||||||
|
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 200)
|
||||||
|
const int BLOCK_DIM_X = 16;
|
||||||
|
const int BLOCK_DIM_Y = 16;
|
||||||
|
const int PATCH_PER_BLOCK = 4;
|
||||||
|
const int HALO_SIZE = KSIZE <= 16 ? 1 : 2;
|
||||||
|
#else
|
||||||
|
const int BLOCK_DIM_X = 16;
|
||||||
|
const int BLOCK_DIM_Y = 8;
|
||||||
|
const int PATCH_PER_BLOCK = 2;
|
||||||
|
const int HALO_SIZE = 2;
|
||||||
|
#endif
|
||||||
|
|
||||||
|
typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type sum_t;
|
||||||
|
|
||||||
|
__shared__ sum_t smem[(PATCH_PER_BLOCK + 2 * HALO_SIZE) * BLOCK_DIM_Y][BLOCK_DIM_X];
|
||||||
|
|
||||||
|
const int x = blockIdx.x * BLOCK_DIM_X + threadIdx.x;
|
||||||
|
|
||||||
|
if (x >= src.cols)
|
||||||
|
return;
|
||||||
|
|
||||||
|
const T* src_col = src.ptr() + x;
|
||||||
|
|
||||||
|
const int yStart = blockIdx.y * (BLOCK_DIM_Y * PATCH_PER_BLOCK) + threadIdx.y;
|
||||||
|
|
||||||
|
if (blockIdx.y > 0)
|
||||||
|
{
|
||||||
|
//Upper halo
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < HALO_SIZE; ++j)
|
||||||
|
smem[threadIdx.y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart - (HALO_SIZE - j) * BLOCK_DIM_Y, x));
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
//Upper halo
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < HALO_SIZE; ++j)
|
||||||
|
smem[threadIdx.y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_low(yStart - (HALO_SIZE - j) * BLOCK_DIM_Y, src_col, src.step));
|
||||||
|
}
|
||||||
|
|
||||||
|
if (blockIdx.y + 2 < gridDim.y)
|
||||||
|
{
|
||||||
|
//Main data
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
||||||
|
smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart + j * BLOCK_DIM_Y, x));
|
||||||
|
|
||||||
|
//Lower halo
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < HALO_SIZE; ++j)
|
||||||
|
smem[threadIdx.y + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_Y, x));
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
//Main data
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
||||||
|
smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_high(yStart + j * BLOCK_DIM_Y, src_col, src.step));
|
||||||
|
|
||||||
|
//Lower halo
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < HALO_SIZE; ++j)
|
||||||
|
smem[threadIdx.y + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_high(yStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_Y, src_col, src.step));
|
||||||
|
}
|
||||||
|
|
||||||
|
__syncthreads();
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
||||||
|
{
|
||||||
|
const int y = yStart + j * BLOCK_DIM_Y;
|
||||||
|
|
||||||
|
if (y < src.rows)
|
||||||
|
{
|
||||||
|
sum_t sum = VecTraits<sum_t>::all(0);
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int k = 0; k < KSIZE; ++k)
|
||||||
|
sum = sum + smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y - anchor + k][threadIdx.x] * c_kernel[k];
|
||||||
|
|
||||||
|
dst(y, x) = saturate_cast<D>(sum);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
template <int KSIZE, typename T, typename D, template<typename> class B>
|
||||||
|
void caller(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
int BLOCK_DIM_X;
|
||||||
|
int BLOCK_DIM_Y;
|
||||||
|
int PATCH_PER_BLOCK;
|
||||||
|
|
||||||
|
if (cc >= 20)
|
||||||
|
{
|
||||||
|
BLOCK_DIM_X = 16;
|
||||||
|
BLOCK_DIM_Y = 16;
|
||||||
|
PATCH_PER_BLOCK = 4;
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
BLOCK_DIM_X = 16;
|
||||||
|
BLOCK_DIM_Y = 8;
|
||||||
|
PATCH_PER_BLOCK = 2;
|
||||||
|
}
|
||||||
|
|
||||||
|
const dim3 block(BLOCK_DIM_X, BLOCK_DIM_Y);
|
||||||
|
const dim3 grid(divUp(src.cols, BLOCK_DIM_X), divUp(src.rows, BLOCK_DIM_Y * PATCH_PER_BLOCK));
|
||||||
|
|
||||||
|
B<T> brd(src.rows);
|
||||||
|
|
||||||
|
linearColumnFilter<KSIZE, T, D><<<grid, block, 0, stream>>>(src, dst, anchor, brd);
|
||||||
|
|
||||||
|
cudaSafeCall( cudaGetLastError() );
|
||||||
|
|
||||||
|
if (stream == 0)
|
||||||
|
cudaSafeCall( cudaDeviceSynchronize() );
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template <typename T, typename D>
|
||||||
|
void linearColumn(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
typedef void (*caller_t)(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream);
|
||||||
|
|
||||||
|
static const caller_t callers[5][33] =
|
||||||
|
{
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdColReflect101>,
|
||||||
|
::caller< 2, T, D, BrdColReflect101>,
|
||||||
|
::caller< 3, T, D, BrdColReflect101>,
|
||||||
|
::caller< 4, T, D, BrdColReflect101>,
|
||||||
|
::caller< 5, T, D, BrdColReflect101>,
|
||||||
|
::caller< 6, T, D, BrdColReflect101>,
|
||||||
|
::caller< 7, T, D, BrdColReflect101>,
|
||||||
|
::caller< 8, T, D, BrdColReflect101>,
|
||||||
|
::caller< 9, T, D, BrdColReflect101>,
|
||||||
|
::caller<10, T, D, BrdColReflect101>,
|
||||||
|
::caller<11, T, D, BrdColReflect101>,
|
||||||
|
::caller<12, T, D, BrdColReflect101>,
|
||||||
|
::caller<13, T, D, BrdColReflect101>,
|
||||||
|
::caller<14, T, D, BrdColReflect101>,
|
||||||
|
::caller<15, T, D, BrdColReflect101>,
|
||||||
|
::caller<16, T, D, BrdColReflect101>,
|
||||||
|
::caller<17, T, D, BrdColReflect101>,
|
||||||
|
::caller<18, T, D, BrdColReflect101>,
|
||||||
|
::caller<19, T, D, BrdColReflect101>,
|
||||||
|
::caller<20, T, D, BrdColReflect101>,
|
||||||
|
::caller<21, T, D, BrdColReflect101>,
|
||||||
|
::caller<22, T, D, BrdColReflect101>,
|
||||||
|
::caller<23, T, D, BrdColReflect101>,
|
||||||
|
::caller<24, T, D, BrdColReflect101>,
|
||||||
|
::caller<25, T, D, BrdColReflect101>,
|
||||||
|
::caller<26, T, D, BrdColReflect101>,
|
||||||
|
::caller<27, T, D, BrdColReflect101>,
|
||||||
|
::caller<28, T, D, BrdColReflect101>,
|
||||||
|
::caller<29, T, D, BrdColReflect101>,
|
||||||
|
::caller<30, T, D, BrdColReflect101>,
|
||||||
|
::caller<31, T, D, BrdColReflect101>,
|
||||||
|
::caller<32, T, D, BrdColReflect101>
|
||||||
|
},
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdColReplicate>,
|
||||||
|
::caller< 2, T, D, BrdColReplicate>,
|
||||||
|
::caller< 3, T, D, BrdColReplicate>,
|
||||||
|
::caller< 4, T, D, BrdColReplicate>,
|
||||||
|
::caller< 5, T, D, BrdColReplicate>,
|
||||||
|
::caller< 6, T, D, BrdColReplicate>,
|
||||||
|
::caller< 7, T, D, BrdColReplicate>,
|
||||||
|
::caller< 8, T, D, BrdColReplicate>,
|
||||||
|
::caller< 9, T, D, BrdColReplicate>,
|
||||||
|
::caller<10, T, D, BrdColReplicate>,
|
||||||
|
::caller<11, T, D, BrdColReplicate>,
|
||||||
|
::caller<12, T, D, BrdColReplicate>,
|
||||||
|
::caller<13, T, D, BrdColReplicate>,
|
||||||
|
::caller<14, T, D, BrdColReplicate>,
|
||||||
|
::caller<15, T, D, BrdColReplicate>,
|
||||||
|
::caller<16, T, D, BrdColReplicate>,
|
||||||
|
::caller<17, T, D, BrdColReplicate>,
|
||||||
|
::caller<18, T, D, BrdColReplicate>,
|
||||||
|
::caller<19, T, D, BrdColReplicate>,
|
||||||
|
::caller<20, T, D, BrdColReplicate>,
|
||||||
|
::caller<21, T, D, BrdColReplicate>,
|
||||||
|
::caller<22, T, D, BrdColReplicate>,
|
||||||
|
::caller<23, T, D, BrdColReplicate>,
|
||||||
|
::caller<24, T, D, BrdColReplicate>,
|
||||||
|
::caller<25, T, D, BrdColReplicate>,
|
||||||
|
::caller<26, T, D, BrdColReplicate>,
|
||||||
|
::caller<27, T, D, BrdColReplicate>,
|
||||||
|
::caller<28, T, D, BrdColReplicate>,
|
||||||
|
::caller<29, T, D, BrdColReplicate>,
|
||||||
|
::caller<30, T, D, BrdColReplicate>,
|
||||||
|
::caller<31, T, D, BrdColReplicate>,
|
||||||
|
::caller<32, T, D, BrdColReplicate>
|
||||||
|
},
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdColConstant>,
|
||||||
|
::caller< 2, T, D, BrdColConstant>,
|
||||||
|
::caller< 3, T, D, BrdColConstant>,
|
||||||
|
::caller< 4, T, D, BrdColConstant>,
|
||||||
|
::caller< 5, T, D, BrdColConstant>,
|
||||||
|
::caller< 6, T, D, BrdColConstant>,
|
||||||
|
::caller< 7, T, D, BrdColConstant>,
|
||||||
|
::caller< 8, T, D, BrdColConstant>,
|
||||||
|
::caller< 9, T, D, BrdColConstant>,
|
||||||
|
::caller<10, T, D, BrdColConstant>,
|
||||||
|
::caller<11, T, D, BrdColConstant>,
|
||||||
|
::caller<12, T, D, BrdColConstant>,
|
||||||
|
::caller<13, T, D, BrdColConstant>,
|
||||||
|
::caller<14, T, D, BrdColConstant>,
|
||||||
|
::caller<15, T, D, BrdColConstant>,
|
||||||
|
::caller<16, T, D, BrdColConstant>,
|
||||||
|
::caller<17, T, D, BrdColConstant>,
|
||||||
|
::caller<18, T, D, BrdColConstant>,
|
||||||
|
::caller<19, T, D, BrdColConstant>,
|
||||||
|
::caller<20, T, D, BrdColConstant>,
|
||||||
|
::caller<21, T, D, BrdColConstant>,
|
||||||
|
::caller<22, T, D, BrdColConstant>,
|
||||||
|
::caller<23, T, D, BrdColConstant>,
|
||||||
|
::caller<24, T, D, BrdColConstant>,
|
||||||
|
::caller<25, T, D, BrdColConstant>,
|
||||||
|
::caller<26, T, D, BrdColConstant>,
|
||||||
|
::caller<27, T, D, BrdColConstant>,
|
||||||
|
::caller<28, T, D, BrdColConstant>,
|
||||||
|
::caller<29, T, D, BrdColConstant>,
|
||||||
|
::caller<30, T, D, BrdColConstant>,
|
||||||
|
::caller<31, T, D, BrdColConstant>,
|
||||||
|
::caller<32, T, D, BrdColConstant>
|
||||||
|
},
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdColReflect>,
|
||||||
|
::caller< 2, T, D, BrdColReflect>,
|
||||||
|
::caller< 3, T, D, BrdColReflect>,
|
||||||
|
::caller< 4, T, D, BrdColReflect>,
|
||||||
|
::caller< 5, T, D, BrdColReflect>,
|
||||||
|
::caller< 6, T, D, BrdColReflect>,
|
||||||
|
::caller< 7, T, D, BrdColReflect>,
|
||||||
|
::caller< 8, T, D, BrdColReflect>,
|
||||||
|
::caller< 9, T, D, BrdColReflect>,
|
||||||
|
::caller<10, T, D, BrdColReflect>,
|
||||||
|
::caller<11, T, D, BrdColReflect>,
|
||||||
|
::caller<12, T, D, BrdColReflect>,
|
||||||
|
::caller<13, T, D, BrdColReflect>,
|
||||||
|
::caller<14, T, D, BrdColReflect>,
|
||||||
|
::caller<15, T, D, BrdColReflect>,
|
||||||
|
::caller<16, T, D, BrdColReflect>,
|
||||||
|
::caller<17, T, D, BrdColReflect>,
|
||||||
|
::caller<18, T, D, BrdColReflect>,
|
||||||
|
::caller<19, T, D, BrdColReflect>,
|
||||||
|
::caller<20, T, D, BrdColReflect>,
|
||||||
|
::caller<21, T, D, BrdColReflect>,
|
||||||
|
::caller<22, T, D, BrdColReflect>,
|
||||||
|
::caller<23, T, D, BrdColReflect>,
|
||||||
|
::caller<24, T, D, BrdColReflect>,
|
||||||
|
::caller<25, T, D, BrdColReflect>,
|
||||||
|
::caller<26, T, D, BrdColReflect>,
|
||||||
|
::caller<27, T, D, BrdColReflect>,
|
||||||
|
::caller<28, T, D, BrdColReflect>,
|
||||||
|
::caller<29, T, D, BrdColReflect>,
|
||||||
|
::caller<30, T, D, BrdColReflect>,
|
||||||
|
::caller<31, T, D, BrdColReflect>,
|
||||||
|
::caller<32, T, D, BrdColReflect>
|
||||||
|
},
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdColWrap>,
|
||||||
|
::caller< 2, T, D, BrdColWrap>,
|
||||||
|
::caller< 3, T, D, BrdColWrap>,
|
||||||
|
::caller< 4, T, D, BrdColWrap>,
|
||||||
|
::caller< 5, T, D, BrdColWrap>,
|
||||||
|
::caller< 6, T, D, BrdColWrap>,
|
||||||
|
::caller< 7, T, D, BrdColWrap>,
|
||||||
|
::caller< 8, T, D, BrdColWrap>,
|
||||||
|
::caller< 9, T, D, BrdColWrap>,
|
||||||
|
::caller<10, T, D, BrdColWrap>,
|
||||||
|
::caller<11, T, D, BrdColWrap>,
|
||||||
|
::caller<12, T, D, BrdColWrap>,
|
||||||
|
::caller<13, T, D, BrdColWrap>,
|
||||||
|
::caller<14, T, D, BrdColWrap>,
|
||||||
|
::caller<15, T, D, BrdColWrap>,
|
||||||
|
::caller<16, T, D, BrdColWrap>,
|
||||||
|
::caller<17, T, D, BrdColWrap>,
|
||||||
|
::caller<18, T, D, BrdColWrap>,
|
||||||
|
::caller<19, T, D, BrdColWrap>,
|
||||||
|
::caller<20, T, D, BrdColWrap>,
|
||||||
|
::caller<21, T, D, BrdColWrap>,
|
||||||
|
::caller<22, T, D, BrdColWrap>,
|
||||||
|
::caller<23, T, D, BrdColWrap>,
|
||||||
|
::caller<24, T, D, BrdColWrap>,
|
||||||
|
::caller<25, T, D, BrdColWrap>,
|
||||||
|
::caller<26, T, D, BrdColWrap>,
|
||||||
|
::caller<27, T, D, BrdColWrap>,
|
||||||
|
::caller<28, T, D, BrdColWrap>,
|
||||||
|
::caller<29, T, D, BrdColWrap>,
|
||||||
|
::caller<30, T, D, BrdColWrap>,
|
||||||
|
::caller<31, T, D, BrdColWrap>,
|
||||||
|
::caller<32, T, D, BrdColWrap>
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
::loadKernel(kernel, ksize, stream);
|
||||||
|
|
||||||
|
callers[brd_type][ksize]((PtrStepSz<T>)src, (PtrStepSz<D>)dst, anchor, cc, stream);
|
||||||
|
}
|
||||||
|
}
|
||||||
53
modules/gpu/src/cuda/row_filter.0.cu
Normal file
53
modules/gpu/src/cuda/row_filter.0.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "row_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearRow<uchar, float>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/row_filter.1.cu
Normal file
53
modules/gpu/src/cuda/row_filter.1.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "row_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearRow<uchar3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/row_filter.2.cu
Normal file
53
modules/gpu/src/cuda/row_filter.2.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "row_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearRow<uchar4, float4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/row_filter.3.cu
Normal file
53
modules/gpu/src/cuda/row_filter.3.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "row_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearRow<short3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/row_filter.4.cu
Normal file
53
modules/gpu/src/cuda/row_filter.4.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "row_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearRow<int, float>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/row_filter.5.cu
Normal file
53
modules/gpu/src/cuda/row_filter.5.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "row_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearRow<float, float>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/row_filter.6.cu
Normal file
53
modules/gpu/src/cuda/row_filter.6.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "row_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearRow<float3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
53
modules/gpu/src/cuda/row_filter.7.cu
Normal file
53
modules/gpu/src/cuda/row_filter.7.cu
Normal file
@ -0,0 +1,53 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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*/
|
||||||
|
|
||||||
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
|
#include "row_filter.h"
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template void linearRow<float4, float4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
#endif /* CUDA_DISABLER */
|
||||||
@ -1,390 +0,0 @@
|
|||||||
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
|
||||||
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
|
||||||
// Copyright (C) 1993-2011, NVIDIA 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*/
|
|
||||||
|
|
||||||
#if !defined CUDA_DISABLER
|
|
||||||
|
|
||||||
#include "internal_shared.hpp"
|
|
||||||
#include "opencv2/gpu/device/saturate_cast.hpp"
|
|
||||||
#include "opencv2/gpu/device/vec_math.hpp"
|
|
||||||
#include "opencv2/gpu/device/limits.hpp"
|
|
||||||
#include "opencv2/gpu/device/border_interpolate.hpp"
|
|
||||||
#include "opencv2/gpu/device/static_check.hpp"
|
|
||||||
|
|
||||||
namespace cv { namespace gpu { namespace device
|
|
||||||
{
|
|
||||||
namespace row_filter
|
|
||||||
{
|
|
||||||
#define MAX_KERNEL_SIZE 32
|
|
||||||
|
|
||||||
__constant__ float c_kernel[MAX_KERNEL_SIZE];
|
|
||||||
|
|
||||||
void loadKernel(const float* kernel, int ksize, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
if (stream == 0)
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) );
|
|
||||||
else
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbolAsync(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );
|
|
||||||
}
|
|
||||||
|
|
||||||
template <int KSIZE, typename T, typename D, typename B>
|
|
||||||
__global__ void linearRowFilter(const PtrStepSz<T> src, PtrStep<D> dst, const int anchor, const B brd)
|
|
||||||
{
|
|
||||||
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 200)
|
|
||||||
const int BLOCK_DIM_X = 32;
|
|
||||||
const int BLOCK_DIM_Y = 8;
|
|
||||||
const int PATCH_PER_BLOCK = 4;
|
|
||||||
const int HALO_SIZE = 1;
|
|
||||||
#else
|
|
||||||
const int BLOCK_DIM_X = 32;
|
|
||||||
const int BLOCK_DIM_Y = 4;
|
|
||||||
const int PATCH_PER_BLOCK = 4;
|
|
||||||
const int HALO_SIZE = 1;
|
|
||||||
#endif
|
|
||||||
|
|
||||||
typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type sum_t;
|
|
||||||
|
|
||||||
__shared__ sum_t smem[BLOCK_DIM_Y][(PATCH_PER_BLOCK + 2 * HALO_SIZE) * BLOCK_DIM_X];
|
|
||||||
|
|
||||||
const int y = blockIdx.y * BLOCK_DIM_Y + threadIdx.y;
|
|
||||||
|
|
||||||
if (y >= src.rows)
|
|
||||||
return;
|
|
||||||
|
|
||||||
const T* src_row = src.ptr(y);
|
|
||||||
|
|
||||||
const int xStart = blockIdx.x * (PATCH_PER_BLOCK * BLOCK_DIM_X) + threadIdx.x;
|
|
||||||
|
|
||||||
if (blockIdx.x > 0)
|
|
||||||
{
|
|
||||||
//Load left halo
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < HALO_SIZE; ++j)
|
|
||||||
smem[threadIdx.y][threadIdx.x + j * BLOCK_DIM_X] = saturate_cast<sum_t>(src_row[xStart - (HALO_SIZE - j) * BLOCK_DIM_X]);
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
//Load left halo
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < HALO_SIZE; ++j)
|
|
||||||
smem[threadIdx.y][threadIdx.x + j * BLOCK_DIM_X] = saturate_cast<sum_t>(brd.at_low(xStart - (HALO_SIZE - j) * BLOCK_DIM_X, src_row));
|
|
||||||
}
|
|
||||||
|
|
||||||
if (blockIdx.x + 2 < gridDim.x)
|
|
||||||
{
|
|
||||||
//Load main data
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
|
||||||
smem[threadIdx.y][threadIdx.x + HALO_SIZE * BLOCK_DIM_X + j * BLOCK_DIM_X] = saturate_cast<sum_t>(src_row[xStart + j * BLOCK_DIM_X]);
|
|
||||||
|
|
||||||
//Load right halo
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < HALO_SIZE; ++j)
|
|
||||||
smem[threadIdx.y][threadIdx.x + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_X + j * BLOCK_DIM_X] = saturate_cast<sum_t>(src_row[xStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_X]);
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
//Load main data
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
|
||||||
smem[threadIdx.y][threadIdx.x + HALO_SIZE * BLOCK_DIM_X + j * BLOCK_DIM_X] = saturate_cast<sum_t>(brd.at_high(xStart + j * BLOCK_DIM_X, src_row));
|
|
||||||
|
|
||||||
//Load right halo
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < HALO_SIZE; ++j)
|
|
||||||
smem[threadIdx.y][threadIdx.x + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_X + j * BLOCK_DIM_X] = saturate_cast<sum_t>(brd.at_high(xStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_X, src_row));
|
|
||||||
}
|
|
||||||
|
|
||||||
__syncthreads();
|
|
||||||
|
|
||||||
#pragma unroll
|
|
||||||
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
|
||||||
{
|
|
||||||
const int x = xStart + j * BLOCK_DIM_X;
|
|
||||||
|
|
||||||
if (x < src.cols)
|
|
||||||
{
|
|
||||||
sum_t sum = VecTraits<sum_t>::all(0);
|
|
||||||
|
|
||||||
#pragma unroll
|
|
||||||
for (int k = 0; k < KSIZE; ++k)
|
|
||||||
sum = sum + smem[threadIdx.y][threadIdx.x + HALO_SIZE * BLOCK_DIM_X + j * BLOCK_DIM_X - anchor + k] * c_kernel[k];
|
|
||||||
|
|
||||||
dst(y, x) = saturate_cast<D>(sum);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
template <int KSIZE, typename T, typename D, template<typename> class B>
|
|
||||||
void linearRowFilter_caller(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
int BLOCK_DIM_X;
|
|
||||||
int BLOCK_DIM_Y;
|
|
||||||
int PATCH_PER_BLOCK;
|
|
||||||
|
|
||||||
if (cc >= 20)
|
|
||||||
{
|
|
||||||
BLOCK_DIM_X = 32;
|
|
||||||
BLOCK_DIM_Y = 8;
|
|
||||||
PATCH_PER_BLOCK = 4;
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
BLOCK_DIM_X = 32;
|
|
||||||
BLOCK_DIM_Y = 4;
|
|
||||||
PATCH_PER_BLOCK = 4;
|
|
||||||
}
|
|
||||||
|
|
||||||
const dim3 block(BLOCK_DIM_X, BLOCK_DIM_Y);
|
|
||||||
const dim3 grid(divUp(src.cols, BLOCK_DIM_X * PATCH_PER_BLOCK), divUp(src.rows, BLOCK_DIM_Y));
|
|
||||||
|
|
||||||
B<T> brd(src.cols);
|
|
||||||
|
|
||||||
linearRowFilter<KSIZE, T, D><<<grid, block, 0, stream>>>(src, dst, anchor, brd);
|
|
||||||
cudaSafeCall( cudaGetLastError() );
|
|
||||||
|
|
||||||
if (stream == 0)
|
|
||||||
cudaSafeCall( cudaDeviceSynchronize() );
|
|
||||||
}
|
|
||||||
|
|
||||||
template <typename T, typename D>
|
|
||||||
void linearRowFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
typedef void (*caller_t)(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream);
|
|
||||||
|
|
||||||
static const caller_t callers[5][33] =
|
|
||||||
{
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearRowFilter_caller< 1, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller< 2, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller< 3, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller< 4, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller< 5, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller< 6, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller< 7, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller< 8, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller< 9, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<10, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<11, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<12, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<13, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<14, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<15, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<16, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<17, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<18, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<19, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<20, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<21, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<22, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<23, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<24, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<25, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<26, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<27, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<28, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<29, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<30, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<31, T, D, BrdRowReflect101>,
|
|
||||||
linearRowFilter_caller<32, T, D, BrdRowReflect101>
|
|
||||||
},
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearRowFilter_caller< 1, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller< 2, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller< 3, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller< 4, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller< 5, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller< 6, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller< 7, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller< 8, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller< 9, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<10, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<11, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<12, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<13, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<14, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<15, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<16, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<17, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<18, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<19, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<20, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<21, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<22, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<23, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<24, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<25, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<26, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<27, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<28, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<29, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<30, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<31, T, D, BrdRowReplicate>,
|
|
||||||
linearRowFilter_caller<32, T, D, BrdRowReplicate>
|
|
||||||
},
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearRowFilter_caller< 1, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller< 2, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller< 3, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller< 4, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller< 5, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller< 6, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller< 7, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller< 8, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller< 9, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<10, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<11, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<12, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<13, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<14, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<15, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<16, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<17, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<18, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<19, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<20, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<21, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<22, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<23, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<24, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<25, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<26, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<27, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<28, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<29, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<30, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<31, T, D, BrdRowConstant>,
|
|
||||||
linearRowFilter_caller<32, T, D, BrdRowConstant>
|
|
||||||
},
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearRowFilter_caller< 1, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller< 2, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller< 3, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller< 4, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller< 5, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller< 6, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller< 7, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller< 8, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller< 9, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<10, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<11, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<12, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<13, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<14, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<15, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<16, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<17, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<18, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<19, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<20, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<21, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<22, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<23, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<24, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<25, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<26, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<27, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<28, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<29, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<30, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<31, T, D, BrdRowReflect>,
|
|
||||||
linearRowFilter_caller<32, T, D, BrdRowReflect>
|
|
||||||
},
|
|
||||||
{
|
|
||||||
0,
|
|
||||||
linearRowFilter_caller< 1, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller< 2, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller< 3, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller< 4, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller< 5, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller< 6, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller< 7, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller< 8, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller< 9, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<10, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<11, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<12, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<13, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<14, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<15, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<16, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<17, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<18, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<19, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<20, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<21, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<22, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<23, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<24, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<25, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<26, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<27, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<28, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<29, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<30, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<31, T, D, BrdRowWrap>,
|
|
||||||
linearRowFilter_caller<32, T, D, BrdRowWrap>
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
loadKernel(kernel, ksize, stream);
|
|
||||||
|
|
||||||
callers[brd_type][ksize]((PtrStepSz<T>)src, (PtrStepSz<D>)dst, anchor, cc, stream);
|
|
||||||
}
|
|
||||||
|
|
||||||
template void linearRowFilter_gpu<uchar , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearRowFilter_gpu<uchar3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearRowFilter_gpu<uchar4, float4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearRowFilter_gpu<short3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearRowFilter_gpu<int , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearRowFilter_gpu<float , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearRowFilter_gpu<float3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
template void linearRowFilter_gpu<float4, float4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
} // namespace row_filter
|
|
||||||
}}} // namespace cv { namespace gpu { namespace device
|
|
||||||
|
|
||||||
|
|
||||||
#endif /* CUDA_DISABLER */
|
|
||||||
377
modules/gpu/src/cuda/row_filter.h
Normal file
377
modules/gpu/src/cuda/row_filter.h
Normal file
@ -0,0 +1,377 @@
|
|||||||
|
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||||
|
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
||||||
|
// Copyright (C) 1993-2011, NVIDIA 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 "opencv2/gpu/device/common.hpp"
|
||||||
|
#include "opencv2/gpu/device/saturate_cast.hpp"
|
||||||
|
#include "opencv2/gpu/device/vec_math.hpp"
|
||||||
|
#include "opencv2/gpu/device/border_interpolate.hpp"
|
||||||
|
|
||||||
|
using namespace cv::gpu;
|
||||||
|
using namespace cv::gpu::device;
|
||||||
|
|
||||||
|
namespace
|
||||||
|
{
|
||||||
|
#define MAX_KERNEL_SIZE 32
|
||||||
|
|
||||||
|
__constant__ float c_kernel[MAX_KERNEL_SIZE];
|
||||||
|
|
||||||
|
void loadKernel(const float* kernel, int ksize, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
if (stream == 0)
|
||||||
|
cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) );
|
||||||
|
else
|
||||||
|
cudaSafeCall( cudaMemcpyToSymbolAsync(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );
|
||||||
|
}
|
||||||
|
|
||||||
|
template <int KSIZE, typename T, typename D, typename B>
|
||||||
|
__global__ void linearRowFilter(const PtrStepSz<T> src, PtrStep<D> dst, const int anchor, const B brd)
|
||||||
|
{
|
||||||
|
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 200)
|
||||||
|
const int BLOCK_DIM_X = 32;
|
||||||
|
const int BLOCK_DIM_Y = 8;
|
||||||
|
const int PATCH_PER_BLOCK = 4;
|
||||||
|
const int HALO_SIZE = 1;
|
||||||
|
#else
|
||||||
|
const int BLOCK_DIM_X = 32;
|
||||||
|
const int BLOCK_DIM_Y = 4;
|
||||||
|
const int PATCH_PER_BLOCK = 4;
|
||||||
|
const int HALO_SIZE = 1;
|
||||||
|
#endif
|
||||||
|
|
||||||
|
typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type sum_t;
|
||||||
|
|
||||||
|
__shared__ sum_t smem[BLOCK_DIM_Y][(PATCH_PER_BLOCK + 2 * HALO_SIZE) * BLOCK_DIM_X];
|
||||||
|
|
||||||
|
const int y = blockIdx.y * BLOCK_DIM_Y + threadIdx.y;
|
||||||
|
|
||||||
|
if (y >= src.rows)
|
||||||
|
return;
|
||||||
|
|
||||||
|
const T* src_row = src.ptr(y);
|
||||||
|
|
||||||
|
const int xStart = blockIdx.x * (PATCH_PER_BLOCK * BLOCK_DIM_X) + threadIdx.x;
|
||||||
|
|
||||||
|
if (blockIdx.x > 0)
|
||||||
|
{
|
||||||
|
//Load left halo
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < HALO_SIZE; ++j)
|
||||||
|
smem[threadIdx.y][threadIdx.x + j * BLOCK_DIM_X] = saturate_cast<sum_t>(src_row[xStart - (HALO_SIZE - j) * BLOCK_DIM_X]);
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
//Load left halo
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < HALO_SIZE; ++j)
|
||||||
|
smem[threadIdx.y][threadIdx.x + j * BLOCK_DIM_X] = saturate_cast<sum_t>(brd.at_low(xStart - (HALO_SIZE - j) * BLOCK_DIM_X, src_row));
|
||||||
|
}
|
||||||
|
|
||||||
|
if (blockIdx.x + 2 < gridDim.x)
|
||||||
|
{
|
||||||
|
//Load main data
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
||||||
|
smem[threadIdx.y][threadIdx.x + HALO_SIZE * BLOCK_DIM_X + j * BLOCK_DIM_X] = saturate_cast<sum_t>(src_row[xStart + j * BLOCK_DIM_X]);
|
||||||
|
|
||||||
|
//Load right halo
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < HALO_SIZE; ++j)
|
||||||
|
smem[threadIdx.y][threadIdx.x + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_X + j * BLOCK_DIM_X] = saturate_cast<sum_t>(src_row[xStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_X]);
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
//Load main data
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
||||||
|
smem[threadIdx.y][threadIdx.x + HALO_SIZE * BLOCK_DIM_X + j * BLOCK_DIM_X] = saturate_cast<sum_t>(brd.at_high(xStart + j * BLOCK_DIM_X, src_row));
|
||||||
|
|
||||||
|
//Load right halo
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < HALO_SIZE; ++j)
|
||||||
|
smem[threadIdx.y][threadIdx.x + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_X + j * BLOCK_DIM_X] = saturate_cast<sum_t>(brd.at_high(xStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_X, src_row));
|
||||||
|
}
|
||||||
|
|
||||||
|
__syncthreads();
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int j = 0; j < PATCH_PER_BLOCK; ++j)
|
||||||
|
{
|
||||||
|
const int x = xStart + j * BLOCK_DIM_X;
|
||||||
|
|
||||||
|
if (x < src.cols)
|
||||||
|
{
|
||||||
|
sum_t sum = VecTraits<sum_t>::all(0);
|
||||||
|
|
||||||
|
#pragma unroll
|
||||||
|
for (int k = 0; k < KSIZE; ++k)
|
||||||
|
sum = sum + smem[threadIdx.y][threadIdx.x + HALO_SIZE * BLOCK_DIM_X + j * BLOCK_DIM_X - anchor + k] * c_kernel[k];
|
||||||
|
|
||||||
|
dst(y, x) = saturate_cast<D>(sum);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
template <int KSIZE, typename T, typename D, template<typename> class B>
|
||||||
|
void caller(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
int BLOCK_DIM_X;
|
||||||
|
int BLOCK_DIM_Y;
|
||||||
|
int PATCH_PER_BLOCK;
|
||||||
|
|
||||||
|
if (cc >= 20)
|
||||||
|
{
|
||||||
|
BLOCK_DIM_X = 32;
|
||||||
|
BLOCK_DIM_Y = 8;
|
||||||
|
PATCH_PER_BLOCK = 4;
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
BLOCK_DIM_X = 32;
|
||||||
|
BLOCK_DIM_Y = 4;
|
||||||
|
PATCH_PER_BLOCK = 4;
|
||||||
|
}
|
||||||
|
|
||||||
|
const dim3 block(BLOCK_DIM_X, BLOCK_DIM_Y);
|
||||||
|
const dim3 grid(divUp(src.cols, BLOCK_DIM_X * PATCH_PER_BLOCK), divUp(src.rows, BLOCK_DIM_Y));
|
||||||
|
|
||||||
|
B<T> brd(src.cols);
|
||||||
|
|
||||||
|
linearRowFilter<KSIZE, T, D><<<grid, block, 0, stream>>>(src, dst, anchor, brd);
|
||||||
|
cudaSafeCall( cudaGetLastError() );
|
||||||
|
|
||||||
|
if (stream == 0)
|
||||||
|
cudaSafeCall( cudaDeviceSynchronize() );
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
namespace filter
|
||||||
|
{
|
||||||
|
template <typename T, typename D>
|
||||||
|
void linearRow(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
typedef void (*caller_t)(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream);
|
||||||
|
|
||||||
|
static const caller_t callers[5][33] =
|
||||||
|
{
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdRowReflect101>,
|
||||||
|
::caller< 2, T, D, BrdRowReflect101>,
|
||||||
|
::caller< 3, T, D, BrdRowReflect101>,
|
||||||
|
::caller< 4, T, D, BrdRowReflect101>,
|
||||||
|
::caller< 5, T, D, BrdRowReflect101>,
|
||||||
|
::caller< 6, T, D, BrdRowReflect101>,
|
||||||
|
::caller< 7, T, D, BrdRowReflect101>,
|
||||||
|
::caller< 8, T, D, BrdRowReflect101>,
|
||||||
|
::caller< 9, T, D, BrdRowReflect101>,
|
||||||
|
::caller<10, T, D, BrdRowReflect101>,
|
||||||
|
::caller<11, T, D, BrdRowReflect101>,
|
||||||
|
::caller<12, T, D, BrdRowReflect101>,
|
||||||
|
::caller<13, T, D, BrdRowReflect101>,
|
||||||
|
::caller<14, T, D, BrdRowReflect101>,
|
||||||
|
::caller<15, T, D, BrdRowReflect101>,
|
||||||
|
::caller<16, T, D, BrdRowReflect101>,
|
||||||
|
::caller<17, T, D, BrdRowReflect101>,
|
||||||
|
::caller<18, T, D, BrdRowReflect101>,
|
||||||
|
::caller<19, T, D, BrdRowReflect101>,
|
||||||
|
::caller<20, T, D, BrdRowReflect101>,
|
||||||
|
::caller<21, T, D, BrdRowReflect101>,
|
||||||
|
::caller<22, T, D, BrdRowReflect101>,
|
||||||
|
::caller<23, T, D, BrdRowReflect101>,
|
||||||
|
::caller<24, T, D, BrdRowReflect101>,
|
||||||
|
::caller<25, T, D, BrdRowReflect101>,
|
||||||
|
::caller<26, T, D, BrdRowReflect101>,
|
||||||
|
::caller<27, T, D, BrdRowReflect101>,
|
||||||
|
::caller<28, T, D, BrdRowReflect101>,
|
||||||
|
::caller<29, T, D, BrdRowReflect101>,
|
||||||
|
::caller<30, T, D, BrdRowReflect101>,
|
||||||
|
::caller<31, T, D, BrdRowReflect101>,
|
||||||
|
::caller<32, T, D, BrdRowReflect101>
|
||||||
|
},
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdRowReplicate>,
|
||||||
|
::caller< 2, T, D, BrdRowReplicate>,
|
||||||
|
::caller< 3, T, D, BrdRowReplicate>,
|
||||||
|
::caller< 4, T, D, BrdRowReplicate>,
|
||||||
|
::caller< 5, T, D, BrdRowReplicate>,
|
||||||
|
::caller< 6, T, D, BrdRowReplicate>,
|
||||||
|
::caller< 7, T, D, BrdRowReplicate>,
|
||||||
|
::caller< 8, T, D, BrdRowReplicate>,
|
||||||
|
::caller< 9, T, D, BrdRowReplicate>,
|
||||||
|
::caller<10, T, D, BrdRowReplicate>,
|
||||||
|
::caller<11, T, D, BrdRowReplicate>,
|
||||||
|
::caller<12, T, D, BrdRowReplicate>,
|
||||||
|
::caller<13, T, D, BrdRowReplicate>,
|
||||||
|
::caller<14, T, D, BrdRowReplicate>,
|
||||||
|
::caller<15, T, D, BrdRowReplicate>,
|
||||||
|
::caller<16, T, D, BrdRowReplicate>,
|
||||||
|
::caller<17, T, D, BrdRowReplicate>,
|
||||||
|
::caller<18, T, D, BrdRowReplicate>,
|
||||||
|
::caller<19, T, D, BrdRowReplicate>,
|
||||||
|
::caller<20, T, D, BrdRowReplicate>,
|
||||||
|
::caller<21, T, D, BrdRowReplicate>,
|
||||||
|
::caller<22, T, D, BrdRowReplicate>,
|
||||||
|
::caller<23, T, D, BrdRowReplicate>,
|
||||||
|
::caller<24, T, D, BrdRowReplicate>,
|
||||||
|
::caller<25, T, D, BrdRowReplicate>,
|
||||||
|
::caller<26, T, D, BrdRowReplicate>,
|
||||||
|
::caller<27, T, D, BrdRowReplicate>,
|
||||||
|
::caller<28, T, D, BrdRowReplicate>,
|
||||||
|
::caller<29, T, D, BrdRowReplicate>,
|
||||||
|
::caller<30, T, D, BrdRowReplicate>,
|
||||||
|
::caller<31, T, D, BrdRowReplicate>,
|
||||||
|
::caller<32, T, D, BrdRowReplicate>
|
||||||
|
},
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdRowConstant>,
|
||||||
|
::caller< 2, T, D, BrdRowConstant>,
|
||||||
|
::caller< 3, T, D, BrdRowConstant>,
|
||||||
|
::caller< 4, T, D, BrdRowConstant>,
|
||||||
|
::caller< 5, T, D, BrdRowConstant>,
|
||||||
|
::caller< 6, T, D, BrdRowConstant>,
|
||||||
|
::caller< 7, T, D, BrdRowConstant>,
|
||||||
|
::caller< 8, T, D, BrdRowConstant>,
|
||||||
|
::caller< 9, T, D, BrdRowConstant>,
|
||||||
|
::caller<10, T, D, BrdRowConstant>,
|
||||||
|
::caller<11, T, D, BrdRowConstant>,
|
||||||
|
::caller<12, T, D, BrdRowConstant>,
|
||||||
|
::caller<13, T, D, BrdRowConstant>,
|
||||||
|
::caller<14, T, D, BrdRowConstant>,
|
||||||
|
::caller<15, T, D, BrdRowConstant>,
|
||||||
|
::caller<16, T, D, BrdRowConstant>,
|
||||||
|
::caller<17, T, D, BrdRowConstant>,
|
||||||
|
::caller<18, T, D, BrdRowConstant>,
|
||||||
|
::caller<19, T, D, BrdRowConstant>,
|
||||||
|
::caller<20, T, D, BrdRowConstant>,
|
||||||
|
::caller<21, T, D, BrdRowConstant>,
|
||||||
|
::caller<22, T, D, BrdRowConstant>,
|
||||||
|
::caller<23, T, D, BrdRowConstant>,
|
||||||
|
::caller<24, T, D, BrdRowConstant>,
|
||||||
|
::caller<25, T, D, BrdRowConstant>,
|
||||||
|
::caller<26, T, D, BrdRowConstant>,
|
||||||
|
::caller<27, T, D, BrdRowConstant>,
|
||||||
|
::caller<28, T, D, BrdRowConstant>,
|
||||||
|
::caller<29, T, D, BrdRowConstant>,
|
||||||
|
::caller<30, T, D, BrdRowConstant>,
|
||||||
|
::caller<31, T, D, BrdRowConstant>,
|
||||||
|
::caller<32, T, D, BrdRowConstant>
|
||||||
|
},
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdRowReflect>,
|
||||||
|
::caller< 2, T, D, BrdRowReflect>,
|
||||||
|
::caller< 3, T, D, BrdRowReflect>,
|
||||||
|
::caller< 4, T, D, BrdRowReflect>,
|
||||||
|
::caller< 5, T, D, BrdRowReflect>,
|
||||||
|
::caller< 6, T, D, BrdRowReflect>,
|
||||||
|
::caller< 7, T, D, BrdRowReflect>,
|
||||||
|
::caller< 8, T, D, BrdRowReflect>,
|
||||||
|
::caller< 9, T, D, BrdRowReflect>,
|
||||||
|
::caller<10, T, D, BrdRowReflect>,
|
||||||
|
::caller<11, T, D, BrdRowReflect>,
|
||||||
|
::caller<12, T, D, BrdRowReflect>,
|
||||||
|
::caller<13, T, D, BrdRowReflect>,
|
||||||
|
::caller<14, T, D, BrdRowReflect>,
|
||||||
|
::caller<15, T, D, BrdRowReflect>,
|
||||||
|
::caller<16, T, D, BrdRowReflect>,
|
||||||
|
::caller<17, T, D, BrdRowReflect>,
|
||||||
|
::caller<18, T, D, BrdRowReflect>,
|
||||||
|
::caller<19, T, D, BrdRowReflect>,
|
||||||
|
::caller<20, T, D, BrdRowReflect>,
|
||||||
|
::caller<21, T, D, BrdRowReflect>,
|
||||||
|
::caller<22, T, D, BrdRowReflect>,
|
||||||
|
::caller<23, T, D, BrdRowReflect>,
|
||||||
|
::caller<24, T, D, BrdRowReflect>,
|
||||||
|
::caller<25, T, D, BrdRowReflect>,
|
||||||
|
::caller<26, T, D, BrdRowReflect>,
|
||||||
|
::caller<27, T, D, BrdRowReflect>,
|
||||||
|
::caller<28, T, D, BrdRowReflect>,
|
||||||
|
::caller<29, T, D, BrdRowReflect>,
|
||||||
|
::caller<30, T, D, BrdRowReflect>,
|
||||||
|
::caller<31, T, D, BrdRowReflect>,
|
||||||
|
::caller<32, T, D, BrdRowReflect>
|
||||||
|
},
|
||||||
|
{
|
||||||
|
0,
|
||||||
|
::caller< 1, T, D, BrdRowWrap>,
|
||||||
|
::caller< 2, T, D, BrdRowWrap>,
|
||||||
|
::caller< 3, T, D, BrdRowWrap>,
|
||||||
|
::caller< 4, T, D, BrdRowWrap>,
|
||||||
|
::caller< 5, T, D, BrdRowWrap>,
|
||||||
|
::caller< 6, T, D, BrdRowWrap>,
|
||||||
|
::caller< 7, T, D, BrdRowWrap>,
|
||||||
|
::caller< 8, T, D, BrdRowWrap>,
|
||||||
|
::caller< 9, T, D, BrdRowWrap>,
|
||||||
|
::caller<10, T, D, BrdRowWrap>,
|
||||||
|
::caller<11, T, D, BrdRowWrap>,
|
||||||
|
::caller<12, T, D, BrdRowWrap>,
|
||||||
|
::caller<13, T, D, BrdRowWrap>,
|
||||||
|
::caller<14, T, D, BrdRowWrap>,
|
||||||
|
::caller<15, T, D, BrdRowWrap>,
|
||||||
|
::caller<16, T, D, BrdRowWrap>,
|
||||||
|
::caller<17, T, D, BrdRowWrap>,
|
||||||
|
::caller<18, T, D, BrdRowWrap>,
|
||||||
|
::caller<19, T, D, BrdRowWrap>,
|
||||||
|
::caller<20, T, D, BrdRowWrap>,
|
||||||
|
::caller<21, T, D, BrdRowWrap>,
|
||||||
|
::caller<22, T, D, BrdRowWrap>,
|
||||||
|
::caller<23, T, D, BrdRowWrap>,
|
||||||
|
::caller<24, T, D, BrdRowWrap>,
|
||||||
|
::caller<25, T, D, BrdRowWrap>,
|
||||||
|
::caller<26, T, D, BrdRowWrap>,
|
||||||
|
::caller<27, T, D, BrdRowWrap>,
|
||||||
|
::caller<28, T, D, BrdRowWrap>,
|
||||||
|
::caller<29, T, D, BrdRowWrap>,
|
||||||
|
::caller<30, T, D, BrdRowWrap>,
|
||||||
|
::caller<31, T, D, BrdRowWrap>,
|
||||||
|
::caller<32, T, D, BrdRowWrap>
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
loadKernel(kernel, ksize, stream);
|
||||||
|
|
||||||
|
callers[brd_type][ksize]((PtrStepSz<T>)src, (PtrStepSz<D>)dst, anchor, cc, stream);
|
||||||
|
}
|
||||||
|
}
|
||||||
@ -830,20 +830,14 @@ void cv::gpu::filter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& ke
|
|||||||
////////////////////////////////////////////////////////////////////////////////////////////////////
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||||
// Separable Linear Filter
|
// Separable Linear Filter
|
||||||
|
|
||||||
namespace cv { namespace gpu { namespace device
|
namespace filter
|
||||||
{
|
{
|
||||||
namespace row_filter
|
template <typename T, typename D>
|
||||||
{
|
void linearRow(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
template <typename T, typename D>
|
|
||||||
void linearRowFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
}
|
|
||||||
|
|
||||||
namespace column_filter
|
template <typename T, typename D>
|
||||||
{
|
void linearColumn(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
||||||
template <typename T, typename D>
|
}
|
||||||
void linearColumnFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
|
|
||||||
}
|
|
||||||
}}}
|
|
||||||
|
|
||||||
namespace
|
namespace
|
||||||
{
|
{
|
||||||
@ -899,8 +893,6 @@ namespace
|
|||||||
|
|
||||||
Ptr<BaseRowFilter_GPU> cv::gpu::getLinearRowFilter_GPU(int srcType, int bufType, const Mat& rowKernel, int anchor, int borderType)
|
Ptr<BaseRowFilter_GPU> cv::gpu::getLinearRowFilter_GPU(int srcType, int bufType, const Mat& rowKernel, int anchor, int borderType)
|
||||||
{
|
{
|
||||||
using namespace ::cv::gpu::device::row_filter;
|
|
||||||
|
|
||||||
static const nppFilter1D_t nppFilter1D_callers[] = {0, nppiFilterRow_8u_C1R, 0, 0, nppiFilterRow_8u_C4R};
|
static const nppFilter1D_t nppFilter1D_callers[] = {0, nppiFilterRow_8u_C1R, 0, 0, nppiFilterRow_8u_C4R};
|
||||||
|
|
||||||
if ((bufType == srcType) && (srcType == CV_8UC1 || srcType == CV_8UC4))
|
if ((bufType == srcType) && (srcType == CV_8UC1 || srcType == CV_8UC4))
|
||||||
@ -940,28 +932,28 @@ Ptr<BaseRowFilter_GPU> cv::gpu::getLinearRowFilter_GPU(int srcType, int bufType,
|
|||||||
switch (srcType)
|
switch (srcType)
|
||||||
{
|
{
|
||||||
case CV_8UC1:
|
case CV_8UC1:
|
||||||
func = linearRowFilter_gpu<uchar, float>;
|
func = filter::linearRow<uchar, float>;
|
||||||
break;
|
break;
|
||||||
case CV_8UC3:
|
case CV_8UC3:
|
||||||
func = linearRowFilter_gpu<uchar3, float3>;
|
func = filter::linearRow<uchar3, float3>;
|
||||||
break;
|
break;
|
||||||
case CV_8UC4:
|
case CV_8UC4:
|
||||||
func = linearRowFilter_gpu<uchar4, float4>;
|
func = filter::linearRow<uchar4, float4>;
|
||||||
break;
|
break;
|
||||||
case CV_16SC3:
|
case CV_16SC3:
|
||||||
func = linearRowFilter_gpu<short3, float3>;
|
func = filter::linearRow<short3, float3>;
|
||||||
break;
|
break;
|
||||||
case CV_32SC1:
|
case CV_32SC1:
|
||||||
func = linearRowFilter_gpu<int, float>;
|
func = filter::linearRow<int, float>;
|
||||||
break;
|
break;
|
||||||
case CV_32FC1:
|
case CV_32FC1:
|
||||||
func = linearRowFilter_gpu<float, float>;
|
func = filter::linearRow<float, float>;
|
||||||
break;
|
break;
|
||||||
case CV_32FC3:
|
case CV_32FC3:
|
||||||
func = linearRowFilter_gpu<float3, float3>;
|
func = filter::linearRow<float3, float3>;
|
||||||
break;
|
break;
|
||||||
case CV_32FC4:
|
case CV_32FC4:
|
||||||
func = linearRowFilter_gpu<float4, float4>;
|
func = filter::linearRow<float4, float4>;
|
||||||
break;
|
break;
|
||||||
}
|
}
|
||||||
|
|
||||||
@ -1020,8 +1012,6 @@ namespace
|
|||||||
|
|
||||||
Ptr<BaseColumnFilter_GPU> cv::gpu::getLinearColumnFilter_GPU(int bufType, int dstType, const Mat& columnKernel, int anchor, int borderType)
|
Ptr<BaseColumnFilter_GPU> cv::gpu::getLinearColumnFilter_GPU(int bufType, int dstType, const Mat& columnKernel, int anchor, int borderType)
|
||||||
{
|
{
|
||||||
using namespace ::cv::gpu::device::column_filter;
|
|
||||||
|
|
||||||
static const nppFilter1D_t nppFilter1D_callers[] = {0, nppiFilterColumn_8u_C1R, 0, 0, nppiFilterColumn_8u_C4R};
|
static const nppFilter1D_t nppFilter1D_callers[] = {0, nppiFilterColumn_8u_C1R, 0, 0, nppiFilterColumn_8u_C4R};
|
||||||
|
|
||||||
if ((bufType == dstType) && (bufType == CV_8UC1 || bufType == CV_8UC4))
|
if ((bufType == dstType) && (bufType == CV_8UC1 || bufType == CV_8UC4))
|
||||||
@ -1061,28 +1051,28 @@ Ptr<BaseColumnFilter_GPU> cv::gpu::getLinearColumnFilter_GPU(int bufType, int ds
|
|||||||
switch (dstType)
|
switch (dstType)
|
||||||
{
|
{
|
||||||
case CV_8UC1:
|
case CV_8UC1:
|
||||||
func = linearColumnFilter_gpu<float, uchar>;
|
func = filter::linearColumn<float, uchar>;
|
||||||
break;
|
break;
|
||||||
case CV_8UC3:
|
case CV_8UC3:
|
||||||
func = linearColumnFilter_gpu<float3, uchar3>;
|
func = filter::linearColumn<float3, uchar3>;
|
||||||
break;
|
break;
|
||||||
case CV_8UC4:
|
case CV_8UC4:
|
||||||
func = linearColumnFilter_gpu<float4, uchar4>;
|
func = filter::linearColumn<float4, uchar4>;
|
||||||
break;
|
break;
|
||||||
case CV_16SC3:
|
case CV_16SC3:
|
||||||
func = linearColumnFilter_gpu<float3, short3>;
|
func = filter::linearColumn<float3, short3>;
|
||||||
break;
|
break;
|
||||||
case CV_32SC1:
|
case CV_32SC1:
|
||||||
func = linearColumnFilter_gpu<float, int>;
|
func = filter::linearColumn<float, int>;
|
||||||
break;
|
break;
|
||||||
case CV_32FC1:
|
case CV_32FC1:
|
||||||
func = linearColumnFilter_gpu<float, float>;
|
func = filter::linearColumn<float, float>;
|
||||||
break;
|
break;
|
||||||
case CV_32FC3:
|
case CV_32FC3:
|
||||||
func = linearColumnFilter_gpu<float3, float3>;
|
func = filter::linearColumn<float3, float3>;
|
||||||
break;
|
break;
|
||||||
case CV_32FC4:
|
case CV_32FC4:
|
||||||
func = linearColumnFilter_gpu<float4, float4>;
|
func = filter::linearColumn<float4, float4>;
|
||||||
break;
|
break;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
Loading…
Reference in New Issue
Block a user