mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-26 20:20:20 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

minor tab fix

Browse Source
This commit is contained in:
Marina Kolpakova 2012-04-24 09:30:21 +00:00
parent f1a0ab365d
commit 99aed4713f
2 changed files with 59 additions and 59 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

92
modules/core/include/opencv2/core/devmem2d.hpp
View File

@ -45,14 +45,14 @@
#ifdef __cplusplus #ifdef __cplusplus
#ifdef __CUDACC__ #ifdef __CUDACC__
#define __CV_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__ #define __CV_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__
#else #else
#define __CV_GPU_HOST_DEVICE__ #define __CV_GPU_HOST_DEVICE__
#endif #endif
namespace cv namespace cv
{ {
namespace gpu namespace gpu
{ {
// Simple lightweight structures that encapsulates information about an image on device. // Simple lightweight structures that encapsulates information about an image on device.
@ -61,88 +61,88 @@ namespace cv
template <bool expr> struct StaticAssert; template <bool expr> struct StaticAssert;
template <> struct StaticAssert<true> {static __CV_GPU_HOST_DEVICE__ void check(){}}; template <> struct StaticAssert<true> {static __CV_GPU_HOST_DEVICE__ void check(){}};
template<typename T> struct DevPtr template<typename T> struct DevPtr
{ {
typedef T elem_type; typedef T elem_type;
typedef int index_type; typedef int index_type;
enum { elem_size = sizeof(elem_type) }; enum { elem_size = sizeof(elem_type) };
T* data; T* data;
__CV_GPU_HOST_DEVICE__ DevPtr() : data(0) {} __CV_GPU_HOST_DEVICE__ DevPtr() : data(0) {}
__CV_GPU_HOST_DEVICE__ DevPtr(T* data_) : data(data_) {} __CV_GPU_HOST_DEVICE__ DevPtr(T* data_) : data(data_) {}
__CV_GPU_HOST_DEVICE__ size_t elemSize() const { return elem_size; } __CV_GPU_HOST_DEVICE__ size_t elemSize() const { return elem_size; }
__CV_GPU_HOST_DEVICE__ operator T*() { return data; } __CV_GPU_HOST_DEVICE__ operator T*() { return data; }
__CV_GPU_HOST_DEVICE__ operator const T*() const { return data; } __CV_GPU_HOST_DEVICE__ operator const T*() const { return data; }
}; };
template<typename T> struct PtrSz : public DevPtr<T> template<typename T> struct PtrSz : public DevPtr<T>
{ {
__CV_GPU_HOST_DEVICE__ PtrSz() : size(0) {} __CV_GPU_HOST_DEVICE__ PtrSz() : size(0) {}
__CV_GPU_HOST_DEVICE__ PtrSz(T* data_, size_t size_) : DevPtr<T>(data_), size(size_) {} __CV_GPU_HOST_DEVICE__ PtrSz(T* data_, size_t size_) : DevPtr<T>(data_), size(size_) {}
size_t size; size_t size;
}; };
template<typename T> struct PtrStep : public DevPtr<T> template<typename T> struct PtrStep : public DevPtr<T>
{ {
__CV_GPU_HOST_DEVICE__ PtrStep() : step(0) {} __CV_GPU_HOST_DEVICE__ PtrStep() : step(0) {}
__CV_GPU_HOST_DEVICE__ PtrStep(T* data_, size_t step_) : DevPtr<T>(data_), step(step_) {} __CV_GPU_HOST_DEVICE__ PtrStep(T* data_, size_t step_) : DevPtr<T>(data_), step(step_) {}
/** \brief stride between two consecutive rows in bytes. Step is stored always and everywhere in bytes!!! */ /** \brief stride between two consecutive rows in bytes. Step is stored always and everywhere in bytes!!! */
size_t step; size_t step;
__CV_GPU_HOST_DEVICE__ T* ptr(int y = 0) { return ( T*)( ( char*)DevPtr<T>::data + y * step); } __CV_GPU_HOST_DEVICE__ T* ptr(int y = 0) { return ( T*)( ( char*)DevPtr<T>::data + y * step); }
__CV_GPU_HOST_DEVICE__ const T* ptr(int y = 0) const { return (const T*)( (const char*)DevPtr<T>::data + y * step); } __CV_GPU_HOST_DEVICE__ const T* ptr(int y = 0) const { return (const T*)( (const char*)DevPtr<T>::data + y * step); }
__CV_GPU_HOST_DEVICE__ T& operator ()(int y, int x) { return ptr(y)[x]; } __CV_GPU_HOST_DEVICE__ T& operator ()(int y, int x) { return ptr(y)[x]; }
__CV_GPU_HOST_DEVICE__ const T& operator ()(int y, int x) const { return ptr(y)[x]; } __CV_GPU_HOST_DEVICE__ const T& operator ()(int y, int x) const { return ptr(y)[x]; }
}; };
template <typename T> struct PtrStepSz : public PtrStep<T> template <typename T> struct PtrStepSz : public PtrStep<T>
{ {
__CV_GPU_HOST_DEVICE__ PtrStepSz() : cols(0), rows(0) {} __CV_GPU_HOST_DEVICE__ PtrStepSz() : cols(0), rows(0) {}
__CV_GPU_HOST_DEVICE__ PtrStepSz(int rows_, int cols_, T* data_, size_t step_) __CV_GPU_HOST_DEVICE__ PtrStepSz(int rows_, int cols_, T* data_, size_t step_)
: PtrStep<T>(data_, step_), cols(cols_), rows(rows_) {} : PtrStep<T>(data_, step_), cols(cols_), rows(rows_) {}
int cols; int cols;
int rows; int rows;
}; };
template <typename T> struct DevMem2D_ : public PtrStepSz<T> template <typename T> struct DevMem2D_ : public PtrStepSz<T>
{ {
DevMem2D_() {} DevMem2D_() {}
DevMem2D_(int rows_, int cols_, T* data_, size_t step_) : PtrStepSz<T>(rows_, cols_, data_, step_) {} DevMem2D_(int rows_, int cols_, T* data_, size_t step_) : PtrStepSz<T>(rows_, cols_, data_, step_) {}
template <typename U> template <typename U>
explicit DevMem2D_(const DevMem2D_<U>& d) : PtrStepSz<T>(d.rows, d.cols, (T*)d.data, d.step) {} explicit DevMem2D_(const DevMem2D_<U>& d) : PtrStepSz<T>(d.rows, d.cols, (T*)d.data, d.step) {}
}; };
template<typename T> struct PtrElemStep_ : public PtrStep<T> template<typename T> struct PtrElemStep_ : public PtrStep<T>
{ {
PtrElemStep_(const DevMem2D_<T>& mem) : PtrStep<T>(mem.data, mem.step) PtrElemStep_(const DevMem2D_<T>& mem) : PtrStep<T>(mem.data, mem.step)
{ {
StaticAssert<256 % sizeof(T) == 0>::check(); StaticAssert<256 % sizeof(T) == 0>::check();
PtrStep<T>::step /= PtrStep<T>::elem_size; PtrStep<T>::step /= PtrStep<T>::elem_size;
} }
__CV_GPU_HOST_DEVICE__ T* ptr(int y = 0) { return PtrStep<T>::data + y * PtrStep<T>::step; } __CV_GPU_HOST_DEVICE__ T* ptr(int y = 0) { return PtrStep<T>::data + y * PtrStep<T>::step; }
__CV_GPU_HOST_DEVICE__ const T* ptr(int y = 0) const { return PtrStep<T>::data + y * PtrStep<T>::step; } __CV_GPU_HOST_DEVICE__ const T* ptr(int y = 0) const { return PtrStep<T>::data + y * PtrStep<T>::step; }
__CV_GPU_HOST_DEVICE__ T& operator ()(int y, int x) { return ptr(y)[x]; } __CV_GPU_HOST_DEVICE__ T& operator ()(int y, int x) { return ptr(y)[x]; }
__CV_GPU_HOST_DEVICE__ const T& operator ()(int y, int x) const { return ptr(y)[x]; } __CV_GPU_HOST_DEVICE__ const T& operator ()(int y, int x) const { return ptr(y)[x]; }
}; };
template<typename T> struct PtrStep_ : public PtrStep<T> template<typename T> struct PtrStep_ : public PtrStep<T>
{ {
PtrStep_() {} PtrStep_() {}
PtrStep_(const DevMem2D_<T>& mem) : PtrStep<T>(mem.data, mem.step) {} PtrStep_(const DevMem2D_<T>& mem) : PtrStep<T>(mem.data, mem.step) {}
}; };
typedef DevMem2D_<unsigned char> DevMem2Db; typedef DevMem2D_<unsigned char> DevMem2Db;
typedef DevMem2Db DevMem2D; typedef DevMem2Db DevMem2D;
typedef DevMem2D_<float> DevMem2Df; typedef DevMem2D_<float> DevMem2Df;
typedef DevMem2D_<int> DevMem2Di; typedef DevMem2D_<int> DevMem2Di;
@ -152,8 +152,8 @@ namespace cv
typedef PtrElemStep_<unsigned char> PtrElemStep; typedef PtrElemStep_<unsigned char> PtrElemStep;
typedef PtrElemStep_<float> PtrElemStepf; typedef PtrElemStep_<float> PtrElemStepf;
typedef PtrElemStep_<int> PtrElemStepi; typedef PtrElemStep_<int> PtrElemStepi;
} }
} }
#endif // __cplusplus #endif // __cplusplus

