From eb449968b6f506d1e5a363768ebda5efb0f668b2 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Tue, 27 Aug 2013 12:59:09 +0400 Subject: [PATCH] removed unroll_detail.hpp --- modules/cudaarithm/src/cuda/unroll_detail.hpp | 135 -------------------------- 1 file changed, 135 deletions(-) delete mode 100644 modules/cudaarithm/src/cuda/unroll_detail.hpp diff --git a/modules/cudaarithm/src/cuda/unroll_detail.hpp b/modules/cudaarithm/src/cuda/unroll_detail.hpp deleted file mode 100644 index 544d838..0000000 --- a/modules/cudaarithm/src/cuda/unroll_detail.hpp +++ /dev/null @@ -1,135 +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. -// 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*/ - -#ifndef __UNROLL_DETAIL_HPP__ -#define __UNROLL_DETAIL_HPP__ - -#include -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/vec_traits.hpp" - -namespace detail -{ - template struct Unroll; - template <> struct Unroll<1> - { - template - static __device__ __forceinline__ volatile R* smem_tuple(R* smem) - { - return smem; - } - - template - static __device__ __forceinline__ R& tie(R& val) - { - return val; - } - - template - static __device__ __forceinline__ const Op& op(const Op& op) - { - return op; - } - }; - template <> struct Unroll<2> - { - template - static __device__ __forceinline__ thrust::tuple smem_tuple(R* smem) - { - return cv::cuda::device::smem_tuple(smem, smem + BLOCK_SIZE); - } - - template - static __device__ __forceinline__ thrust::tuple::elem_type&, typename cv::cuda::device::VecTraits::elem_type&> tie(R& val) - { - return thrust::tie(val.x, val.y); - } - - template - static __device__ __forceinline__ const thrust::tuple op(const Op& op) - { - return thrust::make_tuple(op, op); - } - }; - template <> struct Unroll<3> - { - template - static __device__ __forceinline__ thrust::tuple smem_tuple(R* smem) - { - return cv::cuda::device::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE); - } - - template - static __device__ __forceinline__ thrust::tuple::elem_type&, typename cv::cuda::device::VecTraits::elem_type&, typename cv::cuda::device::VecTraits::elem_type&> tie(R& val) - { - return thrust::tie(val.x, val.y, val.z); - } - - template - static __device__ __forceinline__ const thrust::tuple op(const Op& op) - { - return thrust::make_tuple(op, op, op); - } - }; - template <> struct Unroll<4> - { - template - static __device__ __forceinline__ thrust::tuple smem_tuple(R* smem) - { - return cv::cuda::device::smem_tuple(smem, smem + BLOCK_SIZE, smem + 2 * BLOCK_SIZE, smem + 3 * BLOCK_SIZE); - } - - template - static __device__ __forceinline__ thrust::tuple::elem_type&, typename cv::cuda::device::VecTraits::elem_type&, typename cv::cuda::device::VecTraits::elem_type&, typename cv::cuda::device::VecTraits::elem_type&> tie(R& val) - { - return thrust::tie(val.x, val.y, val.z, val.w); - } - - template - static __device__ __forceinline__ const thrust::tuple op(const Op& op) - { - return thrust::make_tuple(op, op, op, op); - } - }; -} - -#endif // __UNROLL_DETAIL_HPP__ -- 2.7.4