Skip to content

Commit 752c01c

Browse files
asmorkalovasmorkalov
committed
Got rid of tuple in polar-cart operations.
1 parent 041b84d commit 752c01c

File tree

3 files changed

+213
-18
lines changed

3 files changed

+213
-18
lines changed

‎modules/cudaarithm/src/cuda/polar_cart.cu

Lines changed: 4 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -132,24 +132,10 @@ void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, Outpu
132132
GpuMat_<float> magc(mag.reshape(1));
133133
GpuMat_<float> anglec(angle.reshape(1));
134134

135-
// if (angleInDegrees)
136-
// {
137-
// gridTransformTuple(zipPtr(xc, yc),
138-
// tie(magc, anglec),
139-
// make_tuple(
140-
// binaryTupleAdapter<0, 1>(magnitude_func<float>()),
141-
// binaryTupleAdapter<0, 1>(direction_func<float, true>())),
142-
// stream);
143-
// }
144-
// else
145-
// {
146-
// gridTransformTuple(zipPtr(xc, yc),
147-
// tie(magc, anglec),
148-
// make_tuple(
149-
// binaryTupleAdapter<0, 1>(magnitude_func<float>()),
150-
// binaryTupleAdapter<0, 1>(direction_func<float, false>())),
151-
// stream);
152-
// }
135+
if (angleInDegrees)
136+
gridTransformBinary(xc, yc, magc, anglec, magnitude_func<float>(), direction_func<float, true>(), stream);
137+
else
138+
gridTransformBinary(xc, yc, magc, anglec, magnitude_func<float>(), direction_func<float, false>(), stream);
153139

154140
syncOutput(mag, _mag, stream);
155141
syncOutput(angle, _angle, stream);

‎modules/cudev/include/opencv2/cudev/grid/detail/transform.hpp

Lines changed: 118 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -179,6 +179,22 @@ namespace grid_transform_detail
179179
dst(y, x) = saturate_cast<DstType>(op(src1(y, x), src2(y, x)));
180180
}
181181

182+
// transformSimple, 2 outputs
183+
184+
template <class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class BinOp1, class BinOp2, class MaskPtr>
185+
__global__ void transformSimple(const SrcPtr1 src1, const SrcPtr2 src2, GlobPtr<DstType1> dst1, GlobPtr<DstType2> dst2,
186+
const BinOp1 op1, const BinOp2 op2, const MaskPtr mask, const int rows, const int cols)
187+
{
188+
const int x = blockIdx.x * blockDim.x + threadIdx.x;
189+
const int y = blockIdx.y * blockDim.y + threadIdx.y;
190+
191+
if (x >= cols || y >= rows || !mask(y, x))
192+
return;
193+
194+
dst1(y, x) = saturate_cast<DstType1>(op1(src1(y, x), src2(y, x)));
195+
dst2(y, x) = saturate_cast<DstType2>(op2(src1(y, x), src2(y, x)));
196+
}
197+
182198
// transformSmart
183199

184200
template <int SHIFT, typename SrcType, typename DstType, class UnOp, class MaskPtr>
@@ -248,6 +264,51 @@ namespace grid_transform_detail
248264
}
249265
}
250266

267+
// transformSmart, 2 outputs
268+
269+
template <int SHIFT, typename SrcType1, typename SrcType2, typename DstType1, typename DstType2, class BinOp1, class BinOp2, class MaskPtr>
270+
__global__ void transformSmart(const GlobPtr<SrcType1> src1_, const GlobPtr<SrcType2> src2_,
271+
GlobPtr<DstType1> dst1_, GlobPtr<DstType2> dst2_,
272+
const BinOp1 op1, const BinOp2 op2, const MaskPtr mask, const int rows, const int cols)
273+
{
274+
typedef typename MakeVec<SrcType1, SHIFT>::type read_type1;
275+
typedef typename MakeVec<SrcType2, SHIFT>::type read_type2;
276+
typedef typename MakeVec<DstType1, SHIFT>::type write_type1;
277+
typedef typename MakeVec<DstType2, SHIFT>::type write_type2;
278+
279+
const int x = blockIdx.x * blockDim.x + threadIdx.x;
280+
const int y = blockIdx.y * blockDim.y + threadIdx.y;
281+
const int x_shifted = x * SHIFT;
282+
283+
if (y < rows)
284+
{
285+
const SrcType1* src1 = src1_.row(y);
286+
const SrcType2* src2 = src2_.row(y);
287+
DstType1* dst1 = dst1_.row(y);
288+
DstType2* dst2 = dst2_.row(y);
289+
290+
if (x_shifted + SHIFT - 1 < cols)
291+
{
292+
const read_type1 src1_n_el = ((const read_type1*)src1)[x];
293+
const read_type2 src2_n_el = ((const read_type2*)src2)[x];
294+
295+
OpUnroller<SHIFT>::unroll(src1_n_el, src2_n_el, ((write_type1*)dst1)[x], op1, mask, x_shifted, y);
296+
OpUnroller<SHIFT>::unroll(src1_n_el, src2_n_el, ((write_type2*)dst2)[x], op2, mask, x_shifted, y);
297+
}
298+
else
299+
{
300+
for (int real_x = x_shifted; real_x < cols; ++real_x)
301+
{
302+
if (mask(y, real_x))
303+
{
304+
dst1[real_x] = op1(src1[real_x], src2[real_x]);
305+
dst2[real_x] = op2(src1[real_x], src2[real_x]);
306+
}
307+
}
308+
}
309+
}
310+
}
311+
251312
// TransformDispatcher
252313

