/*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*/

#if !defined CUDA_DISABLER

#include "internal_shared.hpp"

namespace cv { namespace gpu { namespace device
{
    namespace split_merge
    {
        template <typename T, size_t elem_size = sizeof(T)>
        struct TypeTraits
        {
            typedef T type;
            typedef T type2;
            typedef T type3;
            typedef T type4;
        };

        template <typename T>
        struct TypeTraits<T, 1>
        {
            typedef char type;
            typedef char2 type2;
            typedef char3 type3;
            typedef char4 type4;
        };

        template <typename T>
        struct TypeTraits<T, 2>
        {
            typedef short type;
            typedef short2 type2;
            typedef short3 type3;
            typedef short4 type4;
        };

        template <typename T>
        struct TypeTraits<T, 4>
        {
            typedef int type;
            typedef int2 type2;
            typedef int3 type3;
            typedef int4 type4;
        };

        template <typename T>
        struct TypeTraits<T, 8>
        {
            typedef double type;
            typedef double2 type2;
            //typedef double3 type3;
            //typedef double4 type3;
        };

        typedef void (*MergeFunction)(const PtrStepSzb* src, PtrStepSzb& dst, const cudaStream_t& stream);
        typedef void (*SplitFunction)(const PtrStepSzb& src, PtrStepSzb* dst, const cudaStream_t& stream);

        //------------------------------------------------------------
        // Merge

        template <typename T>
        __global__ void mergeC2_(const uchar* src0, size_t src0_step,
                                 const uchar* src1, size_t src1_step,
                                 int rows, int cols, uchar* dst, size_t dst_step)
        {
            typedef typename TypeTraits<T>::type2 dst_type;

            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const T* src0_y = (const T*)(src0 + y * src0_step);
            const T* src1_y = (const T*)(src1 + y * src1_step);
            dst_type* dst_y = (dst_type*)(dst + y * dst_step);

            if (x < cols && y < rows)
            {
                dst_type dst_elem;
                dst_elem.x = src0_y[x];
                dst_elem.y = src1_y[x];
                dst_y[x] = dst_elem;
            }
        }


        template <typename T>
        __global__ void mergeC3_(const uchar* src0, size_t src0_step,
                                 const uchar* src1, size_t src1_step,
                                 const uchar* src2, size_t src2_step,
                                 int rows, int cols, uchar* dst, size_t dst_step)
        {
            typedef typename TypeTraits<T>::type3 dst_type;

            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const T* src0_y = (const T*)(src0 + y * src0_step);
            const T* src1_y = (const T*)(src1 + y * src1_step);
            const T* src2_y = (const T*)(src2 + y * src2_step);
            dst_type* dst_y = (dst_type*)(dst + y * dst_step);

            if (x < cols && y < rows)
            {
                dst_type dst_elem;
                dst_elem.x = src0_y[x];
                dst_elem.y = src1_y[x];
                dst_elem.z = src2_y[x];
                dst_y[x] = dst_elem;
            }
        }


        template <>
        __global__ void mergeC3_<double>(const uchar* src0, size_t src0_step,
                                 const uchar* src1, size_t src1_step,
                                 const uchar* src2, size_t src2_step,
                                 int rows, int cols, uchar* dst, size_t dst_step)
        {
            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const double* src0_y = (const double*)(src0 + y * src0_step);
            const double* src1_y = (const double*)(src1 + y * src1_step);
            const double* src2_y = (const double*)(src2 + y * src2_step);
            double* dst_y = (double*)(dst + y * dst_step);

            if (x < cols && y < rows)
            {
                dst_y[3 * x] = src0_y[x];
                dst_y[3 * x + 1] = src1_y[x];
                dst_y[3 * x + 2] = src2_y[x];
            }
        }


