mirror of
https://github.com/opencv/opencv.git
synced 2026-07-29 23:33:05 +04:00
Merge branch '4.x' into '5.x'
This commit is contained in:
Maksim Shabunin
@@ -917,11 +917,8 @@ CV_EXPORTS_W void reduceArgMax(InputArray src, OutputArray dst, int axis, bool l
|
||||
|
||||
The function cv::minMaxIdx finds the minimum and maximum element values and their positions. The
|
||||
extremums are searched across the whole array or, if mask is not an empty array, in the specified
|
||||
array region. The function does not work with multi-channel arrays. If you need to find minimum or
|
||||
maximum elements across all the channels, use Mat::reshape first to reinterpret the array as
|
||||
single-channel. Or you may extract the particular channel using either extractImageCOI, or
|
||||
mixChannels, or split. In case of a sparse matrix, the minimum is found among non-zero elements
|
||||
only.
|
||||
array region. In case of a sparse matrix, the minimum is found among non-zero elements
|
||||
only. Multi-channel input is supported without mask and extremums indexes (should be nullptr).
|
||||
@note When minIdx is not NULL, it must have at least 2 elements (as well as maxIdx), even if src is
|
||||
a single-row or single-column matrix. In OpenCV (following MATLAB) each array has at least 2
|
||||
dimensions, i.e. single-column matrix is Mx1 matrix (and therefore minIdx/maxIdx will be
|
||||
@@ -1825,7 +1822,7 @@ CV_EXPORTS_W void transpose(InputArray src, OutputArray dst);
|
||||
* @note Input should be continuous single-channel matrix.
|
||||
* @param src input array.
|
||||
* @param order a permutation of [0,1,..,N-1] where N is the number of axes of src.
|
||||
* The i’th axis of dst will correspond to the axis numbered order[i] of the input.
|
||||
* The i'th axis of dst will correspond to the axis numbered order[i] of the input.
|
||||
* @param dst output array of the same type as src.
|
||||
*/
|
||||
CV_EXPORTS_W void transposeND(InputArray src, const std::vector<int>& order, OutputArray dst);
|
||||
|
||||
@@ -134,6 +134,22 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
val = smem[tid];
|
||||
}
|
||||
|
||||
template <typename T, class Op>
|
||||
__device__ __forceinline__ void merge(volatile T* smem, T& val, unsigned int tid, unsigned int delta, const Op& op)
|
||||
{
|
||||
T reg = smem[tid + delta];
|
||||
smem[tid] = val = op(val, reg);
|
||||
}
|
||||
|
||||
template <typename T, class Op>
|
||||
__device__ __forceinline__ void mergeShfl(T& val, unsigned int delta, unsigned int width, const Op& op)
|
||||
{
|
||||
T reg = shfl_down(val, delta, width);
|
||||
val = op(val, reg);
|
||||
}
|
||||
|
||||
#if (CUDART_VERSION < 12040) // details: https://github.com/opencv/opencv_contrib/issues/3690
|
||||
template <typename P0, typename P1, typename P2, typename P3, typename P4, typename P5, typename P6, typename P7, typename P8, typename P9,
|
||||
typename R0, typename R1, typename R2, typename R3, typename R4, typename R5, typename R6, typename R7, typename R8, typename R9>
|
||||
__device__ __forceinline__ void loadToSmem(const thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>& smem,
|
||||
@@ -142,6 +158,7 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9> >::value>::loadToSmem(smem, val, tid);
|
||||
}
|
||||
|
||||
template <typename P0, typename P1, typename P2, typename P3, typename P4, typename P5, typename P6, typename P7, typename P8, typename P9,
|
||||
typename R0, typename R1, typename R2, typename R3, typename R4, typename R5, typename R6, typename R7, typename R8, typename R9>
|
||||
__device__ __forceinline__ void loadFromSmem(const thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>& smem,
|
||||
@@ -151,18 +168,6 @@ namespace cv { namespace cuda { namespace device
|
||||
For<0, thrust::tuple_size<thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9> >::value>::loadFromSmem(smem, val, tid);
|
||||
}
|
||||
|
||||
template <typename T, class Op>
|
||||
__device__ __forceinline__ void merge(volatile T* smem, T& val, unsigned int tid, unsigned int delta, const Op& op)
|
||||
{
|
||||
T reg = smem[tid + delta];
|
||||
smem[tid] = val = op(val, reg);
|
||||
}
|
||||
template <typename T, class Op>
|
||||
__device__ __forceinline__ void mergeShfl(T& val, unsigned int delta, unsigned int width, const Op& op)
|
||||
{
|
||||
T reg = shfl_down(val, delta, width);
|
||||
val = op(val, reg);
|
||||
}
|
||||
template <typename P0, typename P1, typename P2, typename P3, typename P4, typename P5, typename P6, typename P7, typename P8, typename P9,
|
||||
typename R0, typename R1, typename R2, typename R3, typename R4, typename R5, typename R6, typename R7, typename R8, typename R9,
|
||||
class Op0, class Op1, class Op2, class Op3, class Op4, class Op5, class Op6, class Op7, class Op8, class Op9>
|
||||
@@ -183,7 +188,31 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9> >::value>::mergeShfl(val, delta, width, op);
|
||||
}
|
||||
#else
|
||||
template <typename... P, typename... R>
|
||||
__device__ __forceinline__ void loadToSmem(const thrust::tuple<P...>& smem, const thrust::tuple<R...>& val, unsigned int tid)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<P...> >::value>::loadToSmem(smem, val, tid);
|
||||
}
|
||||
|
||||
template <typename... P, typename... R>
|
||||
__device__ __forceinline__ void loadFromSmem(const thrust::tuple<P...>& smem, const thrust::tuple<R...>& val, unsigned int tid)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<P...> >::value>::loadFromSmem(smem, val, tid);
|
||||
}
|
||||
|
||||
template <typename... P, typename... R, class... Op>
|
||||
__device__ __forceinline__ void merge(const thrust::tuple<P...>& smem, const thrust::tuple<R...>& val, unsigned int tid, unsigned int delta, const thrust::tuple<Op...>& op)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<P...> >::value>::merge(smem, val, tid, delta, op);
|
||||
}
|
||||
|
||||
template <typename... R, class... Op>
|
||||
__device__ __forceinline__ void mergeShfl(const thrust::tuple<R...>& val, unsigned int delta, unsigned int width, const thrust::tuple<Op...>& op)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<R...> >::value>::mergeShfl(val, delta, width, op);
|
||||
}
|
||||
#endif
|
||||
template <unsigned int N> struct Generic
|
||||
{
|
||||
template <typename Pointer, typename Reference, class Op>
|
||||
|
||||
@@ -177,6 +177,8 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
data = smem[tid];
|
||||
}
|
||||
|
||||
#if (CUDART_VERSION < 12040)
|
||||
template <typename VP0, typename VP1, typename VP2, typename VP3, typename VP4, typename VP5, typename VP6, typename VP7, typename VP8, typename VP9,
|
||||
typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9>
|
||||
__device__ __forceinline__ void loadToSmem(const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& smem,
|
||||
@@ -193,9 +195,18 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9> >::value>::loadFromSmem(smem, data, tid);
|
||||
}
|
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// copyVals
|
||||
#else
|
||||
template <typename... VP, typename... VR>
|
||||
__device__ __forceinline__ void loadToSmem(const thrust::tuple<VP...>& smem, const thrust::tuple<VR...>& data, unsigned int tid)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VP...> >::value>::loadToSmem(smem, data, tid);
|
||||
}
|
||||
template <typename... VP, typename... VR>
|
||||
__device__ __forceinline__ void loadFromSmem(const thrust::tuple<VP...>& smem, const thrust::tuple<VR...>& data, unsigned int tid)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VP...> >::value>::loadFromSmem(smem, data, tid);
|
||||
}
|
||||
#endif
|
||||
|
||||
template <typename V>
|
||||
__device__ __forceinline__ void copyValsShfl(V& val, unsigned int delta, int width)
|
||||
@@ -207,24 +218,6 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
svals[tid] = val = svals[tid + delta];
|
||||
}
|
||||
template <typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9>
|
||||
__device__ __forceinline__ void copyValsShfl(const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val,
|
||||
unsigned int delta,
|
||||
int width)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9> >::value>::copyShfl(val, delta, width);
|
||||
}
|
||||
template <typename VP0, typename VP1, typename VP2, typename VP3, typename VP4, typename VP5, typename VP6, typename VP7, typename VP8, typename VP9,
|
||||
typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9>
|
||||
__device__ __forceinline__ void copyVals(const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals,
|
||||
const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val,
|
||||
unsigned int tid, unsigned int delta)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9> >::value>::copy(svals, val, tid, delta);
|
||||
}
|
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// merge
|
||||
|
||||
template <typename K, typename V, class Cmp>
|
||||
__device__ __forceinline__ void mergeShfl(K& key, V& val, const Cmp& cmp, unsigned int delta, int width)
|
||||
@@ -248,6 +241,24 @@ namespace cv { namespace cuda { namespace device
|
||||
copyVals(svals, val, tid, delta);
|
||||
}
|
||||
}
|
||||
|
||||
#if (CUDART_VERSION < 12040) // details: https://github.com/opencv/opencv_contrib/issues/3690
|
||||
template <typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9>
|
||||
__device__ __forceinline__ void copyValsShfl(const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val,
|
||||
unsigned int delta,
|
||||
int width)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9> >::value>::copyShfl(val, delta, width);
|
||||
}
|
||||
template <typename VP0, typename VP1, typename VP2, typename VP3, typename VP4, typename VP5, typename VP6, typename VP7, typename VP8, typename VP9,
|
||||
typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9>
|
||||
__device__ __forceinline__ void copyVals(const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals,
|
||||
const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val,
|
||||
unsigned int tid, unsigned int delta)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9> >::value>::copy(svals, val, tid, delta);
|
||||
}
|
||||
|
||||
template <typename K,
|
||||
typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9,
|
||||
class Cmp>
|
||||
@@ -305,7 +316,61 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9> >::value>::merge(skeys, key, svals, val, cmp, tid, delta);
|
||||
}
|
||||
#else
|
||||
template <typename... VR>
|
||||
__device__ __forceinline__ void copyValsShfl(const thrust::tuple<VR...>& val, unsigned int delta, int width)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VR...> >::value>::copyShfl(val, delta, width);
|
||||
}
|
||||
template <typename... VP, typename... VR>
|
||||
__device__ __forceinline__ void copyVals(const thrust::tuple<VP...>& svals, const thrust::tuple<VR...>& val, unsigned int tid, unsigned int delta)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VP...> >::value>::copy(svals, val, tid, delta);
|
||||
}
|
||||
|
||||
template <typename K, typename... VR, class Cmp>
|
||||
__device__ __forceinline__ void mergeShfl(K& key, const thrust::tuple<VR...>& val, const Cmp& cmp, unsigned int delta, int width)
|
||||
{
|
||||
K reg = shfl_down(key, delta, width);
|
||||
|
||||
if (cmp(reg, key))
|
||||
{
|
||||
key = reg;
|
||||
copyValsShfl(val, delta, width);
|
||||
}
|
||||
}
|
||||
template <typename K, typename... VP, typename... VR, class Cmp>
|
||||
__device__ __forceinline__ void merge(volatile K* skeys, K& key, const thrust::tuple<VP...>& svals,
|
||||
const thrust::tuple<VR...>& val, const Cmp& cmp, unsigned int tid, unsigned int delta)
|
||||
{
|
||||
K reg = skeys[tid + delta];
|
||||
|
||||
if (cmp(reg, key))
|
||||
{
|
||||
skeys[tid] = key = reg;
|
||||
copyVals(svals, val, tid, delta);
|
||||
}
|
||||
}
|
||||
template <typename... KR, typename... VR, class... Cmp>
|
||||
__device__ __forceinline__ void mergeShfl(const thrust::tuple<KR...>& key,
|
||||
const thrust::tuple<VR...>& val,
|
||||
const thrust::tuple<Cmp...>& cmp,
|
||||
unsigned int delta, int width)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<KR...> >::value>::mergeShfl(key, val, cmp, delta, width);
|
||||
}
|
||||
template <typename... KP, typename... KR, typename... VP, typename... VR, class... Cmp>
|
||||
__device__ __forceinline__ void merge(const thrust::tuple<KP...>& skeys,
|
||||
const thrust::tuple<KR...>& key,
|
||||
const thrust::tuple<VP...>& svals,
|
||||
const thrust::tuple<VR...>& val,
|
||||
const thrust::tuple<Cmp...>& cmp,
|
||||
unsigned int tid, unsigned int delta)
|
||||
{
|
||||
For<0, thrust::tuple_size<thrust::tuple<VP...> >::value>::merge(skeys, key, svals, val, cmp, tid, delta);
|
||||
}
|
||||
|
||||
#endif
|
||||
//////////////////////////////////////////////////////
|
||||
// Generic
|
||||
|
||||
|
||||
@@ -64,6 +64,12 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
reduce_detail::Dispatcher<N>::reductor::template reduce<volatile T*, T&, const Op&>(smem, val, tid, op);
|
||||
}
|
||||
template <unsigned int N, typename K, typename V, class Cmp>
|
||||
__device__ __forceinline__ void reduceKeyVal(volatile K* skeys, K& key, volatile V* svals, V& val, unsigned int tid, const Cmp& cmp)
|
||||
{
|
||||
reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, volatile V*, V&, const Cmp&>(skeys, key, svals, val, tid, cmp);
|
||||
}
|
||||
#if (CUDART_VERSION < 12040) // details: https://github.com/opencv/opencv_contrib/issues/3690
|
||||
template <int N,
|
||||
typename P0, typename P1, typename P2, typename P3, typename P4, typename P5, typename P6, typename P7, typename P8, typename P9,
|
||||
typename R0, typename R1, typename R2, typename R3, typename R4, typename R5, typename R6, typename R7, typename R8, typename R9,
|
||||
@@ -79,11 +85,6 @@ namespace cv { namespace cuda { namespace device
|
||||
const thrust::tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>&>(smem, val, tid, op);
|
||||
}
|
||||
|
||||
template <unsigned int N, typename K, typename V, class Cmp>
|
||||
__device__ __forceinline__ void reduceKeyVal(volatile K* skeys, K& key, volatile V* svals, V& val, unsigned int tid, const Cmp& cmp)
|
||||
{
|
||||
reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, volatile V*, V&, const Cmp&>(skeys, key, svals, val, tid, cmp);
|
||||
}
|
||||
template <unsigned int N,
|
||||
typename K,
|
||||
typename VP0, typename VP1, typename VP2, typename VP3, typename VP4, typename VP5, typename VP6, typename VP7, typename VP8, typename VP9,
|
||||
@@ -99,6 +100,7 @@ namespace cv { namespace cuda { namespace device
|
||||
const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&,
|
||||
const Cmp&>(skeys, key, svals, val, tid, cmp);
|
||||
}
|
||||
|
||||
template <unsigned int N,
|
||||
typename KP0, typename KP1, typename KP2, typename KP3, typename KP4, typename KP5, typename KP6, typename KP7, typename KP8, typename KP9,
|
||||
typename KR0, typename KR1, typename KR2, typename KR3, typename KR4, typename KR5, typename KR6, typename KR7, typename KR8, typename KR9,
|
||||
@@ -120,6 +122,25 @@ namespace cv { namespace cuda { namespace device
|
||||
const thrust::tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>&
|
||||
>(skeys, key, svals, val, tid, cmp);
|
||||
}
|
||||
#else
|
||||
template <int N, typename... P, typename... R, class... Op>
|
||||
__device__ __forceinline__ void reduce(const thrust::tuple<P...>& smem, const thrust::tuple<R...>& val, unsigned int tid, const thrust::tuple<Op...>& op)
|
||||
{
|
||||
reduce_detail::Dispatcher<N>::reductor::template reduce<const thrust::tuple<P...>&, const thrust::tuple<R...>&, const thrust::tuple<Op...>&>(smem, val, tid, op);
|
||||
}
|
||||
|
||||
template <unsigned int N, typename K, typename... VP, typename... VR, class Cmp>
|
||||
__device__ __forceinline__ void reduceKeyVal(volatile K* skeys, K& key, const thrust::tuple<VP...>& svals, const thrust::tuple<VR...>& val, unsigned int tid, const Cmp& cmp)
|
||||
{
|
||||
reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, const thrust::tuple<VP...>&, const thrust::tuple<VR...>&, const Cmp&>(skeys, key, svals, val, tid, cmp);
|
||||
}
|
||||
|
||||
template <unsigned int N, typename... KP, typename... KR, typename... VP, typename... VR, class... Cmp>
|
||||
__device__ __forceinline__ void reduceKeyVal(const thrust::tuple<KP...>& skeys, const thrust::tuple<KR...>& key, const thrust::tuple<VP...>& svals, const thrust::tuple<VR...>& val, unsigned int tid, const thrust::tuple<Cmp...>& cmp)
|
||||
{
|
||||
reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<const thrust::tuple<KP...>&, const thrust::tuple<KR...>&, const thrust::tuple<VP...>&, const thrust::tuple<VR...>&, const thrust::tuple<Cmp...>&>(skeys, key, svals, val, tid, cmp);
|
||||
}
|
||||
#endif
|
||||
|
||||
// smem_tuple
|
||||
|
||||
|
||||
@@ -201,6 +201,14 @@ namespace cv {
|
||||
# define CV_ICC __INTEL_COMPILER
|
||||
#endif
|
||||
|
||||
#if defined _WIN32
|
||||
# define CV_CDECL __cdecl
|
||||
# define CV_STDCALL __stdcall
|
||||
#else
|
||||
# define CV_CDECL
|
||||
# define CV_STDCALL
|
||||
#endif
|
||||
|
||||
#ifndef CV_INLINE
|
||||
# if defined __cplusplus
|
||||
# define CV_INLINE static inline
|
||||
@@ -483,6 +491,7 @@ Cv64suf;
|
||||
* Matrix type (Mat) *
|
||||
\****************************************************************************************/
|
||||
|
||||
#define CV_MAX_DIM 32
|
||||
#define CV_MAT_CN_MASK ((CV_CN_MAX - 1) << CV_CN_SHIFT)
|
||||
#define CV_MAT_CN(flags) ((((flags) & CV_MAT_CN_MASK) >> CV_CN_SHIFT) + 1)
|
||||
#define CV_MAT_TYPE_MASK (CV_DEPTH_MAX*CV_CN_MAX - 1)
|
||||
@@ -513,6 +522,13 @@ Cv64suf;
|
||||
# define MAX(a,b) ((a) < (b) ? (b) : (a))
|
||||
#endif
|
||||
|
||||
/** min & max without jumps */
|
||||
#define CV_IMIN(a, b) ((a) ^ (((a)^(b)) & (((a) < (b)) - 1)))
|
||||
#define CV_IMAX(a, b) ((a) ^ (((a)^(b)) & (((a) > (b)) - 1)))
|
||||
#define CV_SWAP(a,b,t) ((t) = (a), (a) = (b), (b) = (t))
|
||||
#define CV_CMP(a,b) (((a) > (b)) - ((a) < (b)))
|
||||
#define CV_SIGN(a) CV_CMP((a),0)
|
||||
|
||||
///////////////////////////////////////// Enum operators ///////////////////////////////////////
|
||||
|
||||
/**
|
||||
|
||||
@@ -200,9 +200,14 @@ inline static vuint32mf2_t vmul(const vuint32mf2_t & op1, uint32_t op2, size_t v
|
||||
return vmul_vx_u32mf2(op1, op2, vl);
|
||||
}
|
||||
|
||||
inline static vuint32mf2_t vreinterpret_u32mf2(vint32mf2_t val)
|
||||
inline static vuint32mf2_t vreinterpret_u32mf2(const vint32mf2_t& val)
|
||||
{
|
||||
return vreinterpret_v_i32mf2_u32mf2(val);
|
||||
}
|
||||
|
||||
inline static vuint32mf2_t vreinterpret_u32mf2(const vuint16mf2_t& val)
|
||||
{
|
||||
return vreinterpret_v_u16mf2_u32mf2(val);
|
||||
}
|
||||
|
||||
#endif //OPENCV_HAL_INTRIN_RVV_COMPAT_OVERLOAD_HPP
|
||||
|
||||
@@ -1528,6 +1528,26 @@ OPENCV_HAL_IMPL_RVV_ZIP(v_uint32, vuint32m2_t, u32, 32, 64, OPENCV_HAL_NOP, OPEN
|
||||
OPENCV_HAL_IMPL_RVV_ZIP(v_int32, vint32m2_t, i32, 32, 64, vreinterpret_u32m2, vreinterpret_u32m1)
|
||||
OPENCV_HAL_IMPL_RVV_ZIP(v_float32, vfloat32m2_t, f32, 32, 64, vreinterpret_u32m2, vreinterpret_u32m1)
|
||||
|
||||
#if CV_SIMD_SCALABLE_64F
|
||||
inline void v_zip(const v_float64& a0, const v_float64& a1, v_float64& b0, v_float64& b1) { \
|
||||
vuint16mf4_t idx0 = vid_v_u16mf4(VTraits<v_float64>::vlanes());
|
||||
vuint16mf4_t idx1 = vadd(idx0, VTraits<v_float64>::vlanes(), VTraits<v_float64>::vlanes());
|
||||
vuint16mf2_t idx = vreinterpret_u16mf2(( \
|
||||
vor(vzext_vf2(idx0, VTraits<v_float64>::vlanes()), \
|
||||
vreinterpret_u32mf2(vslide1up(vreinterpret_u16mf2(vzext_vf2(idx1, VTraits<v_float64>::vlanes())), 0, VTraits<v_uint32>::vlanes())), \
|
||||
VTraits<v_uint32>::vlanes())));
|
||||
#if 0
|
||||
vfloat64m2_t temp = __riscv_vcreate_v_f64m1_f64m2(a0, a1);
|
||||
#else // TODO: clean up when RVV Intrinsic is frozen.
|
||||
vfloat64m2_t temp = vlmul_ext_f64m2(a0);
|
||||
temp = vset(temp, 1, a1);
|
||||
#endif
|
||||
temp = vrgatherei16(temp, idx, VTraits<v_float64>::vlanes()*2);
|
||||
b0 = vget_f64m1(temp, 0); \
|
||||
b1 = vget_f64m1(temp, 1); \
|
||||
}
|
||||
#endif
|
||||
|
||||
#define OPENCV_HAL_IMPL_RVV_UNPACKS(_Tpvec, width) \
|
||||
inline _Tpvec v_combine_low(const _Tpvec& a, const _Tpvec& b) \
|
||||
{ \
|
||||
@@ -1859,12 +1879,14 @@ inline v_int32 v_trunc(const v_float32& a)
|
||||
#if CV_SIMD_SCALABLE_64F
|
||||
inline v_int32 v_round(const v_float64& a)
|
||||
{
|
||||
return vfncvt_x(vlmul_ext_f64m2(vfadd(a, 1e-6, VTraits<v_float64>::vlanes())), VTraits<v_float32>::vlanes());
|
||||
return vfncvt_x(vlmul_ext_f64m2(a), VTraits<v_float32>::vlanes());
|
||||
}
|
||||
|
||||
inline v_int32 v_round(const v_float64& a, const v_float64& b)
|
||||
{
|
||||
return vfncvt_x(vset(vlmul_ext_f64m2(vfadd(a, 1e-6, VTraits<v_float64>::vlanes())), 1, b), VTraits<v_float32>::vlanes());
|
||||
// return vfncvt_x(vset(vlmul_ext_f64m2(vfadd(a, 1e-6, VTraits<v_float64>::vlanes())), 1, b), VTraits<v_float32>::vlanes());
|
||||
// Fix https://github.com/opencv/opencv/issues/24746
|
||||
return vfncvt_x(vset(vlmul_ext_f64m2(a), 1, b), VTraits<v_float32>::vlanes());
|
||||
}
|
||||
|
||||
inline v_int32 v_floor(const v_float64& a)
|
||||
|
||||
@@ -90,13 +90,7 @@
|
||||
#include <float.h>
|
||||
#endif // SKIP_INCLUDES
|
||||
|
||||
#if defined _WIN32
|
||||
# define CV_CDECL __cdecl
|
||||
# define CV_STDCALL __stdcall
|
||||
#else
|
||||
# define CV_CDECL
|
||||
# define CV_STDCALL
|
||||
#endif
|
||||
|
||||
|
||||
#ifndef CV_DEFAULT
|
||||
# ifdef __cplusplus
|
||||
@@ -203,21 +197,13 @@ enum {
|
||||
* Common macros and inline functions *
|
||||
\****************************************************************************************/
|
||||
|
||||
#define CV_SWAP(a,b,t) ((t) = (a), (a) = (b), (b) = (t))
|
||||
|
||||
/** min & max without jumps */
|
||||
#define CV_IMIN(a, b) ((a) ^ (((a)^(b)) & (((a) < (b)) - 1)))
|
||||
|
||||
#define CV_IMAX(a, b) ((a) ^ (((a)^(b)) & (((a) > (b)) - 1)))
|
||||
|
||||
/** absolute value without jumps */
|
||||
#ifndef __cplusplus
|
||||
# define CV_IABS(a) (((a) ^ ((a) < 0 ? -1 : 0)) - ((a) < 0 ? -1 : 0))
|
||||
#else
|
||||
# define CV_IABS(a) abs(a)
|
||||
#endif
|
||||
#define CV_CMP(a,b) (((a) > (b)) - ((a) < (b)))
|
||||
#define CV_SIGN(a) CV_CMP((a),0)
|
||||
|
||||
|
||||
#define cvInvSqrt(value) ((float)(1./sqrt(value)))
|
||||
#define cvSqrt(value) ((float)sqrt(value))
|
||||
@@ -675,8 +661,6 @@ CV_INLINE int cvIplDepth( int type )
|
||||
#define CV_MATND_MAGIC_VAL 0x42430000
|
||||
#define CV_TYPE_NAME_MATND "opencv-nd-matrix"
|
||||
|
||||
#define CV_MAX_DIM 32
|
||||
|
||||
#ifdef __cplusplus
|
||||
typedef struct CvMatND CvMatND;
|
||||
CV_EXPORTS CvMatND cvMatND(const cv::Mat& m);
|
||||
|
||||
Reference in New Issue
Block a user