253314
template <bool UseSmart, class Policy> struct TransformDispatcher;
@@ -279,6 +340,20 @@ namespace grid_transform_detail
279340
if (stream == 0)
280341
CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
281342
}
343+
344+
template <class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class BinOp1, class BinOp2, class MaskPtr>
345+
__host__ static void call(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtr<DstType1>& dst1, const GlobPtr<DstType2>& dst2,
346+
const BinOp1& op1, const BinOp2& op2, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
347+
{
348+
const dim3 block(Policy::block_size_x, Policy::block_size_y);
349+
const dim3 grid(divUp(cols, block.x), divUp(rows, block.y));
350+
351+
transformSimple<<<grid, block, 0, stream>>>(src1, src2, dst1, dst2, op1, op2, mask, rows, cols);
352+
CV_CUDEV_SAFE_CALL( cudaGetLastError() );
353+
354+
if (stream == 0)
355+
CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
356+
}
282357
};
283358

284359
template <class Policy> struct TransformDispatcher<true, Policy>
@@ -336,6 +411,33 @@ namespace grid_transform_detail
336411
if (stream == 0)
337412
CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
338413
}
414+
415+
template <typename SrcType1, typename SrcType2, typename DstType1, typename DstType2, class BinOp1, class BinOp2, class MaskPtr>
416+
__host__ static void call(const GlobPtr<SrcType1>& src1, const GlobPtr<SrcType2>& src2,
417+
const GlobPtr<DstType1>& dst1, const GlobPtr<DstType2>& dst2,
418+
const BinOp1& op1, const BinOp2& op2, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
419+
{
420+
if (Policy::shift == 1 ||
421+
!isAligned(src1.data, Policy::shift * sizeof(SrcType1)) || !isAligned(src1.step, Policy::shift * sizeof(SrcType1)) ||
422+
!isAligned(src2.data, Policy::shift * sizeof(SrcType2)) || !isAligned(src2.step, Policy::shift * sizeof(SrcType2)) ||
423+
!isAligned(dst1.data, Policy::shift * sizeof(DstType1)) || !isAligned(dst1.step, Policy::shift * sizeof(DstType1))||
424+
!isAligned(dst2.data, Policy::shift * sizeof(DstType2)) || !isAligned(dst2.step, Policy::shift * sizeof(DstType2))
425+
)
426+
{
427+
TransformDispatcher<false, Policy>::call(src1, src2, dst1, dst2, op1, op2, mask, rows, cols, stream);
428+
return;
429+
}
430+
431+
const dim3 block(Policy::block_size_x, Policy::block_size_y);
432+
const dim3 grid(divUp(cols, block.x * Policy::shift), divUp(rows, block.y));
433+
434+
transformSmart<Policy::shift><<<grid, block, 0, stream>>>(src1, src2, dst1, dst2, op1, op2, mask, rows, cols);
435+
CV_CUDEV_SAFE_CALL( cudaGetLastError() );
436+
437+
if (stream == 0)
438+
CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
439+
}
440+
339441
};
340442

341443
template <class Policy, class SrcPtr, typename DstType, class UnOp, class MaskPtr>
@@ -350,6 +452,13 @@ namespace grid_transform_detail
350452
TransformDispatcher<false, Policy>::call(src1, src2, dst, op, mask, rows, cols, stream);
351453
}
352454

455+
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class BinOp1, class BinOp2, class MaskPtr>
456+
__host__ void transform_binary(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtr<DstType1>& dst1, const GlobPtr<DstType2>& dst2,
457+
const BinOp1& op1, const BinOp2& op2, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
458+
{
459+
TransformDispatcher<false, Policy>::call(src1, src2, dst1, dst2, op1, op2, mask, rows, cols, stream);
460+
}
461+
353462
template <class Policy, typename SrcType, typename DstType, class UnOp, class MaskPtr>
354463
__host__ void transform_unary(const GlobPtr<SrcType>& src, const GlobPtr<DstType>& dst, const UnOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
355464
{
@@ -362,6 +471,15 @@ namespace grid_transform_detail
362471
TransformDispatcher<VecTraits<SrcType1>::cn == 1 && VecTraits<SrcType2>::cn == 1 && VecTraits<DstType>::cn == 1 && Policy::shift != 1, Policy>::call(src1, src2, dst, op, mask, rows, cols, stream);
363472
}
364473

474+
template <class Policy, typename SrcType1, typename SrcType2, typename DstType1, typename DstType2, class BinOp1, class BinOp2, class MaskPtr>
475+
__host__ void transform_binary(const GlobPtr<SrcType1>& src1, const GlobPtr<SrcType2>& src2, const GlobPtr<DstType1>& dst1, const GlobPtr<DstType2>& dst2,
476+
const BinOp1& op1, const BinOp2& op2, const MaskPtr& mask, int rows, int cols, cudaStream_t stream)
477+
{
478+
TransformDispatcher<VecTraits<SrcType1>::cn == 1 && VecTraits<SrcType2>::cn == 1 &&
479+
VecTraits<DstType1>::cn == 1 && VecTraits<DstType2>::cn == 1 &&
480+
Policy::shift != 1, Policy>::call(src1, src2, dst1, dst2, op1, op2, mask, rows, cols, stream);
481+
}
482+
365483
// transform_tuple
366484

367485
template <int count> struct Unroll

‎modules/cudev/include/opencv2/cudev/grid/transform.hpp

Lines changed: 91 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -121,6 +121,22 @@ __host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, Gpu
121121
grid_transform_detail::transform_binary<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
122122
}
123123

124+
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class BinOp1, class BinOp2, class MaskPtr>
125+
__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<DstType1>& dst1, GpuMat_<DstType2>& dst2,
126+
const BinOp1& op1, const BinOp2& op2, const MaskPtr& mask, Stream& stream = Stream::Null())
127+
{
128+
const int rows = getRows(src1);
129+
const int cols = getCols(src1);
130+
131+
CV_Assert( getRows(src2) == rows && getCols(src2) == cols );
132+
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
133+
134+
dst1.create(rows, cols);
135+
dst2.create(rows, cols);
136+
137+
grid_transform_detail::transform_binary<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst1), shrinkPtr(dst2), op1, op2, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
138+
}
139+
124140
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType, class BinOp, class MaskPtr>
125141
__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz<DstType>& dst, const BinOp& op, const MaskPtr& mask, Stream& stream = Stream::Null())
126142
{
@@ -134,6 +150,22 @@ __host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, con
134150
grid_transform_detail::transform_binary<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
135151
}
136152

153+
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class BinOp1, class BinOp2, class MaskPtr>
154+
__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz<DstType1>& dst1, const GlobPtrSz<DstType2>& dst2,
155+
const BinOp1& op1, const BinOp2& op2, const MaskPtr& mask, Stream& stream = Stream::Null())
156+
{
157+
const int rows = getRows(src1);
158+
const int cols = getCols(src1);
159+
160+
CV_Assert( getRows(dst1) == rows && getCols(dst1) == cols );
161+
CV_Assert( getRows(dst2) == rows && getCols(dst2) == cols );
162+
CV_Assert( getRows(src2) == rows && getCols(src2) == cols );
163+
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
164+
165+
grid_transform_detail::transform_binary<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst1), shrinkPtr(dst2), op1, op2, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
166+
}
167+
168+
137169
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType, class BinOp>
138170
__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<DstType>& dst, const BinOp& op, Stream& stream = Stream::Null())
139171
{
@@ -147,6 +179,21 @@ __host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, Gpu
147179
grid_transform_detail::transform_binary<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
148180
}
149181

