Commit 1d93ca00 authored by Andrey Morozov's avatar Andrey Morozov

Added files for implementation of operations SetTo()

parent 08cba33f
set(name "gpu")
set(name "gpu")
set(DEPS "opencv_core") set(DEPS "opencv_core")
set(the_target "opencv_${name}") set(the_target "opencv_${name}")
...@@ -15,20 +14,20 @@ include_directories("${CMAKE_CURRENT_SOURCE_DIR}/include" ...@@ -15,20 +14,20 @@ include_directories("${CMAKE_CURRENT_SOURCE_DIR}/include"
"${CMAKE_CURRENT_BINARY_DIR}") "${CMAKE_CURRENT_BINARY_DIR}")
foreach(d ${DEPS}) foreach(d ${DEPS})
if(${d} MATCHES "opencv_") if(${d} MATCHES "opencv_")
string(REPLACE "opencv_" "${CMAKE_CURRENT_SOURCE_DIR}/../" d_dir ${d}) string(REPLACE "opencv_" "${CMAKE_CURRENT_SOURCE_DIR}/../" d_dir ${d})
include_directories("${d_dir}/include") include_directories("${d_dir}/include")
endif() endif()
endforeach() endforeach()
file(GLOB lib_srcs "src/*.cpp") file(GLOB lib_srcs "src/*.cpp")
file(GLOB lib_int_hdrs "src/*.h*") file(GLOB lib_int_hdrs "src/*.h*")
file(GLOB lib_cuda "cuda/*.cu") file(GLOB lib_cuda "cuda/*.cu")
file(GLOB lib_cuda_hdrs "cuda/*.h*") file(GLOB lib_cuda_hdrs "cuda/*.h*")
source_group("Src" FILES ${lib_srcs} ${lib_int_hdrs}) source_group("Src" FILES ${lib_srcs} ${lib_int_hdrs})
source_group("Cuda" FILES ${lib_cuda} ${lib_cuda_hdrs}) source_group("Cuda" FILES ${lib_cuda} ${lib_cuda_hdrs})
file(GLOB lib_hdrs "include/opencv2/${name}/*.h*") file(GLOB lib_hdrs "include/opencv2/${name}/*.h*")
source_group("Include" FILES ${lib_hdrs}) source_group("Include" FILES ${lib_hdrs})
if (HAVE_CUDA) if (HAVE_CUDA)
...@@ -38,11 +37,11 @@ if (HAVE_CUDA) ...@@ -38,11 +37,11 @@ if (HAVE_CUDA)
if (UNIX OR APPLE) if (UNIX OR APPLE)
set (CUDA_NVCC_FLAGS "-Xcompiler;-fPIC") set (CUDA_NVCC_FLAGS "-Xcompiler;-fPIC")
endif() endif()
CUDA_COMPILE(cuda_objs ${lib_cuda}) CUDA_COMPILE(cuda_objs ${lib_cuda})
#CUDA_BUILD_CLEAN_TARGET() #CUDA_BUILD_CLEAN_TARGET()
endif() endif()
add_library(${the_target} ${lib_srcs} ${lib_hdrs} ${lib_int_hdrs} ${lib_cuda} ${lib_cuda_hdrs} ${cuda_objs}) add_library(${the_target} ${lib_srcs} ${lib_hdrs} ${lib_int_hdrs} ${lib_cuda} ${lib_cuda_hdrs} ${cuda_objs})
...@@ -51,7 +50,7 @@ if(PCHSupport_FOUND) ...@@ -51,7 +50,7 @@ if(PCHSupport_FOUND)
if(${CMAKE_GENERATOR} MATCHES "Visual*" OR ${CMAKE_GENERATOR} MATCHES "Xcode*") if(${CMAKE_GENERATOR} MATCHES "Visual*" OR ${CMAKE_GENERATOR} MATCHES "Xcode*")
if(${CMAKE_GENERATOR} MATCHES "Visual*") if(${CMAKE_GENERATOR} MATCHES "Visual*")
set(${the_target}_pch "src/precomp.cpp") set(${the_target}_pch "src/precomp.cpp")
endif() endif()
add_native_precompiled_header(${the_target} ${pch_header}) add_native_precompiled_header(${the_target} ${pch_header})
elseif(CMAKE_COMPILER_IS_GNUCXX AND ${CMAKE_GENERATOR} MATCHES ".*Makefiles") elseif(CMAKE_COMPILER_IS_GNUCXX AND ${CMAKE_GENERATOR} MATCHES ".*Makefiles")
add_precompiled_header(${the_target} ${pch_header}) add_precompiled_header(${the_target} ${pch_header})
......
...@@ -48,18 +48,21 @@ ...@@ -48,18 +48,21 @@
namespace cv namespace cv
{ {
namespace gpu namespace gpu
{ {
typedef unsigned char uchar; typedef unsigned char uchar;
typedef unsigned short ushort; typedef unsigned short ushort;
typedef unsigned int uint; typedef unsigned int uint;
extern "C" void error( const char *error_string, const char *file, const int line, const char *func = ""); extern "C" void error( const char *error_string, const char *file, const int line, const char *func = "");
namespace impl namespace impl
{ {
static inline int divUp(int a, int b) { return (a % b == 0) ? a/b : a/b + 1; } static inline int divUp(int a, int b) { return (a % b == 0) ? a/b : a/b + 1; }
extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, DevMem2D& disp, int maxdisp, DevMem2D_<uint>& minSSD_buf); extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, DevMem2D& disp, int maxdisp, DevMem2D_<uint>& minSSD_buf);
extern "C" void set_to_without_mask (const DevMem2D& mat, const double * scalar, int depth, int channels);
extern "C" void set_to_with_mask (const DevMem2D& mat, const double * scalar, const DevMem2D& mask, int depth, int channels);
} }
} }
} }
......
/*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 "cuda_shared.hpp"
#include "cuda_runtime.h"
__constant__ float scalar_d[4];
namespace mat_operators
{
template <typename T, int channels>
__global__ void kernel_set_to_without_mask(T * mat)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
mat[i * sizeof(T)] = static_cast<T>(scalar_d[i % channels]);
}
}
extern "C" void cv::gpu::impl::set_to_with_mask(const DevMem2D& mat, const double * scalar, const DevMem2D& mask, int depth, int channels)
{
}
extern "C" void cv::gpu::impl::set_to_without_mask(const DevMem2D& mat, const double * scalar, int depth, int channels)
{
scalar_d[0] = scalar[0];
scalar_d[1] = scalar[1];
scalar_d[2] = scalar[2];
scalar_d[3] = scalar[3];
int numBlocks = mat.rows * mat.step / 256;
dim3 threadsPerBlock(256);
if (channels == 1)
{
if (depth == 1) ::mat_operators::kernel_set_to_without_mask<unsigned char, 1><<<numBlocks,threadsPerBlock>>>(mat.ptr);
if (depth == 2) ::mat_operators::kernel_set_to_without_mask<unsigned short, 1><<<numBlocks,threadsPerBlock>>>((unsigned short *)mat.ptr);
if (depth == 4) ::mat_operators::kernel_set_to_without_mask<unsigned int, 1><<<numBlocks,threadsPerBlock>>>((unsigned int *)mat.ptr);
}
if (channels == 2)
{
if (depth == 1) ::mat_operators::kernel_set_to_without_mask<unsigned char, 2><<<numBlocks,threadsPerBlock>>>(mat.ptr);
if (depth == 2) ::mat_operators::kernel_set_to_without_mask<unsigned short, 2><<<numBlocks,threadsPerBlock>>>((unsigned short *)mat.ptr);
if (depth == 4) ::mat_operators::kernel_set_to_without_mask<unsigned int, 2><<<numBlocks,threadsPerBlock>>>((unsigned int *)mat.ptr);
}
if (channels == 3)
{
if (depth == 1) ::mat_operators::kernel_set_to_without_mask<unsigned char, 3><<<numBlocks,threadsPerBlock>>>(mat.ptr);
if (depth == 2) ::mat_operators::kernel_set_to_without_mask<unsigned short, 3><<<numBlocks,threadsPerBlock>>>((unsigned short *)mat.ptr);
if (depth == 4) ::mat_operators::kernel_set_to_without_mask<unsigned int, 3><<<numBlocks,threadsPerBlock>>>((unsigned int *)mat.ptr);
}
}
...@@ -41,23 +41,23 @@ ...@@ -41,23 +41,23 @@
//M*/ //M*/
#include "precomp.hpp" #include "precomp.hpp"
#include "opencv2/gpu/stream_access.hpp" //#include "opencv2/gpu/stream_access.hpp"
using namespace cv; using namespace cv;
using namespace cv::gpu; using namespace cv::gpu;
cv::gpu::CudaStream::CudaStream() : impl( (Impl*)fastMalloc(sizeof(Impl)) ) cv::gpu::CudaStream::CudaStream() //: impl( (Impl*)fastMalloc(sizeof(Impl)) )
{ {
//cudaSafeCall( cudaStreamCreate( &impl->stream) ); //cudaSafeCall( cudaStreamCreate( &impl->stream) );
} }
cv::gpu::CudaStream::~CudaStream() cv::gpu::CudaStream::~CudaStream()
{ {
if (impl) if (impl)
{ {
cudaSafeCall( cudaStreamDestroy( *(cudaStream_t*)impl ) ); cudaSafeCall( cudaStreamDestroy( *(cudaStream_t*)impl ) );
cv::fastFree( impl ); cv::fastFree( impl );
} }
} }
bool cv::gpu::CudaStream::queryIfComplete() bool cv::gpu::CudaStream::queryIfComplete()
...@@ -70,8 +70,8 @@ bool cv::gpu::CudaStream::queryIfComplete() ...@@ -70,8 +70,8 @@ bool cv::gpu::CudaStream::queryIfComplete()
//if (err == cudaErrorNotReady) //if (err == cudaErrorNotReady)
// return false; // return false;
////cudaErrorInvalidResourceHandle ////cudaErrorInvalidResourceHandle
//cudaSafeCall( err ); //cudaSafeCall( err );
return true; return true;
} }
void cv::gpu::CudaStream::waitForCompletion() void cv::gpu::CudaStream::waitForCompletion()
...@@ -81,7 +81,7 @@ void cv::gpu::CudaStream::waitForCompletion() ...@@ -81,7 +81,7 @@ void cv::gpu::CudaStream::waitForCompletion()
void cv::gpu::CudaStream::enqueueDownload(const GpuMat& src, Mat& dst) void cv::gpu::CudaStream::enqueueDownload(const GpuMat& src, Mat& dst)
{ {
// cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToHost, // cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToHost,
} }
void cv::gpu::CudaStream::enqueueUpload(const Mat& src, GpuMat& dst) void cv::gpu::CudaStream::enqueueUpload(const Mat& src, GpuMat& dst)
{ {
...@@ -109,4 +109,4 @@ void cv::gpu::CudaStream::enqueueConvert(const GpuMat& src, GpuMat& dst, int typ ...@@ -109,4 +109,4 @@ void cv::gpu::CudaStream::enqueueConvert(const GpuMat& src, GpuMat& dst, int typ
//struct cudaStream_t& cv::gpu::CudaStream::getStream() { return stream; } //struct cudaStream_t& cv::gpu::CudaStream::getStream() { return stream; }
...@@ -68,26 +68,42 @@ void GpuMat::copyTo( GpuMat& m ) const ...@@ -68,26 +68,42 @@ void GpuMat::copyTo( GpuMat& m ) const
cudaSafeCall( cudaMemcpy2D(m.data, m.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice) ); cudaSafeCall( cudaMemcpy2D(m.data, m.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice) );
cudaSafeCall( cudaThreadSynchronize() ); cudaSafeCall( cudaThreadSynchronize() );
} }
void GpuMat::copyTo( GpuMat& /*m*/, const GpuMat&/* mask */) const void GpuMat::copyTo( GpuMat& /*m*/, const GpuMat&/* mask */) const
{ {
CV_Assert(!"Not implemented"); CV_Assert(!"Not implemented");
} }
void GpuMat::convertTo( GpuMat& /*m*/, int /*rtype*/, double /*alpha*/, double /*beta*/ ) const void GpuMat::convertTo( GpuMat& /*m*/, int /*rtype*/, double /*alpha*/, double /*beta*/ ) const
{ {
CV_Assert(!"Not implemented"); CV_Assert(!"Not implemented");
} }
GpuMat& GpuMat::operator = (const Scalar& /*s*/) GpuMat& GpuMat::operator = (const Scalar& s)
{ {
CV_Assert(!"Not implemented"); CV_Assert(!"Not implemented");
cv::gpu::impl::set_to_without_mask(*this, s.val, this->depth(), this->channels());
return *this; return *this;
} }
GpuMat& GpuMat::setTo(const Scalar& /*s*/, const GpuMat& /*mask*/) GpuMat& GpuMat::setTo(const Scalar& s, const GpuMat& mask)
{ {
CV_Assert(!"Not implemented"); CV_Assert(!"Not implemented");
CV_DbgAssert(!this->empty());
this->channels();
this->depth();
if (mask.empty())
{
cv::gpu::impl::set_to_without_mask(*this, s.val, this->depth(), this->channels());
}
else
{
cv::gpu::impl::set_to_with_mask(*this, s.val, mask, this->depth(), this->channels());
}
return *this; return *this;
} }
...@@ -147,8 +163,8 @@ void GpuMat::create(int _rows, int _cols, int _type) ...@@ -147,8 +163,8 @@ void GpuMat::create(int _rows, int _cols, int _type)
rows = _rows; rows = _rows;
cols = _cols; cols = _cols;
size_t esz = elemSize(); size_t esz = elemSize();
void *dev_ptr; void *dev_ptr;
cudaSafeCall( cudaMallocPitch(&dev_ptr, &step, esz * cols, rows) ); cudaSafeCall( cudaMallocPitch(&dev_ptr, &step, esz * cols, rows) );
...@@ -157,10 +173,10 @@ void GpuMat::create(int _rows, int _cols, int _type) ...@@ -157,10 +173,10 @@ void GpuMat::create(int _rows, int _cols, int _type)
int64 _nettosize = (int64)step*rows; int64 _nettosize = (int64)step*rows;
size_t nettosize = (size_t)_nettosize; size_t nettosize = (size_t)_nettosize;
datastart = data = (uchar*)dev_ptr; datastart = data = (uchar*)dev_ptr;
dataend = data + nettosize; dataend = data + nettosize;
refcount = (int*)fastMalloc(sizeof(*refcount)); refcount = (int*)fastMalloc(sizeof(*refcount));
*refcount = 1; *refcount = 1;
} }
...@@ -171,7 +187,7 @@ void GpuMat::release() ...@@ -171,7 +187,7 @@ void GpuMat::release()
if( refcount && CV_XADD(refcount, -1) == 1 ) if( refcount && CV_XADD(refcount, -1) == 1 )
{ {
fastFree(refcount); fastFree(refcount);
cudaSafeCall( cudaFree(datastart) ); cudaSafeCall( cudaFree(datastart) );
} }
data = datastart = dataend = 0; data = datastart = dataend = 0;
step = rows = cols = 0; step = rows = cols = 0;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment