mirror of https://github.com/opencv/opencv.git
module reorganization: added folder with pure device functions, cuda_shared.hpp renamed to internal_shared.hpp
Anatoly Baksheev
14 years ago
parent
fadd19b976
commit
652fb1212e
33 changed files with 1579 additions and 1762 deletions
@ -1,134 +1,144 @@ |
||||
|
||||
set(name "gpu") |
||||
set(DEPS "opencv_core" "opencv_imgproc" "opencv_objdetect" "opencv_features2d" "opencv_flann") |
||||
|
||||
set(OPENCV_LINKER_LIBS ${OPENCV_LINKER_LIBS} opencv_gpu) |
||||
|
||||
set(the_target "opencv_${name}") |
||||
|
||||
project(${the_target}) |
||||
|
||||
add_definitions(-DCVAPI_EXPORTS) |
||||
|
||||
include_directories("${CMAKE_CURRENT_SOURCE_DIR}/include" |
||||
"${CMAKE_CURRENT_SOURCE_DIR}/src/cuda" |
||||
"${CMAKE_CURRENT_SOURCE_DIR}/src" |
||||
"${CMAKE_CURRENT_BINARY_DIR}") |
||||
|
||||
foreach(d ${DEPS}) |
||||
if(${d} MATCHES "opencv_") |
||||
string(REPLACE "opencv_" "${CMAKE_CURRENT_SOURCE_DIR}/../" d_dir ${d}) |
||||
include_directories("${d_dir}/include") |
||||
endif() |
||||
endforeach() |
||||
|
||||
file(GLOB lib_srcs "src/*.cpp") |
||||
file(GLOB lib_int_hdrs "src/*.h*") |
||||
file(GLOB lib_cuda "src/cuda/*.cu*") |
||||
file(GLOB lib_cuda_hdrs "src/cuda/*.h*") |
||||
source_group("Src\\Host" FILES ${lib_srcs} ${lib_int_hdrs}) |
||||
source_group("Src\\Cuda" FILES ${lib_cuda} ${lib_cuda_hdrs}) |
||||
|
||||
file(GLOB lib_hdrs "include/opencv2/${name}/*.h*") |
||||
source_group("Include" FILES ${lib_hdrs}) |
||||
|
||||
if (HAVE_CUDA) |
||||
get_filename_component(_path_to_findnpp "${CMAKE_CURRENT_LIST_FILE}" PATH) |
||||
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${_path_to_findnpp}) |
||||
find_package(NPP 3.2.16 REQUIRED) |
||||
message(STATUS "NPP detected: " ${NPP_VERSION}) |
||||
|
||||
include_directories(${CUDA_INCLUDE_DIRS} ${CUDA_NPP_INCLUDES}) |
||||
|
||||
if (UNIX OR APPLE) |
||||
set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-Xcompiler;-fPIC;") |
||||
#set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}" "-fPIC") |
||||
endif() |
||||
|
||||
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-keep") |
||||
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-Xcompiler;/EHsc-;") |
||||
|
||||
string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") |
||||
string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") |
||||
string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") |
||||
|
||||
if(MSVC) |
||||
#string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") |
||||
#string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") |
||||
#string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") |
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4211 /wd4201 /wd4100 /wd4505 /wd4408") |
||||
|
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_C_FLAGS "${CMAKE_C_FLAGS}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") |
||||
endif() |
||||
|
||||
CUDA_COMPILE(cuda_objs ${lib_cuda}) |
||||
#CUDA_BUILD_CLEAN_TARGET() |
||||
endif() |
||||
|
||||
|
||||
add_library(${the_target} ${lib_srcs} ${lib_hdrs} ${lib_int_hdrs} ${lib_cuda} ${lib_cuda_hdrs} ${cuda_objs}) |
||||
|
||||
|
||||
if(PCHSupport_FOUND) |
||||
set(pch_header ${CMAKE_CURRENT_SOURCE_DIR}/src/precomp.hpp) |
||||
if(${CMAKE_GENERATOR} MATCHES "Visual*" OR ${CMAKE_GENERATOR} MATCHES "Xcode*") |
||||
if(${CMAKE_GENERATOR} MATCHES "Visual*") |
||||
set(${the_target}_pch "src/precomp.cpp") |
||||
endif() |
||||
add_native_precompiled_header(${the_target} ${pch_header}) |
||||
elseif(CMAKE_COMPILER_IS_GNUCXX AND ${CMAKE_GENERATOR} MATCHES ".*Makefiles") |
||||
add_precompiled_header(${the_target} ${pch_header}) |
||||
endif() |
||||
endif() |
||||
|
||||
# For dynamic link numbering convenions |
||||
set_target_properties(${the_target} PROPERTIES |
||||
VERSION ${OPENCV_VERSION} |
||||
SOVERSION ${OPENCV_SOVERSION} |
||||
OUTPUT_NAME "${the_target}${OPENCV_DLLVERSION}" |
||||
) |
||||
|
||||
# Additional target properties |
||||
set_target_properties(${the_target} PROPERTIES |
||||
DEBUG_POSTFIX "${OPENCV_DEBUG_POSTFIX}" |
||||
ARCHIVE_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib/" |
||||
RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/bin/" |
||||
INSTALL_NAME_DIR "${CMAKE_INSTALL_PREFIX}/lib" |
||||
) |
||||
|
||||
# Add the required libraries for linking: |
||||
target_link_libraries(${the_target} ${OPENCV_LINKER_LIBS} ${IPP_LIBS} ${DEPS}) |
||||
|
||||
if (HAVE_CUDA) |
||||
target_link_libraries(${the_target} ${CUDA_LIBRARIES} ${CUDA_NPP_LIBRARIES}) |
||||
endif() |
||||
|
||||
if(MSVC) |
||||
if(CMAKE_CROSSCOMPILING) |
||||
set_target_properties(${the_target} PROPERTIES |
||||
LINK_FLAGS "/NODEFAULTLIB:secchk" |
||||
) |
||||
endif() |
||||
set_target_properties(${the_target} PROPERTIES |
||||
LINK_FLAGS "/NODEFAULTLIB:libc" |
||||
) |
||||
endif() |
||||
|
||||
# Dependencies of this target: |
||||
add_dependencies(${the_target} ${DEPS}) |
||||
|
||||
install(TARGETS ${the_target} |
||||
RUNTIME DESTINATION bin COMPONENT main |
||||
LIBRARY DESTINATION lib COMPONENT main |
||||
ARCHIVE DESTINATION lib COMPONENT main) |
||||
|
||||
install(FILES ${lib_hdrs} |
||||
DESTINATION include/opencv2/${name} |
||||
COMPONENT main) |
||||
|
||||
|
||||
|
||||
set(name "gpu") |
||||
|
||||
#"opencv_features2d" "opencv_flann" "opencv_objdetect" - only headers needed |
||||
set(DEPS "opencv_core" "opencv_imgproc" "opencv_objdetect" "opencv_features2d" "opencv_flann") |
||||
|
||||
set(OPENCV_LINKER_LIBS ${OPENCV_LINKER_LIBS} opencv_gpu) |
||||
|
||||
set(the_target "opencv_${name}") |
||||
|
||||
project(${the_target}) |
||||
|
||||
add_definitions(-DCVAPI_EXPORTS) |
||||
|
||||
include_directories("${CMAKE_CURRENT_SOURCE_DIR}/include" |
||||
"${CMAKE_CURRENT_SOURCE_DIR}/src/cuda" |
||||
"${CMAKE_CURRENT_SOURCE_DIR}/src" |
||||
"${CMAKE_CURRENT_BINARY_DIR}") |
||||
|
||||
foreach(d ${DEPS}) |
||||
if(${d} MATCHES "opencv_") |
||||
string(REPLACE "opencv_" "${CMAKE_CURRENT_SOURCE_DIR}/../" d_dir ${d}) |
||||
include_directories("${d_dir}/include") |
||||
endif() |
||||
endforeach() |
||||
|
||||
file(GLOB lib_srcs "src/*.cpp") |
||||
file(GLOB lib_int_hdrs "src/*.h*") |
||||
file(GLOB lib_cuda "src/cuda/*.cu*") |
||||
file(GLOB lib_cuda_hdrs "src/cuda/*.h*") |
||||
source_group("Src\\Host" FILES ${lib_srcs} ${lib_int_hdrs}) |
||||
source_group("Src\\Cuda" FILES ${lib_cuda} ${lib_cuda_hdrs}) |
||||
|
||||
file(GLOB lib_hdrs "include/opencv2/${name}/*.h*") |
||||
source_group("Include" FILES ${lib_hdrs}) |
||||
|
||||
#file(GLOB lib_device_hdrs "include/opencv2/${name}/device/*.h*") |
||||
file(GLOB lib_device_hdrs "src/opencv2/gpu/device/*.h*") |
||||
source_group("Device" FILES ${lib_device_hdrs}) |
||||
|
||||
if (HAVE_CUDA) |
||||
get_filename_component(_path_to_findnpp "${CMAKE_CURRENT_LIST_FILE}" PATH) |
||||
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${_path_to_findnpp}) |
||||
find_package(NPP 3.2.16 REQUIRED) |
||||
message(STATUS "NPP detected: " ${NPP_VERSION}) |
||||
|
||||
include_directories(${CUDA_INCLUDE_DIRS} ${CUDA_NPP_INCLUDES}) |
||||
|
||||
if (UNIX OR APPLE) |
||||
set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-Xcompiler;-fPIC;") |
||||
#set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}" "-fPIC") |
||||
endif() |
||||
|
||||
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-keep") |
||||
#set (CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-Xcompiler;/EHsc-;") |
||||
|
||||
string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") |
||||
string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") |
||||
string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") |
||||
|
||||
if(MSVC) |
||||
#string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") |
||||
#string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") |
||||
#string(REPLACE "/W4" "/W3" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") |
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4211 /wd4201 /wd4100 /wd4505 /wd4408") |
||||
|
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_C_FLAGS "${CMAKE_C_FLAGS}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") |
||||
string(REPLACE "/EHsc-" "/EHs" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") |
||||
endif() |
||||
|
||||
CUDA_COMPILE(cuda_objs ${lib_cuda}) |
||||
#CUDA_BUILD_CLEAN_TARGET() |
||||
endif() |
||||
|
||||
|
||||
add_library(${the_target} ${lib_srcs} ${lib_hdrs} ${lib_int_hdrs} ${lib_cuda} ${lib_cuda_hdrs} ${lib_device_hdrs} ${cuda_objs}) |
||||
|
||||
|
||||
if(PCHSupport_FOUND) |
||||
set(pch_header ${CMAKE_CURRENT_SOURCE_DIR}/src/precomp.hpp) |
||||
if(${CMAKE_GENERATOR} MATCHES "Visual*" OR ${CMAKE_GENERATOR} MATCHES "Xcode*") |
||||
if(${CMAKE_GENERATOR} MATCHES "Visual*") |
||||
set(${the_target}_pch "src/precomp.cpp") |
||||
endif() |
||||
add_native_precompiled_header(${the_target} ${pch_header}) |
||||
elseif(CMAKE_COMPILER_IS_GNUCXX AND ${CMAKE_GENERATOR} MATCHES ".*Makefiles") |
||||
add_precompiled_header(${the_target} ${pch_header}) |
||||
endif() |
||||
endif() |
||||
|
||||
# For dynamic link numbering convenions |
||||
set_target_properties(${the_target} PROPERTIES |
||||
VERSION ${OPENCV_VERSION} |
||||
SOVERSION ${OPENCV_SOVERSION} |
||||
OUTPUT_NAME "${the_target}${OPENCV_DLLVERSION}" |
||||
) |
||||
|
||||
# Additional target properties |
||||
set_target_properties(${the_target} PROPERTIES |
||||
DEBUG_POSTFIX "${OPENCV_DEBUG_POSTFIX}" |
||||
ARCHIVE_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib/" |
||||
RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/bin/" |
||||
INSTALL_NAME_DIR "${CMAKE_INSTALL_PREFIX}/lib" |
||||
) |
||||
|
||||
# Add the required libraries for linking: |
||||
target_link_libraries(${the_target} ${OPENCV_LINKER_LIBS} ${IPP_LIBS} ${DEPS}) |
||||
|
||||
if (HAVE_CUDA) |
||||
target_link_libraries(${the_target} ${CUDA_LIBRARIES} ${CUDA_NPP_LIBRARIES}) |
||||
endif() |
||||
|
||||
if(MSVC) |
||||
if(CMAKE_CROSSCOMPILING) |
||||
set_target_properties(${the_target} PROPERTIES |
||||
LINK_FLAGS "/NODEFAULTLIB:secchk" |
||||
) |
||||
endif() |
||||
set_target_properties(${the_target} PROPERTIES |
||||
LINK_FLAGS "/NODEFAULTLIB:libc" |
||||
) |
||||
endif() |
||||
|
||||
# Dependencies of this target: |
||||
add_dependencies(${the_target} ${DEPS}) |
||||
|
||||
install(TARGETS ${the_target} |
||||
RUNTIME DESTINATION bin COMPONENT main |
||||
LIBRARY DESTINATION lib COMPONENT main |
||||
ARCHIVE DESTINATION lib COMPONENT main) |
||||
|
||||
install(FILES ${lib_hdrs} |
||||
DESTINATION include/opencv2/${name} |
||||
COMPONENT main) |
||||
|
||||
#install(FILES ${lib_device_hdrs} |
||||
# DESTINATION include/opencv2/${name}/device |
||||
# COMPONENT main) |
||||
|
||||
|
||||
|
||||
@ -1,63 +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*/
|
||||
|
||||
#include "internal_shared.hpp" |
||||
#include "border_interpolate.hpp" |
||||
#include "opencv2/gpu/gpu.hpp" |
||||
|
||||
|
||||
bool cv::gpu::tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType) |
||||
{ |
||||
if (cpuBorderType == cv::BORDER_REFLECT101) |
||||
{ |
||||
gpuBorderType = cv::gpu::BORDER_REFLECT101; |
||||
return true; |
||||
} |
||||
|
||||
if (cpuBorderType == cv::BORDER_REPLICATE) |
||||
{ |
||||
gpuBorderType = cv::gpu::BORDER_REPLICATE; |
||||
return true; |
||||
} |
||||
|
||||
return false; |
||||
} |
||||
@ -1,56 +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 __OPENCV_GPU_BORDER_INTERPOLATE_HPP__ |
||||
#define __OPENCV_GPU_BORDER_INTERPOLATE_HPP__ |
||||
|
||||
#include "border_interpolate.hpp" |
||||
|
||||
namespace cv { namespace gpu { |
||||
|
||||
// Converts CPU border extrapolation mode into GPU internal analogue.
|
||||
// Returns true if the GPU analogue exists, false otherwise.
|
||||
bool tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType); |
||||
|
||||
}} |
||||
|
||||
#endif |
||||
@ -1,179 +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 __OPENCV_GPU_BORDER_INTERPOLATE_HPP__ |
||||
#define __OPENCV_GPU_BORDER_INTERPOLATE_HPP__ |
||||
|
||||
#include "../internal_shared.hpp" |
||||
|
||||
namespace cv { namespace gpu { |
||||
|
||||
struct BrdReflect101
|
||||
{ |
||||
BrdReflect101(int len): last(len - 1) {} |
||||
|
||||
__device__ int idx_low(int i) const |
||||
{ |
||||
return abs(i); |
||||
} |
||||
|
||||
__device__ int idx_high(int i) const
|
||||
{ |
||||
return last - abs(last - i); |
||||
} |
||||
|
||||
__device__ int idx(int i) const |
||||
{ |
||||
return abs(idx_high(i)); |
||||
} |
||||
|
||||
bool is_range_safe(int mini, int maxi) const
|
||||
{ |
||||
return -last <= mini && maxi <= 2 * last; |
||||
} |
||||
|
||||
int last; |
||||
}; |
||||
|
||||
|
||||
template <typename T> |
||||
struct BrdRowReflect101: BrdReflect101 |
||||
{ |
||||
BrdRowReflect101(int len): BrdReflect101(len) {} |
||||
|
||||
__device__ float at_low(int i, const T* data) const
|
||||
{ |
||||
return data[idx_low(i)]; |
||||
} |
||||
|
||||
__device__ float at_high(int i, const T* data) const
|
||||
{ |
||||
return data[idx_high(i)]; |
||||
} |
||||
}; |
||||
|
||||
|
||||
template <typename T> |
||||
struct BrdColReflect101: BrdReflect101 |
||||
{ |
||||
BrdColReflect101(int len, int step): BrdReflect101(len), step(step) {} |
||||
|
||||
__device__ float at_low(int i, const T* data) const
|
||||
{ |
||||
return data[idx_low(i) * step]; |
||||
} |
||||
|
||||
__device__ float at_high(int i, const T* data) const
|
||||
{ |
||||
return data[idx_high(i) * step]; |
||||
} |
||||
|
||||
int step; |
||||
}; |
||||
|
||||
|
||||
struct BrdReplicate |
||||
{ |
||||
BrdReplicate(int len): last(len - 1) {} |
||||
|
||||
__device__ int idx_low(int i) const |
||||
{ |
||||
return max(i, 0); |
||||
} |
||||
|
||||
__device__ int idx_high(int i) const
|
||||
{ |
||||
return min(i, last); |
||||
} |
||||
|
||||
__device__ int idx(int i) const |
||||
{ |
||||
return max(min(i, last), 0); |
||||
} |
||||
|
||||
bool is_range_safe(int mini, int maxi) const
|
||||
{ |
||||
return true; |
||||
} |
||||
|
||||
int last; |
||||
}; |
||||
|
||||
|
||||
template <typename T> |
||||
struct BrdRowReplicate: BrdReplicate |
||||
{ |
||||
BrdRowReplicate(int len): BrdReplicate(len) {} |
||||
|
||||
__device__ float at_low(int i, const T* data) const
|
||||
{ |
||||
return data[idx_low(i)]; |
||||
} |
||||
|
||||
__device__ float at_high(int i, const T* data) const
|
||||
{ |
||||
return data[idx_high(i)]; |
||||
} |
||||
}; |
||||
|
||||
|
||||
template <typename T> |
||||
struct BrdColReplicate: BrdReplicate |
||||
{ |
||||
BrdColReplicate(int len, int step): BrdReplicate(len), step(step) {} |
||||
|
||||
__device__ float at_low(int i, const T* data) const
|
||||
{ |
||||
return data[idx_low(i) * step]; |
||||
} |
||||
|
||||
__device__ float at_high(int i, const T* data) const
|
||||
{ |
||||
return data[idx_high(i) * step]; |
||||
} |
||||
|
||||
int step; |
||||
}; |
||||
|
||||
}} |
||||
|
||||
#endif |
||||
@ -1,169 +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 __OPENCV_GPU_SATURATE_CAST_HPP__ |
||||
#define __OPENCV_GPU_SATURATE_CAST_HPP__ |
||||
|
||||
#include "cuda_shared.hpp" |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace gpu |
||||
{ |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(uchar v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(schar v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(ushort v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(short v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(uint v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(int v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(float v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(double v) { return _Tp(v); } |
||||
|
||||
template<> static __device__ uchar saturate_cast<uchar>(schar v) |
||||
{ return (uchar)max((int)v, 0); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(ushort v) |
||||
{ return (uchar)min((uint)v, (uint)UCHAR_MAX); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(int v) |
||||
{ return (uchar)((uint)v <= UCHAR_MAX ? v : v > 0 ? UCHAR_MAX : 0); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(uint v) |
||||
{ return (uchar)min(v, (uint)UCHAR_MAX); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(short v) |
||||
{ return saturate_cast<uchar>((uint)v); } |
||||
|
||||
template<> static __device__ uchar saturate_cast<uchar>(float v) |
||||
{ int iv = __float2int_rn(v); return saturate_cast<uchar>(iv); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(double v) |
||||
{ |
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
int iv = __double2int_rn(v); return saturate_cast<uchar>(iv); |
||||
#else |
||||
return saturate_cast<uchar>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ schar saturate_cast<schar>(uchar v) |
||||
{ return (schar)min((int)v, SCHAR_MAX); } |
||||
template<> static __device__ schar saturate_cast<schar>(ushort v) |
||||
{ return (schar)min((uint)v, (uint)SCHAR_MAX); } |
||||
template<> static __device__ schar saturate_cast<schar>(int v) |
||||
{ |
||||
return (schar)((uint)(v-SCHAR_MIN) <= (uint)UCHAR_MAX ? |
||||
v : v > 0 ? SCHAR_MAX : SCHAR_MIN); |
||||
} |
||||
template<> static __device__ schar saturate_cast<schar>(short v) |
||||
{ return saturate_cast<schar>((int)v); } |
||||
template<> static __device__ schar saturate_cast<schar>(uint v) |
||||
{ return (schar)min(v, (uint)SCHAR_MAX); } |
||||
|
||||
template<> static __device__ schar saturate_cast<schar>(float v) |
||||
{ int iv = __float2int_rn(v); return saturate_cast<schar>(iv); } |
||||
template<> static __device__ schar saturate_cast<schar>(double v) |
||||
{
|
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
int iv = __double2int_rn(v); return saturate_cast<schar>(iv); |
||||
#else |
||||
return saturate_cast<schar>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ ushort saturate_cast<ushort>(schar v) |
||||
{ return (ushort)max((int)v, 0); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(short v) |
||||
{ return (ushort)max((int)v, 0); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(int v) |
||||
{ return (ushort)((uint)v <= (uint)USHRT_MAX ? v : v > 0 ? USHRT_MAX : 0); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(uint v) |
||||
{ return (ushort)min(v, (uint)USHRT_MAX); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(float v) |
||||
{ int iv = __float2int_rn(v); return saturate_cast<ushort>(iv); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(double v) |
||||
{
|
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
int iv = __double2int_rn(v); return saturate_cast<ushort>(iv); |
||||
#else |
||||
return saturate_cast<ushort>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ short saturate_cast<short>(ushort v) |
||||
{ return (short)min((int)v, SHRT_MAX); } |
||||
template<> static __device__ short saturate_cast<short>(int v) |
||||
{ |
||||
return (short)((uint)(v - SHRT_MIN) <= (uint)USHRT_MAX ? |
||||
v : v > 0 ? SHRT_MAX : SHRT_MIN); |
||||
} |
||||
template<> static __device__ short saturate_cast<short>(uint v) |
||||
{ return (short)min(v, (uint)SHRT_MAX); } |
||||
template<> static __device__ short saturate_cast<short>(float v) |
||||
{ int iv = __float2int_rn(v); return saturate_cast<short>(iv); } |
||||
template<> static __device__ short saturate_cast<short>(double v) |
||||
{
|
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
int iv = __double2int_rn(v); return saturate_cast<short>(iv); |
||||
#else |
||||
return saturate_cast<short>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ int saturate_cast<int>(float v) { return __float2int_rn(v); } |
||||
template<> static __device__ int saturate_cast<int>(double v)
|
||||
{ |
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
return __double2int_rn(v); |
||||
#else |
||||
return saturate_cast<int>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ uint saturate_cast<uint>(float v){ return __float2uint_rn(v); } |
||||
template<> static __device__ uint saturate_cast<uint>(double v)
|
||||
{
|
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
return __double2uint_rn(v); |
||||
#else |
||||
return saturate_cast<uint>((float)v); |
||||
#endif |
||||
} |
||||
} |
||||
} |
||||
|
||||
#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */ |
||||
@ -1,936 +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 __OPENCV_GPU_VECMATH_HPP__ |
||||
#define __OPENCV_GPU_VECMATH_HPP__ |
||||
|
||||
#include "cuda_shared.hpp" |
||||
#include "saturate_cast.hpp" |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace gpu |
||||
{ |
||||
template<typename T, int N> struct TypeVec; |
||||
|
||||
template<> struct TypeVec<uchar, 1> { typedef uchar vec_t; }; |
||||
template<> struct TypeVec<uchar1, 1> { typedef uchar1 vec_t; }; |
||||
template<> struct TypeVec<uchar, 2> { typedef uchar2 vec_t; }; |
||||
template<> struct TypeVec<uchar2, 2> { typedef uchar2 vec_t; }; |
||||
template<> struct TypeVec<uchar, 3> { typedef uchar3 vec_t; }; |
||||
template<> struct TypeVec<uchar3, 3> { typedef uchar3 vec_t; }; |
||||
template<> struct TypeVec<uchar, 4> { typedef uchar4 vec_t; }; |
||||
template<> struct TypeVec<uchar4, 4> { typedef uchar4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<char, 1> { typedef char vec_t; }; |
||||
template<> struct TypeVec<char1, 1> { typedef char1 vec_t; }; |
||||
template<> struct TypeVec<char, 2> { typedef char2 vec_t; }; |
||||
template<> struct TypeVec<char2, 2> { typedef char2 vec_t; }; |
||||
template<> struct TypeVec<char, 3> { typedef char3 vec_t; }; |
||||
template<> struct TypeVec<char3, 3> { typedef char3 vec_t; }; |
||||
template<> struct TypeVec<char, 4> { typedef char4 vec_t; }; |
||||
template<> struct TypeVec<char4, 4> { typedef char4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<ushort, 1> { typedef ushort vec_t; }; |
||||
template<> struct TypeVec<ushort1, 1> { typedef ushort1 vec_t; }; |
||||
template<> struct TypeVec<ushort, 2> { typedef ushort2 vec_t; }; |
||||
template<> struct TypeVec<ushort2, 2> { typedef ushort2 vec_t; }; |
||||
template<> struct TypeVec<ushort, 3> { typedef ushort3 vec_t; }; |
||||
template<> struct TypeVec<ushort3, 3> { typedef ushort3 vec_t; }; |
||||
template<> struct TypeVec<ushort, 4> { typedef ushort4 vec_t; }; |
||||
template<> struct TypeVec<ushort4, 4> { typedef ushort4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<short, 1> { typedef short vec_t; }; |
||||
template<> struct TypeVec<short1, 1> { typedef short1 vec_t; }; |
||||
template<> struct TypeVec<short, 2> { typedef short2 vec_t; }; |
||||
template<> struct TypeVec<short2, 2> { typedef short2 vec_t; }; |
||||
template<> struct TypeVec<short, 3> { typedef short3 vec_t; }; |
||||
template<> struct TypeVec<short3, 3> { typedef short3 vec_t; }; |
||||
template<> struct TypeVec<short, 4> { typedef short4 vec_t; }; |
||||
template<> struct TypeVec<short4, 4> { typedef short4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<uint, 1> { typedef uint vec_t; }; |
||||
template<> struct TypeVec<uint1, 1> { typedef uint1 vec_t; }; |
||||
template<> struct TypeVec<uint, 2> { typedef uint2 vec_t; }; |
||||
template<> struct TypeVec<uint2, 2> { typedef uint2 vec_t; }; |
||||
template<> struct TypeVec<uint, 3> { typedef uint3 vec_t; }; |
||||
template<> struct TypeVec<uint3, 3> { typedef uint3 vec_t; }; |
||||
template<> struct TypeVec<uint, 4> { typedef uint4 vec_t; }; |
||||
template<> struct TypeVec<uint4, 4> { typedef uint4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<int, 1> { typedef int vec_t; }; |
||||
template<> struct TypeVec<int1, 1> { typedef int1 vec_t; }; |
||||
template<> struct TypeVec<int, 2> { typedef int2 vec_t; }; |
||||
template<> struct TypeVec<int2, 2> { typedef int2 vec_t; }; |
||||
template<> struct TypeVec<int, 3> { typedef int3 vec_t; }; |
||||
template<> struct TypeVec<int3, 3> { typedef int3 vec_t; }; |
||||
template<> struct TypeVec<int, 4> { typedef int4 vec_t; }; |
||||
template<> struct TypeVec<int4, 4> { typedef int4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<float, 1> { typedef float vec_t; }; |
||||
template<> struct TypeVec<float1, 1> { typedef float1 vec_t; }; |
||||
template<> struct TypeVec<float, 2> { typedef float2 vec_t; }; |
||||
template<> struct TypeVec<float2, 2> { typedef float2 vec_t; }; |
||||
template<> struct TypeVec<float, 3> { typedef float3 vec_t; }; |
||||
template<> struct TypeVec<float3, 3> { typedef float3 vec_t; }; |
||||
template<> struct TypeVec<float, 4> { typedef float4 vec_t; }; |
||||
template<> struct TypeVec<float4, 4> { typedef float4 vec_t; }; |
||||
|
||||
template<typename T> struct VecTraits; |
||||
|
||||
template<> struct VecTraits<uchar>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=1}; |
||||
static __device__ uchar all(uchar v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<uchar1>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=1}; |
||||
static __device__ uchar1 all(uchar v) {return make_uchar1(v);} |
||||
}; |
||||
template<> struct VecTraits<uchar2>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ uchar2 all(uchar v) {return make_uchar2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<uchar3>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ uchar3 all(uchar v) {return make_uchar3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<uchar4>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ uchar4 all(uchar v) {return make_uchar4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<char>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ char all(char v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<char1>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ char1 all(char v) {return make_char1(v);} |
||||
}; |
||||
template<> struct VecTraits<char2>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ char2 all(char v) {return make_char2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<char3>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ char3 all(char v) {return make_char3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<char4>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ char4 all(char v) {return make_char4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<ushort>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ ushort all(ushort v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<ushort1>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ ushort1 all(ushort v) {return make_ushort1(v);} |
||||
}; |
||||
template<> struct VecTraits<ushort2>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ ushort2 all(ushort v) {return make_ushort2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<ushort3>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ ushort3 all(ushort v) {return make_ushort3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<ushort4>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ ushort4 all(ushort v) {return make_ushort4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<short>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ short all(short v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<short1>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ short1 all(short v) {return make_short1(v);} |
||||
}; |
||||
template<> struct VecTraits<short2>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ short2 all(short v) {return make_short2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<short3>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ short3 all(short v) {return make_short3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<short4>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ short4 all(short v) {return make_short4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<uint>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ uint all(uint v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<uint1>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ uint1 all(uint v) {return make_uint1(v);} |
||||
}; |
||||
template<> struct VecTraits<uint2>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ uint2 all(uint v) {return make_uint2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<uint3>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ uint3 all(uint v) {return make_uint3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<uint4>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ uint4 all(uint v) {return make_uint4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<int>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ int all(int v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<int1>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ int1 all(int v) {return make_int1(v);} |
||||
}; |
||||
template<> struct VecTraits<int2>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ int2 all(int v) {return make_int2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<int3>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ int3 all(int v) {return make_int3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<int4>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ int4 all(int v) {return make_int4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<float>
|
||||
{
|
||||
typedef float elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ float all(float v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<float1>
|
||||
{
|
||||
typedef float elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ float1 all(float v) {return make_float1(v);} |
||||
}; |
||||
template<> struct VecTraits<float2>
|
||||
{
|
||||
typedef float elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ float2 all(float v) {return make_float2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<float3>
|
||||
{
|
||||
typedef float elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ float3 all(float v) {return make_float3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<float4>
|
||||
{
|
||||
typedef float elem_t; |
||||
enum {cn=4};
|
||||
static __device__ float4 all(float v) {return make_float4(v, v, v, v);} |
||||
}; |
||||
|
||||
template <int cn, typename VecD> struct SatCast; |
||||
template <typename VecD> struct SatCast<1, VecD> |
||||
{ |
||||
template <typename VecS> |
||||
__device__ VecD operator()(const VecS& v) |
||||
{ |
||||
VecD res;
|
||||
res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x); |
||||
return res; |
||||
} |
||||
}; |
||||
template <typename VecD> struct SatCast<2, VecD> |
||||
{ |
||||
template <typename VecS> |
||||
__device__ VecD operator()(const VecS& v) |
||||
{ |
||||
VecD res;
|
||||
res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y); |
||||
return res; |
||||
} |
||||
}; |
||||
template <typename VecD> struct SatCast<3, VecD> |
||||
{ |
||||
template <typename VecS> |
||||
__device__ VecD operator()(const VecS& v) |
||||
{ |
||||
VecD res;
|
||||
res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.z); |
||||
return res; |
||||
} |
||||
}; |
||||
template <typename VecD> struct SatCast<4, VecD> |
||||
{ |
||||
template <typename VecS> |
||||
__device__ VecD operator()(const VecS& v) |
||||
{ |
||||
VecD res;
|
||||
res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.z); |
||||
res.w = saturate_cast< VecTraits<VecD>::elem_t >(v.w); |
||||
return res; |
||||
} |
||||
}; |
||||
|
||||
template <typename VecD, typename VecS> static __device__ VecD saturate_cast_caller(const VecS& v) |
||||
{ |
||||
SatCast< |
||||
|
||||
VecTraits<VecD>::cn,
|
||||
|
||||
VecD |
||||
>
|
||||
|
||||
cast; |
||||
return cast(v); |
||||
} |
||||
|
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uchar1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const char1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const ushort1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const short1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uint1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const int1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const float1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
|
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uchar2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const char2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const ushort2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const short2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uint2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const int2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const float2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
|
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uchar3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const char3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const ushort3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const short3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uint3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const int3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const float3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
|
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uchar4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const char4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const ushort4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const short4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uint4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const int4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const float4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
|
||||
static __device__ uchar1 operator+(const uchar1& a, const uchar1& b) |
||||
{ |
||||
return make_uchar1(a.x + b.x); |
||||
} |
||||
static __device__ uchar1 operator-(const uchar1& a, const uchar1& b) |
||||
{ |
||||
return make_uchar1(a.x - b.x); |
||||
} |
||||
static __device__ uchar1 operator*(const uchar1& a, const uchar1& b) |
||||
{ |
||||
return make_uchar1(a.x * b.x); |
||||
} |
||||
static __device__ uchar1 operator/(const uchar1& a, const uchar1& b) |
||||
{ |
||||
return make_uchar1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const uchar1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ uchar2 operator+(const uchar2& a, const uchar2& b) |
||||
{ |
||||
return make_uchar2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ uchar2 operator-(const uchar2& a, const uchar2& b) |
||||
{ |
||||
return make_uchar2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ uchar2 operator*(const uchar2& a, const uchar2& b) |
||||
{ |
||||
return make_uchar2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ uchar2 operator/(const uchar2& a, const uchar2& b) |
||||
{ |
||||
return make_uchar2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const uchar2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ uchar3 operator+(const uchar3& a, const uchar3& b) |
||||
{ |
||||
return make_uchar3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ uchar3 operator-(const uchar3& a, const uchar3& b) |
||||
{ |
||||
return make_uchar3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ uchar3 operator*(const uchar3& a, const uchar3& b) |
||||
{ |
||||
return make_uchar3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ uchar3 operator/(const uchar3& a, const uchar3& b) |
||||
{ |
||||
return make_uchar3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const uchar3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ uchar4 operator+(const uchar4& a, const uchar4& b) |
||||
{ |
||||
return make_uchar4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ uchar4 operator-(const uchar4& a, const uchar4& b) |
||||
{ |
||||
return make_uchar4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ uchar4 operator*(const uchar4& a, const uchar4& b) |
||||
{ |
||||
return make_uchar4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ uchar4 operator/(const uchar4& a, const uchar4& b) |
||||
{ |
||||
return make_uchar4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const uchar4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ char1 operator+(const char1& a, const char1& b) |
||||
{ |
||||
return make_char1(a.x + b.x); |
||||
} |
||||
static __device__ char1 operator-(const char1& a, const char1& b) |
||||
{ |
||||
return make_char1(a.x - b.x); |
||||
} |
||||
static __device__ char1 operator*(const char1& a, const char1& b) |
||||
{ |
||||
return make_char1(a.x * b.x); |
||||
} |
||||
static __device__ char1 operator/(const char1& a, const char1& b) |
||||
{ |
||||
return make_char1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const char1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ char2 operator+(const char2& a, const char2& b) |
||||
{ |
||||
return make_char2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ char2 operator-(const char2& a, const char2& b) |
||||
{ |
||||
return make_char2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ char2 operator*(const char2& a, const char2& b) |
||||
{ |
||||
return make_char2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ char2 operator/(const char2& a, const char2& b) |
||||
{ |
||||
return make_char2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const char2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ char3 operator+(const char3& a, const char3& b) |
||||
{ |
||||
return make_char3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ char3 operator-(const char3& a, const char3& b) |
||||
{ |
||||
return make_char3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ char3 operator*(const char3& a, const char3& b) |
||||
{ |
||||
return make_char3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ char3 operator/(const char3& a, const char3& b) |
||||
{ |
||||
return make_char3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const char3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ char4 operator+(const char4& a, const char4& b) |
||||
{ |
||||
return make_char4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ char4 operator-(const char4& a, const char4& b) |
||||
{ |
||||
return make_char4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ char4 operator*(const char4& a, const char4& b) |
||||
{ |
||||
return make_char4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ char4 operator/(const char4& a, const char4& b) |
||||
{ |
||||
return make_char4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const char4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ ushort1 operator+(const ushort1& a, const ushort1& b) |
||||
{ |
||||
return make_ushort1(a.x + b.x); |
||||
} |
||||
static __device__ ushort1 operator-(const ushort1& a, const ushort1& b) |
||||
{ |
||||
return make_ushort1(a.x - b.x); |
||||
} |
||||
static __device__ ushort1 operator*(const ushort1& a, const ushort1& b) |
||||
{ |
||||
return make_ushort1(a.x * b.x); |
||||
} |
||||
static __device__ ushort1 operator/(const ushort1& a, const ushort1& b) |
||||
{ |
||||
return make_ushort1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const ushort1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ ushort2 operator+(const ushort2& a, const ushort2& b) |
||||
{ |
||||
return make_ushort2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ ushort2 operator-(const ushort2& a, const ushort2& b) |
||||
{ |
||||
return make_ushort2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ ushort2 operator*(const ushort2& a, const ushort2& b) |
||||
{ |
||||
return make_ushort2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ ushort2 operator/(const ushort2& a, const ushort2& b) |
||||
{ |
||||
return make_ushort2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const ushort2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ ushort3 operator+(const ushort3& a, const ushort3& b) |
||||
{ |
||||
return make_ushort3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ ushort3 operator-(const ushort3& a, const ushort3& b) |
||||
{ |
||||
return make_ushort3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ ushort3 operator*(const ushort3& a, const ushort3& b) |
||||
{ |
||||
return make_ushort3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ ushort3 operator/(const ushort3& a, const ushort3& b) |
||||
{ |
||||
return make_ushort3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const ushort3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ ushort4 operator+(const ushort4& a, const ushort4& b) |
||||
{ |
||||
return make_ushort4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ ushort4 operator-(const ushort4& a, const ushort4& b) |
||||
{ |
||||
return make_ushort4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ ushort4 operator*(const ushort4& a, const ushort4& b) |
||||
{ |
||||
return make_ushort4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ ushort4 operator/(const ushort4& a, const ushort4& b) |
||||
{ |
||||
return make_ushort4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const ushort4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ short1 operator+(const short1& a, const short1& b) |
||||
{ |
||||
return make_short1(a.x + b.x); |
||||
} |
||||
static __device__ short1 operator-(const short1& a, const short1& b) |
||||
{ |
||||
return make_short1(a.x - b.x); |
||||
} |
||||
static __device__ short1 operator*(const short1& a, const short1& b) |
||||
{ |
||||
return make_short1(a.x * b.x); |
||||
} |
||||
static __device__ short1 operator/(const short1& a, const short1& b) |
||||
{ |
||||
return make_short1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const short1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ short2 operator+(const short2& a, const short2& b) |
||||
{ |
||||
return make_short2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ short2 operator-(const short2& a, const short2& b) |
||||
{ |
||||
return make_short2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ short2 operator*(const short2& a, const short2& b) |
||||
{ |
||||
return make_short2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ short2 operator/(const short2& a, const short2& b) |
||||
{ |
||||
return make_short2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const short2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ short3 operator+(const short3& a, const short3& b) |
||||
{ |
||||
return make_short3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ short3 operator-(const short3& a, const short3& b) |
||||
{ |
||||
return make_short3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ short3 operator*(const short3& a, const short3& b) |
||||
{ |
||||
return make_short3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ short3 operator/(const short3& a, const short3& b) |
||||
{ |
||||
return make_short3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const short3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ short4 operator+(const short4& a, const short4& b) |
||||
{ |
||||
return make_short4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ short4 operator-(const short4& a, const short4& b) |
||||
{ |
||||
return make_short4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ short4 operator*(const short4& a, const short4& b) |
||||
{ |
||||
return make_short4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ short4 operator/(const short4& a, const short4& b) |
||||
{ |
||||
return make_short4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const short4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ int1 operator+(const int1& a, const int1& b) |
||||
{ |
||||
return make_int1(a.x + b.x); |
||||
} |
||||
static __device__ int1 operator-(const int1& a, const int1& b) |
||||
{ |
||||
return make_int1(a.x - b.x); |
||||
} |
||||
static __device__ int1 operator*(const int1& a, const int1& b) |
||||
{ |
||||
return make_int1(a.x * b.x); |
||||
} |
||||
static __device__ int1 operator/(const int1& a, const int1& b) |
||||
{ |
||||
return make_int1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const int1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ int2 operator+(const int2& a, const int2& b) |
||||
{ |
||||
return make_int2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ int2 operator-(const int2& a, const int2& b) |
||||
{ |
||||
return make_int2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ int2 operator*(const int2& a, const int2& b) |
||||
{ |
||||
return make_int2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ int2 operator/(const int2& a, const int2& b) |
||||
{ |
||||
return make_int2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const int2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ int3 operator+(const int3& a, const int3& b) |
||||
{ |
||||
return make_int3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ int3 operator-(const int3& a, const int3& b) |
||||
{ |
||||
return make_int3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ int3 operator*(const int3& a, const int3& b) |
||||
{ |
||||
return make_int3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ int3 operator/(const int3& a, const int3& b) |
||||
{ |
||||
return make_int3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const int3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ int4 operator+(const int4& a, const int4& b) |
||||
{ |
||||
return make_int4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ int4 operator-(const int4& a, const int4& b) |
||||
{ |
||||
return make_int4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ int4 operator*(const int4& a, const int4& b) |
||||
{ |
||||
return make_int4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ int4 operator/(const int4& a, const int4& b) |
||||
{ |
||||
return make_int4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const int4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ float1 operator+(const float1& a, const float1& b) |
||||
{ |
||||
return make_float1(a.x + b.x); |
||||
} |
||||
static __device__ float1 operator-(const float1& a, const float1& b) |
||||
{ |
||||
return make_float1(a.x - b.x); |
||||
} |
||||
static __device__ float1 operator*(const float1& a, const float1& b) |
||||
{ |
||||
return make_float1(a.x * b.x); |
||||
} |
||||
static __device__ float1 operator/(const float1& a, const float1& b) |
||||
{ |
||||
return make_float1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const float1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ float2 operator+(const float2& a, const float2& b) |
||||
{ |
||||
return make_float2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ float2 operator-(const float2& a, const float2& b) |
||||
{ |
||||
return make_float2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ float2 operator*(const float2& a, const float2& b) |
||||
{ |
||||
return make_float2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ float2 operator/(const float2& a, const float2& b) |
||||
{ |
||||
return make_float2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const float2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ float3 operator+(const float3& a, const float3& b) |
||||
{ |
||||
return make_float3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ float3 operator-(const float3& a, const float3& b) |
||||
{ |
||||
return make_float3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ float3 operator*(const float3& a, const float3& b) |
||||
{ |
||||
return make_float3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ float3 operator/(const float3& a, const float3& b) |
||||
{ |
||||
return make_float3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const float3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ float4 operator+(const float4& a, const float4& b) |
||||
{ |
||||
return make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ float4 operator-(const float4& a, const float4& b) |
||||
{ |
||||
return make_float4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ float4 operator*(const float4& a, const float4& b) |
||||
{ |
||||
return make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ float4 operator/(const float4& a, const float4& b) |
||||
{ |
||||
return make_float4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const float4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
} |
||||
} |
||||
|
||||
#endif // __OPENCV_GPU_VECMATH_HPP__
|
||||
@ -1,57 +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 __OPENCV_GPU_INTERNAL_SHARED_HPP__ |
||||
#define __OPENCV_GPU_INTERNAL_SHARED_HPP__ |
||||
|
||||
namespace cv { namespace gpu { |
||||
|
||||
// Internal GPU anlagues of CPU border extrapolation types
|
||||
enum
|
||||
{ |
||||
BORDER_REFLECT101 = 0, |
||||
BORDER_REPLICATE |
||||
}; |
||||
|
||||
}} |
||||
|
||||
#endif |
||||
@ -0,0 +1,176 @@ |
||||
/*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 bpied warranties, including, but not limited to, the bpied
|
||||
// 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*/
|
||||
|
||||
namespace cv
|
||||
{
|
||||
namespace gpu
|
||||
{ |
||||
namespace device |
||||
{ |
||||
struct BrdReflect101
|
||||
{ |
||||
BrdReflect101(int len): last(len - 1) {} |
||||
|
||||
__device__ int idx_low(int i) const |
||||
{ |
||||
return abs(i); |
||||
} |
||||
|
||||
__device__ int idx_high(int i) const
|
||||
{ |
||||
return last - abs(last - i); |
||||
} |
||||
|
||||
__device__ int idx(int i) const |
||||
{ |
||||
return abs(idx_high(i)); |
||||
} |
||||
|
||||
bool is_range_safe(int mini, int maxi) const
|
||||
{ |
||||
return -last <= mini && maxi <= 2 * last; |
||||
} |
||||
|
||||
int last; |
||||
}; |
||||
|
||||
|
||||
template <typename T> |
||||
struct BrdRowReflect101: BrdReflect101 |
||||
{ |
||||
BrdRowReflect101(int len): BrdReflect101(len) {} |
||||
|
||||
__device__ float at_low(int i, const T* data) const
|
||||
{ |
||||
return data[idx_low(i)]; |
||||
} |
||||
|
||||
__device__ float at_high(int i, const T* data) const
|
||||
{ |
||||
return data[idx_high(i)]; |
||||
} |
||||
}; |
||||
|
||||
|
||||
template <typename T> |
||||
struct BrdColReflect101: BrdReflect101 |
||||
{ |
||||
BrdColReflect101(int len, int step): BrdReflect101(len), step(step) {} |
||||
|
||||
__device__ float at_low(int i, const T* data) const
|
||||
{ |
||||
return data[idx_low(i) * step]; |
||||
} |
||||
|
||||
__device__ float at_high(int i, const T* data) const
|
||||
{ |
||||
return data[idx_high(i) * step]; |
||||
} |
||||
|
||||
int step; |
||||
}; |
||||
|
||||
|
||||
struct BrdReplicate |
||||
{ |
||||
BrdReplicate(int len): last(len - 1) {} |
||||
|
||||
__device__ int idx_low(int i) const |
||||
{ |
||||
return max(i, 0); |
||||
} |
||||
|
||||
__device__ int idx_high(int i) const
|
||||
{ |
||||
return min(i, last); |
||||
} |
||||
|
||||
__device__ int idx(int i) const |
||||
{ |
||||
return max(min(i, last), 0); |
||||
} |
||||
|
||||
bool is_range_safe(int mini, int maxi) const
|
||||
{ |
||||
return true; |
||||
} |
||||
|
||||
int last; |
||||
}; |
||||
|
||||
|
||||
template <typename T> |
||||
struct BrdRowReplicate: BrdReplicate |
||||
{ |
||||
BrdRowReplicate(int len): BrdReplicate(len) {} |
||||
|
||||
__device__ float at_low(int i, const T* data) const
|
||||
{ |
||||
return data[idx_low(i)]; |
||||
} |
||||
|
||||
__device__ float at_high(int i, const T* data) const
|
||||
{ |
||||
return data[idx_high(i)]; |
||||
} |
||||
}; |
||||
|
||||
|
||||
template <typename T> |
||||
struct BrdColReplicate: BrdReplicate |
||||
{ |
||||
BrdColReplicate(int len, int step): BrdReplicate(len), step(step) {} |
||||
|
||||
__device__ float at_low(int i, const T* data) const
|
||||
{ |
||||
return data[idx_low(i) * step]; |
||||
} |
||||
|
||||
__device__ float at_high(int i, const T* data) const
|
||||
{ |
||||
return data[idx_high(i) * step]; |
||||
} |
||||
int step; |
||||
}; |
||||
} |
||||
} |
||||
} |
||||
@ -0,0 +1,172 @@ |
||||
/*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 __OPENCV_GPU_SATURATE_CAST_HPP__ |
||||
#define __OPENCV_GPU_SATURATE_CAST_HPP__ |
||||
|
||||
#include "internal_shared.hpp" |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace gpu |
||||
{ |
||||
namespace device |
||||
{ |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(uchar v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(schar v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(ushort v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(short v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(uint v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(int v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(float v) { return _Tp(v); } |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(double v) { return _Tp(v); } |
||||
|
||||
template<> static __device__ uchar saturate_cast<uchar>(schar v) |
||||
{ return (uchar)max((int)v, 0); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(ushort v) |
||||
{ return (uchar)min((uint)v, (uint)UCHAR_MAX); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(int v) |
||||
{ return (uchar)((uint)v <= UCHAR_MAX ? v : v > 0 ? UCHAR_MAX : 0); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(uint v) |
||||
{ return (uchar)min(v, (uint)UCHAR_MAX); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(short v) |
||||
{ return saturate_cast<uchar>((uint)v); } |
||||
|
||||
template<> static __device__ uchar saturate_cast<uchar>(float v) |
||||
{ int iv = __float2int_rn(v); return saturate_cast<uchar>(iv); } |
||||
template<> static __device__ uchar saturate_cast<uchar>(double v) |
||||
{ |
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
int iv = __double2int_rn(v); return saturate_cast<uchar>(iv); |
||||
#else |
||||
return saturate_cast<uchar>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ schar saturate_cast<schar>(uchar v) |
||||
{ return (schar)min((int)v, SCHAR_MAX); } |
||||
template<> static __device__ schar saturate_cast<schar>(ushort v) |
||||
{ return (schar)min((uint)v, (uint)SCHAR_MAX); } |
||||
template<> static __device__ schar saturate_cast<schar>(int v) |
||||
{ |
||||
return (schar)((uint)(v-SCHAR_MIN) <= (uint)UCHAR_MAX ? |
||||
v : v > 0 ? SCHAR_MAX : SCHAR_MIN); |
||||
} |
||||
template<> static __device__ schar saturate_cast<schar>(short v) |
||||
{ return saturate_cast<schar>((int)v); } |
||||
template<> static __device__ schar saturate_cast<schar>(uint v) |
||||
{ return (schar)min(v, (uint)SCHAR_MAX); } |
||||
|
||||
template<> static __device__ schar saturate_cast<schar>(float v) |
||||
{ int iv = __float2int_rn(v); return saturate_cast<schar>(iv); } |
||||
template<> static __device__ schar saturate_cast<schar>(double v) |
||||
{
|
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
int iv = __double2int_rn(v); return saturate_cast<schar>(iv); |
||||
#else |
||||
return saturate_cast<schar>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ ushort saturate_cast<ushort>(schar v) |
||||
{ return (ushort)max((int)v, 0); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(short v) |
||||
{ return (ushort)max((int)v, 0); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(int v) |
||||
{ return (ushort)((uint)v <= (uint)USHRT_MAX ? v : v > 0 ? USHRT_MAX : 0); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(uint v) |
||||
{ return (ushort)min(v, (uint)USHRT_MAX); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(float v) |
||||
{ int iv = __float2int_rn(v); return saturate_cast<ushort>(iv); } |
||||
template<> static __device__ ushort saturate_cast<ushort>(double v) |
||||
{
|
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
int iv = __double2int_rn(v); return saturate_cast<ushort>(iv); |
||||
#else |
||||
return saturate_cast<ushort>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ short saturate_cast<short>(ushort v) |
||||
{ return (short)min((int)v, SHRT_MAX); } |
||||
template<> static __device__ short saturate_cast<short>(int v) |
||||
{ |
||||
return (short)((uint)(v - SHRT_MIN) <= (uint)USHRT_MAX ? |
||||
v : v > 0 ? SHRT_MAX : SHRT_MIN); |
||||
} |
||||
template<> static __device__ short saturate_cast<short>(uint v) |
||||
{ return (short)min(v, (uint)SHRT_MAX); } |
||||
template<> static __device__ short saturate_cast<short>(float v) |
||||
{ int iv = __float2int_rn(v); return saturate_cast<short>(iv); } |
||||
template<> static __device__ short saturate_cast<short>(double v) |
||||
{
|
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
int iv = __double2int_rn(v); return saturate_cast<short>(iv); |
||||
#else |
||||
return saturate_cast<short>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ int saturate_cast<int>(float v) { return __float2int_rn(v); } |
||||
template<> static __device__ int saturate_cast<int>(double v)
|
||||
{ |
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
return __double2int_rn(v); |
||||
#else |
||||
return saturate_cast<int>((float)v); |
||||
#endif |
||||
} |
||||
|
||||
template<> static __device__ uint saturate_cast<uint>(float v){ return __float2uint_rn(v); } |
||||
template<> static __device__ uint saturate_cast<uint>(double v)
|
||||
{
|
||||
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130 |
||||
return __double2uint_rn(v); |
||||
#else |
||||
return saturate_cast<uint>((float)v); |
||||
#endif |
||||
} |
||||
} |
||||
} |
||||
} |
||||
|
||||
#endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */ |
||||
@ -0,0 +1,939 @@ |
||||
/*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 __OPENCV_GPU_VECMATH_HPP__ |
||||
#define __OPENCV_GPU_VECMATH_HPP__ |
||||
|
||||
#include "internal_shared.hpp" |
||||
#include "saturate_cast.hpp" |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace gpu |
||||
{ |
||||
namespace device |
||||
{ |
||||
template<typename T, int N> struct TypeVec; |
||||
|
||||
template<> struct TypeVec<uchar, 1> { typedef uchar vec_t; }; |
||||
template<> struct TypeVec<uchar1, 1> { typedef uchar1 vec_t; }; |
||||
template<> struct TypeVec<uchar, 2> { typedef uchar2 vec_t; }; |
||||
template<> struct TypeVec<uchar2, 2> { typedef uchar2 vec_t; }; |
||||
template<> struct TypeVec<uchar, 3> { typedef uchar3 vec_t; }; |
||||
template<> struct TypeVec<uchar3, 3> { typedef uchar3 vec_t; }; |
||||
template<> struct TypeVec<uchar, 4> { typedef uchar4 vec_t; }; |
||||
template<> struct TypeVec<uchar4, 4> { typedef uchar4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<char, 1> { typedef char vec_t; }; |
||||
template<> struct TypeVec<char1, 1> { typedef char1 vec_t; }; |
||||
template<> struct TypeVec<char, 2> { typedef char2 vec_t; }; |
||||
template<> struct TypeVec<char2, 2> { typedef char2 vec_t; }; |
||||
template<> struct TypeVec<char, 3> { typedef char3 vec_t; }; |
||||
template<> struct TypeVec<char3, 3> { typedef char3 vec_t; }; |
||||
template<> struct TypeVec<char, 4> { typedef char4 vec_t; }; |
||||
template<> struct TypeVec<char4, 4> { typedef char4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<ushort, 1> { typedef ushort vec_t; }; |
||||
template<> struct TypeVec<ushort1, 1> { typedef ushort1 vec_t; }; |
||||
template<> struct TypeVec<ushort, 2> { typedef ushort2 vec_t; }; |
||||
template<> struct TypeVec<ushort2, 2> { typedef ushort2 vec_t; }; |
||||
template<> struct TypeVec<ushort, 3> { typedef ushort3 vec_t; }; |
||||
template<> struct TypeVec<ushort3, 3> { typedef ushort3 vec_t; }; |
||||
template<> struct TypeVec<ushort, 4> { typedef ushort4 vec_t; }; |
||||
template<> struct TypeVec<ushort4, 4> { typedef ushort4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<short, 1> { typedef short vec_t; }; |
||||
template<> struct TypeVec<short1, 1> { typedef short1 vec_t; }; |
||||
template<> struct TypeVec<short, 2> { typedef short2 vec_t; }; |
||||
template<> struct TypeVec<short2, 2> { typedef short2 vec_t; }; |
||||
template<> struct TypeVec<short, 3> { typedef short3 vec_t; }; |
||||
template<> struct TypeVec<short3, 3> { typedef short3 vec_t; }; |
||||
template<> struct TypeVec<short, 4> { typedef short4 vec_t; }; |
||||
template<> struct TypeVec<short4, 4> { typedef short4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<uint, 1> { typedef uint vec_t; }; |
||||
template<> struct TypeVec<uint1, 1> { typedef uint1 vec_t; }; |
||||
template<> struct TypeVec<uint, 2> { typedef uint2 vec_t; }; |
||||
template<> struct TypeVec<uint2, 2> { typedef uint2 vec_t; }; |
||||
template<> struct TypeVec<uint, 3> { typedef uint3 vec_t; }; |
||||
template<> struct TypeVec<uint3, 3> { typedef uint3 vec_t; }; |
||||
template<> struct TypeVec<uint, 4> { typedef uint4 vec_t; }; |
||||
template<> struct TypeVec<uint4, 4> { typedef uint4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<int, 1> { typedef int vec_t; }; |
||||
template<> struct TypeVec<int1, 1> { typedef int1 vec_t; }; |
||||
template<> struct TypeVec<int, 2> { typedef int2 vec_t; }; |
||||
template<> struct TypeVec<int2, 2> { typedef int2 vec_t; }; |
||||
template<> struct TypeVec<int, 3> { typedef int3 vec_t; }; |
||||
template<> struct TypeVec<int3, 3> { typedef int3 vec_t; }; |
||||
template<> struct TypeVec<int, 4> { typedef int4 vec_t; }; |
||||
template<> struct TypeVec<int4, 4> { typedef int4 vec_t; }; |
||||
|
||||
template<> struct TypeVec<float, 1> { typedef float vec_t; }; |
||||
template<> struct TypeVec<float1, 1> { typedef float1 vec_t; }; |
||||
template<> struct TypeVec<float, 2> { typedef float2 vec_t; }; |
||||
template<> struct TypeVec<float2, 2> { typedef float2 vec_t; }; |
||||
template<> struct TypeVec<float, 3> { typedef float3 vec_t; }; |
||||
template<> struct TypeVec<float3, 3> { typedef float3 vec_t; }; |
||||
template<> struct TypeVec<float, 4> { typedef float4 vec_t; }; |
||||
template<> struct TypeVec<float4, 4> { typedef float4 vec_t; }; |
||||
|
||||
template<typename T> struct VecTraits; |
||||
|
||||
template<> struct VecTraits<uchar>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=1}; |
||||
static __device__ uchar all(uchar v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<uchar1>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=1}; |
||||
static __device__ uchar1 all(uchar v) {return make_uchar1(v);} |
||||
}; |
||||
template<> struct VecTraits<uchar2>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ uchar2 all(uchar v) {return make_uchar2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<uchar3>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ uchar3 all(uchar v) {return make_uchar3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<uchar4>
|
||||
{
|
||||
typedef uchar elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ uchar4 all(uchar v) {return make_uchar4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<char>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ char all(char v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<char1>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ char1 all(char v) {return make_char1(v);} |
||||
}; |
||||
template<> struct VecTraits<char2>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ char2 all(char v) {return make_char2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<char3>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ char3 all(char v) {return make_char3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<char4>
|
||||
{
|
||||
typedef char elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ char4 all(char v) {return make_char4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<ushort>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ ushort all(ushort v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<ushort1>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ ushort1 all(ushort v) {return make_ushort1(v);} |
||||
}; |
||||
template<> struct VecTraits<ushort2>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ ushort2 all(ushort v) {return make_ushort2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<ushort3>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ ushort3 all(ushort v) {return make_ushort3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<ushort4>
|
||||
{
|
||||
typedef ushort elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ ushort4 all(ushort v) {return make_ushort4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<short>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ short all(short v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<short1>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ short1 all(short v) {return make_short1(v);} |
||||
}; |
||||
template<> struct VecTraits<short2>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ short2 all(short v) {return make_short2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<short3>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ short3 all(short v) {return make_short3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<short4>
|
||||
{
|
||||
typedef short elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ short4 all(short v) {return make_short4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<uint>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ uint all(uint v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<uint1>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ uint1 all(uint v) {return make_uint1(v);} |
||||
}; |
||||
template<> struct VecTraits<uint2>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ uint2 all(uint v) {return make_uint2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<uint3>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ uint3 all(uint v) {return make_uint3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<uint4>
|
||||
{
|
||||
typedef uint elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ uint4 all(uint v) {return make_uint4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<int>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ int all(int v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<int1>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ int1 all(int v) {return make_int1(v);} |
||||
}; |
||||
template<> struct VecTraits<int2>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ int2 all(int v) {return make_int2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<int3>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ int3 all(int v) {return make_int3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<int4>
|
||||
{
|
||||
typedef int elem_t;
|
||||
enum {cn=4};
|
||||
static __device__ int4 all(int v) {return make_int4(v, v, v, v);} |
||||
}; |
||||
|
||||
template<> struct VecTraits<float>
|
||||
{
|
||||
typedef float elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ float all(float v) {return v;} |
||||
}; |
||||
template<> struct VecTraits<float1>
|
||||
{
|
||||
typedef float elem_t;
|
||||
enum {cn=1};
|
||||
static __device__ float1 all(float v) {return make_float1(v);} |
||||
}; |
||||
template<> struct VecTraits<float2>
|
||||
{
|
||||
typedef float elem_t;
|
||||
enum {cn=2};
|
||||
static __device__ float2 all(float v) {return make_float2(v, v);} |
||||
}; |
||||
template<> struct VecTraits<float3>
|
||||
{
|
||||
typedef float elem_t;
|
||||
enum {cn=3};
|
||||
static __device__ float3 all(float v) {return make_float3(v, v, v);} |
||||
}; |
||||
template<> struct VecTraits<float4>
|
||||
{
|
||||
typedef float elem_t; |
||||
enum {cn=4};
|
||||
static __device__ float4 all(float v) {return make_float4(v, v, v, v);} |
||||
}; |
||||
|
||||
template <int cn, typename VecD> struct SatCast; |
||||
template <typename VecD> struct SatCast<1, VecD> |
||||
{ |
||||
template <typename VecS> |
||||
__device__ VecD operator()(const VecS& v) |
||||
{ |
||||
VecD res;
|
||||
res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x); |
||||
return res; |
||||
} |
||||
}; |
||||
template <typename VecD> struct SatCast<2, VecD> |
||||
{ |
||||
template <typename VecS> |
||||
__device__ VecD operator()(const VecS& v) |
||||
{ |
||||
VecD res;
|
||||
res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y); |
||||
return res; |
||||
} |
||||
}; |
||||
template <typename VecD> struct SatCast<3, VecD> |
||||
{ |
||||
template <typename VecS> |
||||
__device__ VecD operator()(const VecS& v) |
||||
{ |
||||
VecD res;
|
||||
res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.z); |
||||
return res; |
||||
} |
||||
}; |
||||
template <typename VecD> struct SatCast<4, VecD> |
||||
{ |
||||
template <typename VecS> |
||||
__device__ VecD operator()(const VecS& v) |
||||
{ |
||||
VecD res;
|
||||
res.x = saturate_cast< VecTraits<VecD>::elem_t >(v.x); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.y); |
||||
res.y = saturate_cast< VecTraits<VecD>::elem_t >(v.z); |
||||
res.w = saturate_cast< VecTraits<VecD>::elem_t >(v.w); |
||||
return res; |
||||
} |
||||
}; |
||||
|
||||
template <typename VecD, typename VecS> static __device__ VecD saturate_cast_caller(const VecS& v) |
||||
{ |
||||
SatCast< |
||||
|
||||
VecTraits<VecD>::cn,
|
||||
|
||||
VecD |
||||
>
|
||||
|
||||
cast; |
||||
return cast(v); |
||||
} |
||||
|
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uchar1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const char1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const ushort1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const short1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uint1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const int1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const float1& v) {return saturate_cast_caller<_Tp>(v);} |
||||
|
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uchar2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const char2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const ushort2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const short2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uint2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const int2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const float2& v) {return saturate_cast_caller<_Tp>(v);} |
||||
|
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uchar3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const char3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const ushort3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const short3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uint3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const int3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const float3& v) {return saturate_cast_caller<_Tp>(v);} |
||||
|
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uchar4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const char4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const ushort4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const short4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const uint4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const int4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
template<typename _Tp> static __device__ _Tp saturate_cast(const float4& v) {return saturate_cast_caller<_Tp>(v);} |
||||
|
||||
static __device__ uchar1 operator+(const uchar1& a, const uchar1& b) |
||||
{ |
||||
return make_uchar1(a.x + b.x); |
||||
} |
||||
static __device__ uchar1 operator-(const uchar1& a, const uchar1& b) |
||||
{ |
||||
return make_uchar1(a.x - b.x); |
||||
} |
||||
static __device__ uchar1 operator*(const uchar1& a, const uchar1& b) |
||||
{ |
||||
return make_uchar1(a.x * b.x); |
||||
} |
||||
static __device__ uchar1 operator/(const uchar1& a, const uchar1& b) |
||||
{ |
||||
return make_uchar1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const uchar1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ uchar2 operator+(const uchar2& a, const uchar2& b) |
||||
{ |
||||
return make_uchar2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ uchar2 operator-(const uchar2& a, const uchar2& b) |
||||
{ |
||||
return make_uchar2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ uchar2 operator*(const uchar2& a, const uchar2& b) |
||||
{ |
||||
return make_uchar2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ uchar2 operator/(const uchar2& a, const uchar2& b) |
||||
{ |
||||
return make_uchar2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const uchar2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ uchar3 operator+(const uchar3& a, const uchar3& b) |
||||
{ |
||||
return make_uchar3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ uchar3 operator-(const uchar3& a, const uchar3& b) |
||||
{ |
||||
return make_uchar3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ uchar3 operator*(const uchar3& a, const uchar3& b) |
||||
{ |
||||
return make_uchar3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ uchar3 operator/(const uchar3& a, const uchar3& b) |
||||
{ |
||||
return make_uchar3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const uchar3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ uchar4 operator+(const uchar4& a, const uchar4& b) |
||||
{ |
||||
return make_uchar4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ uchar4 operator-(const uchar4& a, const uchar4& b) |
||||
{ |
||||
return make_uchar4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ uchar4 operator*(const uchar4& a, const uchar4& b) |
||||
{ |
||||
return make_uchar4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ uchar4 operator/(const uchar4& a, const uchar4& b) |
||||
{ |
||||
return make_uchar4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const uchar4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ char1 operator+(const char1& a, const char1& b) |
||||
{ |
||||
return make_char1(a.x + b.x); |
||||
} |
||||
static __device__ char1 operator-(const char1& a, const char1& b) |
||||
{ |
||||
return make_char1(a.x - b.x); |
||||
} |
||||
static __device__ char1 operator*(const char1& a, const char1& b) |
||||
{ |
||||
return make_char1(a.x * b.x); |
||||
} |
||||
static __device__ char1 operator/(const char1& a, const char1& b) |
||||
{ |
||||
return make_char1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const char1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ char2 operator+(const char2& a, const char2& b) |
||||
{ |
||||
return make_char2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ char2 operator-(const char2& a, const char2& b) |
||||
{ |
||||
return make_char2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ char2 operator*(const char2& a, const char2& b) |
||||
{ |
||||
return make_char2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ char2 operator/(const char2& a, const char2& b) |
||||
{ |
||||
return make_char2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const char2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ char3 operator+(const char3& a, const char3& b) |
||||
{ |
||||
return make_char3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ char3 operator-(const char3& a, const char3& b) |
||||
{ |
||||
return make_char3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ char3 operator*(const char3& a, const char3& b) |
||||
{ |
||||
return make_char3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ char3 operator/(const char3& a, const char3& b) |
||||
{ |
||||
return make_char3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const char3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ char4 operator+(const char4& a, const char4& b) |
||||
{ |
||||
return make_char4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ char4 operator-(const char4& a, const char4& b) |
||||
{ |
||||
return make_char4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ char4 operator*(const char4& a, const char4& b) |
||||
{ |
||||
return make_char4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ char4 operator/(const char4& a, const char4& b) |
||||
{ |
||||
return make_char4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const char4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ ushort1 operator+(const ushort1& a, const ushort1& b) |
||||
{ |
||||
return make_ushort1(a.x + b.x); |
||||
} |
||||
static __device__ ushort1 operator-(const ushort1& a, const ushort1& b) |
||||
{ |
||||
return make_ushort1(a.x - b.x); |
||||
} |
||||
static __device__ ushort1 operator*(const ushort1& a, const ushort1& b) |
||||
{ |
||||
return make_ushort1(a.x * b.x); |
||||
} |
||||
static __device__ ushort1 operator/(const ushort1& a, const ushort1& b) |
||||
{ |
||||
return make_ushort1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const ushort1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ ushort2 operator+(const ushort2& a, const ushort2& b) |
||||
{ |
||||
return make_ushort2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ ushort2 operator-(const ushort2& a, const ushort2& b) |
||||
{ |
||||
return make_ushort2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ ushort2 operator*(const ushort2& a, const ushort2& b) |
||||
{ |
||||
return make_ushort2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ ushort2 operator/(const ushort2& a, const ushort2& b) |
||||
{ |
||||
return make_ushort2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const ushort2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ ushort3 operator+(const ushort3& a, const ushort3& b) |
||||
{ |
||||
return make_ushort3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ ushort3 operator-(const ushort3& a, const ushort3& b) |
||||
{ |
||||
return make_ushort3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ ushort3 operator*(const ushort3& a, const ushort3& b) |
||||
{ |
||||
return make_ushort3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ ushort3 operator/(const ushort3& a, const ushort3& b) |
||||
{ |
||||
return make_ushort3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const ushort3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ ushort4 operator+(const ushort4& a, const ushort4& b) |
||||
{ |
||||
return make_ushort4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ ushort4 operator-(const ushort4& a, const ushort4& b) |
||||
{ |
||||
return make_ushort4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ ushort4 operator*(const ushort4& a, const ushort4& b) |
||||
{ |
||||
return make_ushort4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ ushort4 operator/(const ushort4& a, const ushort4& b) |
||||
{ |
||||
return make_ushort4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const ushort4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ short1 operator+(const short1& a, const short1& b) |
||||
{ |
||||
return make_short1(a.x + b.x); |
||||
} |
||||
static __device__ short1 operator-(const short1& a, const short1& b) |
||||
{ |
||||
return make_short1(a.x - b.x); |
||||
} |
||||
static __device__ short1 operator*(const short1& a, const short1& b) |
||||
{ |
||||
return make_short1(a.x * b.x); |
||||
} |
||||
static __device__ short1 operator/(const short1& a, const short1& b) |
||||
{ |
||||
return make_short1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const short1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ short2 operator+(const short2& a, const short2& b) |
||||
{ |
||||
return make_short2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ short2 operator-(const short2& a, const short2& b) |
||||
{ |
||||
return make_short2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ short2 operator*(const short2& a, const short2& b) |
||||
{ |
||||
return make_short2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ short2 operator/(const short2& a, const short2& b) |
||||
{ |
||||
return make_short2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const short2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ short3 operator+(const short3& a, const short3& b) |
||||
{ |
||||
return make_short3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ short3 operator-(const short3& a, const short3& b) |
||||
{ |
||||
return make_short3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ short3 operator*(const short3& a, const short3& b) |
||||
{ |
||||
return make_short3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ short3 operator/(const short3& a, const short3& b) |
||||
{ |
||||
return make_short3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const short3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ short4 operator+(const short4& a, const short4& b) |
||||
{ |
||||
return make_short4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ short4 operator-(const short4& a, const short4& b) |
||||
{ |
||||
return make_short4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ short4 operator*(const short4& a, const short4& b) |
||||
{ |
||||
return make_short4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ short4 operator/(const short4& a, const short4& b) |
||||
{ |
||||
return make_short4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const short4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ int1 operator+(const int1& a, const int1& b) |
||||
{ |
||||
return make_int1(a.x + b.x); |
||||
} |
||||
static __device__ int1 operator-(const int1& a, const int1& b) |
||||
{ |
||||
return make_int1(a.x - b.x); |
||||
} |
||||
static __device__ int1 operator*(const int1& a, const int1& b) |
||||
{ |
||||
return make_int1(a.x * b.x); |
||||
} |
||||
static __device__ int1 operator/(const int1& a, const int1& b) |
||||
{ |
||||
return make_int1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const int1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ int2 operator+(const int2& a, const int2& b) |
||||
{ |
||||
return make_int2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ int2 operator-(const int2& a, const int2& b) |
||||
{ |
||||
return make_int2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ int2 operator*(const int2& a, const int2& b) |
||||
{ |
||||
return make_int2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ int2 operator/(const int2& a, const int2& b) |
||||
{ |
||||
return make_int2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const int2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ int3 operator+(const int3& a, const int3& b) |
||||
{ |
||||
return make_int3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ int3 operator-(const int3& a, const int3& b) |
||||
{ |
||||
return make_int3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ int3 operator*(const int3& a, const int3& b) |
||||
{ |
||||
return make_int3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ int3 operator/(const int3& a, const int3& b) |
||||
{ |
||||
return make_int3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const int3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ int4 operator+(const int4& a, const int4& b) |
||||
{ |
||||
return make_int4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ int4 operator-(const int4& a, const int4& b) |
||||
{ |
||||
return make_int4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ int4 operator*(const int4& a, const int4& b) |
||||
{ |
||||
return make_int4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ int4 operator/(const int4& a, const int4& b) |
||||
{ |
||||
return make_int4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const int4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
|
||||
static __device__ float1 operator+(const float1& a, const float1& b) |
||||
{ |
||||
return make_float1(a.x + b.x); |
||||
} |
||||
static __device__ float1 operator-(const float1& a, const float1& b) |
||||
{ |
||||
return make_float1(a.x - b.x); |
||||
} |
||||
static __device__ float1 operator*(const float1& a, const float1& b) |
||||
{ |
||||
return make_float1(a.x * b.x); |
||||
} |
||||
static __device__ float1 operator/(const float1& a, const float1& b) |
||||
{ |
||||
return make_float1(a.x / b.x); |
||||
} |
||||
static __device__ float1 operator*(const float1& a, float s) |
||||
{ |
||||
return make_float1(a.x * s); |
||||
} |
||||
|
||||
static __device__ float2 operator+(const float2& a, const float2& b) |
||||
{ |
||||
return make_float2(a.x + b.x, a.y + b.y); |
||||
} |
||||
static __device__ float2 operator-(const float2& a, const float2& b) |
||||
{ |
||||
return make_float2(a.x - b.x, a.y - b.y); |
||||
} |
||||
static __device__ float2 operator*(const float2& a, const float2& b) |
||||
{ |
||||
return make_float2(a.x * b.x, a.y * b.y); |
||||
} |
||||
static __device__ float2 operator/(const float2& a, const float2& b) |
||||
{ |
||||
return make_float2(a.x / b.x, a.y / b.y); |
||||
} |
||||
static __device__ float2 operator*(const float2& a, float s) |
||||
{ |
||||
return make_float2(a.x * s, a.y * s); |
||||
} |
||||
|
||||
static __device__ float3 operator+(const float3& a, const float3& b) |
||||
{ |
||||
return make_float3(a.x + b.x, a.y + b.y, a.z + b.z); |
||||
} |
||||
static __device__ float3 operator-(const float3& a, const float3& b) |
||||
{ |
||||
return make_float3(a.x - b.x, a.y - b.y, a.z - b.z); |
||||
} |
||||
static __device__ float3 operator*(const float3& a, const float3& b) |
||||
{ |
||||
return make_float3(a.x * b.x, a.y * b.y, a.z * b.z); |
||||
} |
||||
static __device__ float3 operator/(const float3& a, const float3& b) |
||||
{ |
||||
return make_float3(a.x / b.x, a.y / b.y, a.z / b.z); |
||||
} |
||||
static __device__ float3 operator*(const float3& a, float s) |
||||
{ |
||||
return make_float3(a.x * s, a.y * s, a.z * s); |
||||
} |
||||
|
||||
static __device__ float4 operator+(const float4& a, const float4& b) |
||||
{ |
||||
return make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); |
||||
} |
||||
static __device__ float4 operator-(const float4& a, const float4& b) |
||||
{ |
||||
return make_float4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); |
||||
} |
||||
static __device__ float4 operator*(const float4& a, const float4& b) |
||||
{ |
||||
return make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); |
||||
} |
||||
static __device__ float4 operator/(const float4& a, const float4& b) |
||||
{ |
||||
return make_float4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); |
||||
} |
||||
static __device__ float4 operator*(const float4& a, float s) |
||||
{ |
||||
return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); |
||||
} |
||||
}
|
||||
} |
||||
} |
||||
|
||||
#endif // __OPENCV_GPU_VECMATH_HPP__
|
||||
Loading…
Reference in new issue