26
modules/gpu/src/nvidia/NCVBroxOpticalFlow.cu
View File

@ -172,11 +172,11 @@ static void add(float *res, const float *rhs, const int count, cudaStream_t stre
/////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////////
__global__ void scaleVector(float *d_res, const float *d_src, float scale, const int len) __global__ void scaleVector(float *d_res, const float *d_src, float scale, const int len)
{ {
const int pos = blockIdx.x * blockDim.x + threadIdx.x; const int pos = blockIdx.x * blockDim.x + threadIdx.x;
if (pos >= len) return; if (pos >= len) return;
d_res[pos] = d_src[pos] * scale; d_res[pos] = d_src[pos] * scale;
} }
/////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////////
@ -191,10 +191,10 @@ __global__ void scaleVector(float *d_res, const float *d_src, float scale, const
/////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////////
static void ScaleVector(float *d_res, const float *d_src, float scale, const int len, cudaStream_t stream) static void ScaleVector(float *d_res, const float *d_src, float scale, const int len, cudaStream_t stream)
{ {
dim3 threads(256); dim3 threads(256);
dim3 blocks(iDivUp(len, threads.x)); dim3 blocks(iDivUp(len, threads.x));
scaleVector<<<blocks, threads, 0, stream>>>(d_res, d_src, scale, len); scaleVector<<<blocks, threads, 0, stream>>>(d_res, d_src, scale, len);
} }
const int SOR_TILE_WIDTH = 32; const int SOR_TILE_WIDTH = 32;
@ -1128,14 +1128,14 @@ NCVStatus NCVBroxOpticalFlow(const NCVBroxOpticalFlowDescriptor desc,
ncvAssertReturnNcvStat( nppiStResize_32f_C1R (ptrU->ptr(), srcSize, kLevelStride * sizeof (float), srcROI, ncvAssertReturnNcvStat( nppiStResize_32f_C1R (ptrU->ptr(), srcSize, kLevelStride * sizeof (float), srcROI,
ptrUNew->ptr(), dstSize, ns * sizeof (float), dstROI, 1.0f/scale_factor, 1.0f/scale_factor, nppStBicubic) ); ptrUNew->ptr(), dstSize, ns * sizeof (float), dstROI, 1.0f/scale_factor, 1.0f/scale_factor, nppStBicubic) );
ScaleVector(ptrUNew->ptr(), ptrUNew->ptr(), 1.0f/scale_factor, ns * nh, stream); ScaleVector(ptrUNew->ptr(), ptrUNew->ptr(), 1.0f/scale_factor, ns * nh, stream);
ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR); ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR);
ncvAssertReturnNcvStat( nppiStResize_32f_C1R (ptrV->ptr(), srcSize, kLevelStride * sizeof (float), srcROI, ncvAssertReturnNcvStat( nppiStResize_32f_C1R (ptrV->ptr(), srcSize, kLevelStride * sizeof (float), srcROI,
ptrVNew->ptr(), dstSize, ns * sizeof (float), dstROI, 1.0f/scale_factor, 1.0f/scale_factor, nppStBicubic) ); ptrVNew->ptr(), dstSize, ns * sizeof (float), dstROI, 1.0f/scale_factor, 1.0f/scale_factor, nppStBicubic) );
ScaleVector(ptrVNew->ptr(), ptrVNew->ptr(), 1.0f/scale_factor, ns * nh, stream); ScaleVector(ptrVNew->ptr(), ptrVNew->ptr(), 1.0f/scale_factor, ns * nh, stream);
ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR); ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR);
cv::gpu::device::swap<FloatVector*>(ptrU, ptrUNew); cv::gpu::device::swap<FloatVector*>(ptrU, ptrUNew);