182+
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class BinOp1, class BinOp2>
183+
__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<DstType1>& dst1, GpuMat_<DstType2>& dst2,
184+
const BinOp1& op1, const BinOp2& op2, Stream& stream = Stream::Null())
185+
{
186+
const int rows = getRows(src1);
187+
const int cols = getCols(src1);
188+
189+
CV_Assert( getRows(src2) == rows && getCols(src2) == cols );
190+
191+
dst1.create(rows, cols);
192+
dst2.create(rows, cols);
193+
194+
grid_transform_detail::transform_binary<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst1), shrinkPtr(dst2), op1, op2, WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
195+
}
196+
150197
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType, class BinOp>
151198
__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz<DstType>& dst, const BinOp& op, Stream& stream = Stream::Null())
152199
{
@@ -159,6 +206,20 @@ __host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, con
159206
grid_transform_detail::transform_binary<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
160207
}
161208

209+
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class BinOp1, class BinOp2>
210+
__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz<DstType1>& dst1, const GlobPtrSz<DstType2>& dst2,
211+
const BinOp1& op1, const BinOp2& op2, Stream& stream = Stream::Null())
212+
{
213+
const int rows = getRows(src1);
214+
const int cols = getCols(src1);
215+
216+
CV_Assert( getRows(dst1) == rows && getCols(dst1) == cols );
217+
CV_Assert( getRows(dst2) == rows && getCols(dst2) == cols );
218+
CV_Assert( getRows(src2) == rows && getCols(src2) == cols );
219+
220+
grid_transform_detail::transform_binary<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst1), shrinkPtr(dst2), op1, op2, WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
221+
}
222+
162223
template <class Policy, class SrcPtr, typename D0, typename D1, class OpTuple, class MaskPtr>
163224
__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
164225
{
@@ -449,24 +510,54 @@ __host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, GpuM
449510
gridTransformBinary_<DefaultTransformPolicy>(src1, src2, dst, op, mask, stream);
450511
}
451512

513+
template <class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class Op1, class Op2, class MaskPtr>
514+
__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<DstType1>& dst1, GpuMat_<DstType2>& dst2,
515+
const Op1& op1, const Op2& op2, const MaskPtr& mask, Stream& stream = Stream::Null())
516+
{
517+
gridTransformBinary_<DefaultTransformPolicy>(src1, src2, dst1, dst2, op1, op2, mask, stream);
518+
}
519+
452520
template <class SrcPtr1, class SrcPtr2, typename DstType, class Op, class MaskPtr>
453521
__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz<DstType>& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null())
454522
{
455523
gridTransformBinary_<DefaultTransformPolicy>(src1, src2, dst, op, mask, stream);
456524
}
457525

526+
template <class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class Op1, class Op2, class MaskPtr>
527+
__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz<DstType2>& dst1, const GlobPtrSz<DstType2>& dst2,
528+
const Op1& op1, const Op2& op2, const MaskPtr& mask, Stream& stream = Stream::Null())
529+
{
530+
gridTransformBinary_<DefaultTransformPolicy>(src1, src2, dst1, dst2, op1, op2, mask, stream);
531+
}
532+
458533
template <class SrcPtr1, class SrcPtr2, typename DstType, class Op>
459534
__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<DstType>& dst, const Op& op, Stream& stream = Stream::Null())
460535
{
461536
gridTransformBinary_<DefaultTransformPolicy>(src1, src2, dst, op, stream);
462537
}
463538

539+
template <class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class Op1, class Op2>
540+
__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2,
541+
GpuMat_<DstType1>& dst1, GpuMat_<DstType2>& dst2,
542+
const Op1& op1, const Op2& op2, Stream& stream = Stream::Null())
543+
{
544+
gridTransformBinary_<DefaultTransformPolicy>(src1, src2, dst1, dst2, op1, op2, stream);
545+
}
546+
464547
template <class SrcPtr1, class SrcPtr2, typename DstType, class Op>
465548
__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz<DstType>& dst, const Op& op, Stream& stream = Stream::Null())
466549
{
467550
gridTransformBinary_<DefaultTransformPolicy>(src1, src2, dst, op, stream);
468551
}
469552

553+
template <class SrcPtr1, class SrcPtr2, typename DstType1, typename DstType2, class Op1, class Op2>
554+
__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2,
555+
const GlobPtrSz<DstType1>& dst1, const GlobPtrSz<DstType2>& dst2,
556+
const Op1& op1, const Op2& op2, Stream& stream = Stream::Null())
557+
{
558+
gridTransformBinary_<DefaultTransformPolicy>(src1, src2, dst1, dst2, op1, op2, stream);
559+
}
560+
470561
template <class SrcPtr, typename D0, typename D1, class OpTuple, class MaskPtr>
471562
__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
472563
{

0 commit comments

Comments
 (0)