mirror of
https://github.com/opencv/opencv.git
synced 2024-11-24 19:20:28 +08:00
fixed warnings
This commit is contained in:
parent
045a856c24
commit
bbd519be42
@ -389,7 +389,7 @@ namespace grid_reduce_detail
|
|||||||
// glob_reduce
|
// glob_reduce
|
||||||
|
|
||||||
template <class Reductor, int BLOCK_SIZE, int PATCH_X, int PATCH_Y, class SrcPtr, typename ResType, class MaskPtr>
|
template <class Reductor, int BLOCK_SIZE, int PATCH_X, int PATCH_Y, class SrcPtr, typename ResType, class MaskPtr>
|
||||||
__global__ void glob_reduce(const SrcPtr src, ResType* result, const MaskPtr mask, const int rows, const int cols)
|
__global__ void reduce(const SrcPtr src, ResType* result, const MaskPtr mask, const int rows, const int cols)
|
||||||
{
|
{
|
||||||
const int x0 = blockIdx.x * blockDim.x * PATCH_X + threadIdx.x;
|
const int x0 = blockIdx.x * blockDim.x * PATCH_X + threadIdx.x;
|
||||||
const int y0 = blockIdx.y * blockDim.y * PATCH_Y + threadIdx.y;
|
const int y0 = blockIdx.y * blockDim.y * PATCH_Y + threadIdx.y;
|
||||||
@ -413,12 +413,12 @@ namespace grid_reduce_detail
|
|||||||
}
|
}
|
||||||
|
|
||||||
template <class Reductor, class Policy, class SrcPtr, typename ResType, class MaskPtr>
|
template <class Reductor, class Policy, class SrcPtr, typename ResType, class MaskPtr>
|
||||||
__host__ void glob_reduce(const SrcPtr& src, ResType* result, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
|
__host__ void reduce(const SrcPtr& src, ResType* result, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
|
||||||
{
|
{
|
||||||
const dim3 block(Policy::block_size_x, Policy::block_size_y);
|
const dim3 block(Policy::block_size_x, Policy::block_size_y);
|
||||||
const dim3 grid(divUp(cols, block.x * Policy::patch_size_x), divUp(rows, block.y * Policy::patch_size_y));
|
const dim3 grid(divUp(cols, block.x * Policy::patch_size_x), divUp(rows, block.y * Policy::patch_size_y));
|
||||||
|
|
||||||
glob_reduce<Reductor, Policy::block_size_x * Policy::block_size_y, Policy::patch_size_x, Policy::patch_size_y><<<grid, block, 0, stream>>>(src, result, mask, rows, cols);
|
reduce<Reductor, Policy::block_size_x * Policy::block_size_y, Policy::patch_size_x, Policy::patch_size_y><<<grid, block, 0, stream>>>(src, result, mask, rows, cols);
|
||||||
CV_CUDEV_SAFE_CALL( cudaGetLastError() );
|
CV_CUDEV_SAFE_CALL( cudaGetLastError() );
|
||||||
|
|
||||||
if (stream == 0)
|
if (stream == 0)
|
||||||
@ -433,7 +433,7 @@ namespace grid_reduce_detail
|
|||||||
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
||||||
typedef typename VecTraits<ResType>::elem_type res_elem_type;
|
typedef typename VecTraits<ResType>::elem_type res_elem_type;
|
||||||
|
|
||||||
glob_reduce<SumReductor<src_type, ResType>, Policy>(src, (res_elem_type*) result, mask, rows, cols, stream);
|
reduce<SumReductor<src_type, ResType>, Policy>(src, (res_elem_type*) result, mask, rows, cols, stream);
|
||||||
}
|
}
|
||||||
|
|
||||||
template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
|
template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
|
||||||
@ -441,7 +441,7 @@ namespace grid_reduce_detail
|
|||||||
{
|
{
|
||||||
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
||||||
|
|
||||||
glob_reduce<MinMaxReductor<minop<ResType>, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
|
reduce<MinMaxReductor<minop<ResType>, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
|
||||||
}
|
}
|
||||||
|
|
||||||
template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
|
template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
|
||||||
@ -449,7 +449,7 @@ namespace grid_reduce_detail
|
|||||||
{
|
{
|
||||||
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
||||||
|
|
||||||
glob_reduce<MinMaxReductor<maxop<ResType>, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
|
reduce<MinMaxReductor<maxop<ResType>, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
|
||||||
}
|
}
|
||||||
|
|
||||||
template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
|
template <class Policy, class SrcPtr, typename ResType, class MaskPtr>
|
||||||
@ -457,7 +457,7 @@ namespace grid_reduce_detail
|
|||||||
{
|
{
|
||||||
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
||||||
|
|
||||||
glob_reduce<MinMaxReductor<both, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
|
reduce<MinMaxReductor<both, src_type, ResType>, Policy>(src, result, mask, rows, cols, stream);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -61,7 +61,7 @@ __host__ void gridCalcSum_(const SrcPtr& src, GpuMat_<ResType>& dst, const MaskP
|
|||||||
{
|
{
|
||||||
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
||||||
|
|
||||||
CV_StaticAssert( VecTraits<src_type>::cn == VecTraits<ResType>::cn, "" );
|
CV_StaticAssert( unsigned(VecTraits<src_type>::cn) == unsigned(VecTraits<ResType>::cn), "" );
|
||||||
|
|
||||||
dst.create(1, 1);
|
dst.create(1, 1);
|
||||||
dst.setTo(0, stream);
|
dst.setTo(0, stream);
|
||||||
@ -83,7 +83,7 @@ __host__ void gridCalcSum_(const SrcPtr& src, GpuMat_<ResType>& dst, Stream& str
|
|||||||
{
|
{
|
||||||
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
typedef typename PtrTraits<SrcPtr>::value_type src_type;
|
||||||
|
|
||||||
CV_StaticAssert( VecTraits<src_type>::cn == VecTraits<ResType>::cn, "" );
|
CV_StaticAssert( unsigned(VecTraits<src_type>::cn) == unsigned(VecTraits<ResType>::cn), "" );
|
||||||
|
|
||||||
dst.create(1, 1);
|
dst.create(1, 1);
|
||||||
dst.setTo(0, stream);
|
dst.setTo(0, stream);
|
||||||
|
Loading…
Reference in New Issue
Block a user