        template <typename T>
        __global__ void mergeC4_(const uchar* src0, size_t src0_step,
                                 const uchar* src1, size_t src1_step,
                                 const uchar* src2, size_t src2_step,
                                 const uchar* src3, size_t src3_step,
                                 int rows, int cols, uchar* dst, size_t dst_step)
        {
            typedef typename TypeTraits<T>::type4 dst_type;

            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const T* src0_y = (const T*)(src0 + y * src0_step);
            const T* src1_y = (const T*)(src1 + y * src1_step);
            const T* src2_y = (const T*)(src2 + y * src2_step);
            const T* src3_y = (const T*)(src3 + y * src3_step);
            dst_type* dst_y = (dst_type*)(dst + y * dst_step);

            if (x < cols && y < rows)
            {
                dst_type dst_elem;
                dst_elem.x = src0_y[x];
                dst_elem.y = src1_y[x];
                dst_elem.z = src2_y[x];
                dst_elem.w = src3_y[x];
                dst_y[x] = dst_elem;
            }
        }


        template <>
        __global__ void mergeC4_<double>(const uchar* src0, size_t src0_step,
                                 const uchar* src1, size_t src1_step,
                                 const uchar* src2, size_t src2_step,
                                 const uchar* src3, size_t src3_step,
                                 int rows, int cols, uchar* dst, size_t dst_step)
        {
            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const double* src0_y = (const double*)(src0 + y * src0_step);
            const double* src1_y = (const double*)(src1 + y * src1_step);
            const double* src2_y = (const double*)(src2 + y * src2_step);
            const double* src3_y = (const double*)(src3 + y * src3_step);
            double2* dst_y = (double2*)(dst + y * dst_step);

            if (x < cols && y < rows)
            {
                dst_y[2 * x] = make_double2(src0_y[x], src1_y[x]);
                dst_y[2 * x + 1] = make_double2(src2_y[x], src3_y[x]);
            }
        }


        template <typename T>
        static void mergeC2_(const PtrStepSzb* src, PtrStepSzb& dst, const cudaStream_t& stream)
        {
            dim3 block(32, 8);
            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
            mergeC2_<T><<<grid, block, 0, stream>>>(
                    src[0].data, src[0].step,
                    src[1].data, src[1].step,
                    dst.rows, dst.cols, dst.data, dst.step);
            cudaSafeCall( cudaGetLastError() );

            if (stream == 0)
                cudaSafeCall(cudaDeviceSynchronize());
        }


        template <typename T>
        static void mergeC3_(const PtrStepSzb* src, PtrStepSzb& dst, const cudaStream_t& stream)
        {
            dim3 block(32, 8);
            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
            mergeC3_<T><<<grid, block, 0, stream>>>(
                    src[0].data, src[0].step,
                    src[1].data, src[1].step,
                    src[2].data, src[2].step,
                    dst.rows, dst.cols, dst.data, dst.step);
            cudaSafeCall( cudaGetLastError() );

            if (stream == 0)
                cudaSafeCall(cudaDeviceSynchronize());
        }


        template <typename T>
        static void mergeC4_(const PtrStepSzb* src, PtrStepSzb& dst, const cudaStream_t& stream)
        {
            dim3 block(32, 8);
            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
            mergeC4_<T><<<grid, block, 0, stream>>>(
                    src[0].data, src[0].step,
                    src[1].data, src[1].step,
                    src[2].data, src[2].step,
                    src[3].data, src[3].step,
                    dst.rows, dst.cols, dst.data, dst.step);
            cudaSafeCall( cudaGetLastError() );

            if (stream == 0)
                cudaSafeCall(cudaDeviceSynchronize());
        }


        void merge_caller(const PtrStepSzb* src, PtrStepSzb& dst,
                                     int total_channels, size_t elem_size,
                                     const cudaStream_t& stream)
        {
            static MergeFunction merge_func_tbl[] =
            {
                mergeC2_<char>, mergeC2_<short>, mergeC2_<int>, 0, mergeC2_<double>,
                mergeC3_<char>, mergeC3_<short>, mergeC3_<int>, 0, mergeC3_<double>,
                mergeC4_<char>, mergeC4_<short>, mergeC4_<int>, 0, mergeC4_<double>,
            };

            size_t merge_func_id = (total_channels - 2) * 5 + (elem_size >> 1);
            MergeFunction merge_func = merge_func_tbl[merge_func_id];

            if (merge_func == 0)
                cv::gpu::error("Unsupported channel count or data type", __FILE__, __LINE__, "merge_caller");

            merge_func(src, dst, stream);
        }



        //------------------------------------------------------------
        // Split


        template <typename T>
        __global__ void splitC2_(const uchar* src, size_t src_step,
                                int rows, int cols,
                                uchar* dst0, size_t dst0_step,
                                uchar* dst1, size_t dst1_step)
        {
            typedef typename TypeTraits<T>::type2 src_type;

            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const src_type* src_y = (const src_type*)(src + y * src_step);
            T* dst0_y = (T*)(dst0 + y * dst0_step);
            T* dst1_y = (T*)(dst1 + y * dst1_step);

            if (x < cols && y < rows)
            {
                src_type src_elem = src_y[x];
                dst0_y[x] = src_elem.x;
                dst1_y[x] = src_elem.y;
            }
        }


        template <typename T>
        __global__ void splitC3_(const uchar* src, size_t src_step,
                                int rows, int cols,
                                uchar* dst0, size_t dst0_step,
                                uchar* dst1, size_t dst1_step,
                                uchar* dst2, size_t dst2_step)
        {
            typedef typename TypeTraits<T>::type3 src_type;

            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const src_type* src_y = (const src_type*)(src + y * src_step);
            T* dst0_y = (T*)(dst0 + y * dst0_step);
            T* dst1_y = (T*)(dst1 + y * dst1_step);
            T* dst2_y = (T*)(dst2 + y * dst2_step);

            if (x < cols && y < rows)
            {
                src_type src_elem = src_y[x];
                dst0_y[x] = src_elem.x;
                dst1_y[x] = src_elem.y;
                dst2_y[x] = src_elem.z;
            }
        }


        template <>
        __global__ void splitC3_<double>(
                const uchar* src, size_t src_step, int rows, int cols,
                uchar* dst0, size_t dst0_step,
                uchar* dst1, size_t dst1_step,
                uchar* dst2, size_t dst2_step)
        {
            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const double* src_y = (const double*)(src + y * src_step);
            double* dst0_y = (double*)(dst0 + y * dst0_step);
            double* dst1_y = (double*)(dst1 + y * dst1_step);
            double* dst2_y = (double*)(dst2 + y * dst2_step);

            if (x < cols && y < rows)
            {
                dst0_y[x] = src_y[3 * x];
                dst1_y[x] = src_y[3 * x + 1];
                dst2_y[x] = src_y[3 * x + 2];
            }
        }


        template <typename T>
        __global__ void splitC4_(const uchar* src, size_t src_step, int rows, int cols,
                                uchar* dst0, size_t dst0_step,
                                uchar* dst1, size_t dst1_step,
                                uchar* dst2, size_t dst2_step,
                                uchar* dst3, size_t dst3_step)
        {
            typedef typename TypeTraits<T>::type4 src_type;

            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const src_type* src_y = (const src_type*)(src + y * src_step);
            T* dst0_y = (T*)(dst0 + y * dst0_step);
            T* dst1_y = (T*)(dst1 + y * dst1_step);
            T* dst2_y = (T*)(dst2 + y * dst2_step);
            T* dst3_y = (T*)(dst3 + y * dst3_step);

            if (x < cols && y < rows)
            {
                src_type src_elem = src_y[x];
                dst0_y[x] = src_elem.x;
                dst1_y[x] = src_elem.y;
                dst2_y[x] = src_elem.z;
                dst3_y[x] = src_elem.w;
            }
        }


