Added branch compatible with variadic tuple version in Trust (since CUDA 12.4).

Author: asmorkalov

modules/cudev/include/opencv2/cudev/block/detail/reduce.hpp

@@ -154,6 +154,17 @@ namespace block_reduce_detail
         val = smem[tid];
     }
 
+
+    // merge
+
+    template <typename T, class Op>
+    __device__ __forceinline__ void merge(volatile T* smem, T& val, uint tid, uint delta, const Op& op)
+    {
+        T reg = smem[tid + delta];
+        smem[tid] = val = op(val, reg);
+    }
+
+#if (CUDART_VERSION < 12040)
     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 tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>& smem,
@@ -172,15 +183,6 @@ namespace block_reduce_detail
         For<0, tuple_size<tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9> >::value>::loadFromSmem(smem, val, tid);
     }
 
-    // merge
-
-    template <typename T, class Op>
-    __device__ __forceinline__ void merge(volatile T* smem, T& val, uint tid, uint delta, const Op& op)
-    {
-        T reg = smem[tid + delta];
-        smem[tid] = 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>
@@ -214,6 +216,41 @@ namespace block_reduce_detail
     }
 #endif
 
+#else
+    template <typename... P, typename... R>
+    __device__ __forceinline__ void loadToSmem(const tuple<P...>& smem, const tuple<R...>& val, uint tid)
+    {
+        For<0, tuple_size<tuple<P...> >::value>::loadToSmem(smem, val, tid);
+    }
+
+    template <typename... P, typename... R>
+    __device__ __forceinline__ void loadFromSmem(const tuple<P...>& smem, const tuple<R...>& val, uint tid)
+    {
+        For<0, tuple_size<tuple<P...> >::value>::loadFromSmem(smem, val, tid);
+    }
+
+    template <typename P..., typename... R, class... Op>
+    __device__ __forceinline__ void merge(const tuple<P...>& smem, const tuple<R...>& val, uint tid, uint delta, const tuple<Op...>& op)
+    {
+        For<0, tuple_size<tuple<P...> >::value>::merge(smem, val, tid, delta, op);
+    }
+
+    // mergeShfl
+
+    template <typename T, class Op>
+    __device__ __forceinline__ void mergeShfl(T& val, uint delta, uint width, const Op& op)
+    {
+        T reg = shfl_down(val, delta, width);
+        val = op(val, reg);
+    }
+
+    template <typename... R, class... Op>
+    __device__ __forceinline__ void mergeShfl(const tuple<R...>& val, uint delta, uint width, const tuple<Op...>& op)
+    {
+        For<0, tuple_size<tuple<R...> >::value>::mergeShfl(val, delta, width, op);
+    }
+#endif
+
     // Generic
 
     template <int N> struct Generic

modules/cudev/include/opencv2/cudev/block/reduce.hpp

@@ -51,6 +51,7 @@
 #include "../warp/reduce.hpp"
 #include "detail/reduce.hpp"
 #include "detail/reduce_key_val.hpp"
+#include <cuda_runtime_api.h>
 
 namespace cv { namespace cudev {
 
@@ -65,6 +66,7 @@ __device__ __forceinline__ void blockReduce(volatile T* smem, T& val, uint tid,
     block_reduce_detail::Dispatcher<N>::reductor::template reduce<volatile T*, T&, const Op&>(smem, val, tid, op);
 }
 
+#if (CUDART_VERSION < 12040)
 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,
@@ -126,6 +128,39 @@ __device__ __forceinline__ void blockReduceKeyVal(const tuple<KP0, KP1, KP2, KP3
             >(skeys, key, svals, val, tid, cmp);
 }
 
+#else
+
+template <int N, typename... P, typename... R, typename... Op>
+__device__ __forceinline__ void blockReduce(const tuple<P...>& smem,
+                                            const tuple<R...>& val,
+                                            uint tid,
+                                            const tuple<Op..>& op)
+{
+    block_reduce_detail::Dispatcher<N>::reductor::template reduce<const tuple<P...>&, const tuple<R...>&, const tuple<Op...>&>(smem, val, tid, op);
+}
+
+// blockReduceKeyVal
+
+template <int N, typename K, typename V, class Cmp>
+__device__ __forceinline__ void blockReduceKeyVal(volatile K* skeys, K& key, volatile V* svals, V& val, uint tid, const Cmp& cmp)
+{
+    block_reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, volatile V*, V&, const Cmp&>(skeys, key, svals, val, tid, cmp);
+}
+
+template <int N, typename K, typename... VP, typename... VR, class Cmp>
+__device__ __forceinline__ void blockReduceKeyVal(volatile K* skeys, K& key, const tuple<VP...>& svals, const tuple<VR...>& val, uint tid, const Cmp& cmp)
+{
+    block_reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, const tuple<VP...>&, const tuple<VR...>&, const Cmp&>(skeys, key, svals, val, tid, cmp);
+}
+
+template <int N, typename... KP, typename... KR, typename... VP, typename... VR, class... Cmp>
+__device__ __forceinline__ void blockReduceKeyVal(const tuple<KP...>& skeys, const tuple<KR...>& key, const tuple<VP...>& svals, const tuple<VR...>& val, uint tid, const tuple<Cmp...>& cmp)
+{
+    block_reduce_key_val_detail::Dispatcher<N>::reductor::template reduce< const tuple<KP...>&, const tuple<KR...>&, const tuple<VP...>&, const tuple<VR...>&, const tuple<Cmp...>&>(skeys, key, svals, val, tid, cmp);
+}
+
+#endif
+
 //! @}
 
 }}