        template <>
        __global__ void splitC4_<double>(
                const uchar* src, size_t src_step, int rows, int cols,
                uchar* dst0, size_t dst0_step,
                uchar* dst1, size_t dst1_step,
                uchar* dst2, size_t dst2_step,
                uchar* dst3, size_t dst3_step)
        {
            const int x = blockIdx.x * blockDim.x + threadIdx.x;
            const int y = blockIdx.y * blockDim.y + threadIdx.y;

            const double2* src_y = (const double2*)(src + y * src_step);
            double* dst0_y = (double*)(dst0 + y * dst0_step);
            double* dst1_y = (double*)(dst1 + y * dst1_step);
            double* dst2_y = (double*)(dst2 + y * dst2_step);
            double* dst3_y = (double*)(dst3 + y * dst3_step);

            if (x < cols && y < rows)
            {
                double2 src_elem1 = src_y[2 * x];
                double2 src_elem2 = src_y[2 * x + 1];
                dst0_y[x] = src_elem1.x;
                dst1_y[x] = src_elem1.y;
                dst2_y[x] = src_elem2.x;
                dst3_y[x] = src_elem2.y;
            }
        }

        template <typename T>
        static void splitC2_(const PtrStepSzb& src, PtrStepSzb* dst, const cudaStream_t& stream)
        {
            dim3 block(32, 8);
            dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
            splitC2_<T><<<grid, block, 0, stream>>>(
                    src.data, src.step, src.rows, src.cols,
                    dst[0].data, dst[0].step,
                    dst[1].data, dst[1].step);
            cudaSafeCall( cudaGetLastError() );

            if (stream == 0)
                cudaSafeCall(cudaDeviceSynchronize());
        }


        template <typename T>
        static void splitC3_(const PtrStepSzb& src, PtrStepSzb* dst, const cudaStream_t& stream)
        {
            dim3 block(32, 8);
            dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
            splitC3_<T><<<grid, block, 0, stream>>>(
                    src.data, src.step, src.rows, src.cols,
                    dst[0].data, dst[0].step,
                    dst[1].data, dst[1].step,
                    dst[2].data, dst[2].step);
            cudaSafeCall( cudaGetLastError() );

            if (stream == 0)
                cudaSafeCall(cudaDeviceSynchronize());
        }


        template <typename T>
        static void splitC4_(const PtrStepSzb& src, PtrStepSzb* dst, const cudaStream_t& stream)
        {
            dim3 block(32, 8);
            dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
            splitC4_<T><<<grid, block, 0, stream>>>(
                     src.data, src.step, src.rows, src.cols,
                     dst[0].data, dst[0].step,
                     dst[1].data, dst[1].step,
                     dst[2].data, dst[2].step,
                     dst[3].data, dst[3].step);
            cudaSafeCall( cudaGetLastError() );

            if (stream == 0)
                cudaSafeCall(cudaDeviceSynchronize());
        }


        void split_caller(const PtrStepSzb& src, PtrStepSzb* dst, int num_channels, size_t elem_size1, const cudaStream_t& stream)
        {
            static SplitFunction split_func_tbl[] =
            {
                splitC2_<char>, splitC2_<short>, splitC2_<int>, 0, splitC2_<double>,
                splitC3_<char>, splitC3_<short>, splitC3_<int>, 0, splitC3_<double>,
                splitC4_<char>, splitC4_<short>, splitC4_<int>, 0, splitC4_<double>,
            };

            size_t split_func_id = (num_channels - 2) * 5 + (elem_size1 >> 1);
            SplitFunction split_func = split_func_tbl[split_func_id];

            if (split_func == 0)
                cv::gpu::error("Unsupported channel count or data type", __FILE__, __LINE__, "split_caller");

            split_func(src, dst, stream);
        }
    } // namespace split_merge
}}} // namespace cv { namespace gpu { namespace device


#endif /* CUDA_DISABLER */