static __global__ void transformSmart(const PtrStepSz<T1> src1_, const PtrStep<T2> src2_, PtrStep<D> dst_, const Mask mask, const BinOp op) { typedef TransformFunctorTraits<BinOp> ft; typedef typename BinaryReadWriteTraits<T1, T2, D, ft::smart_shift>::read_type1 read_type1; typedef typename BinaryReadWriteTraits<T1, T2, D, ft::smart_shift>::read_type2 read_type2; typedef typename BinaryReadWriteTraits<T1, T2, D, ft::smart_shift>::write_type write_type; const int x = threadIdx.x + blockIdx.x * blockDim.x; const int y = threadIdx.y + blockIdx.y * blockDim.y; const int x_shifted = x * ft::smart_shift; if (y < src1_.rows) { const T1* src1 = src1_.ptr(y); const T2* src2 = src2_.ptr(y); D* dst = dst_.ptr(y); if (x_shifted + ft::smart_shift - 1 < src1_.cols) { const read_type1 src1_n_el = ((const read_type1*)src1)[x]; const read_type2 src2_n_el = ((const read_type2*)src2)[x]; OpUnroller<ft::smart_shift>::unroll(src1_n_el, src2_n_el, ((write_type*)dst)[x], mask, op, x_shifted, y); } else { for (int real_x = x_shifted; real_x < src1_.cols; ++real_x) { if (mask(y, real_x)) dst[real_x] = op(src1[real_x], src2[real_x]); } } } }
static __global__ void transformSimple(const PtrStepSz<T1> src1, const PtrStep<T2> src2, PtrStep<D> dst, const Mask mask, const BinOp op) { const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; if (x < src1.cols && y < src1.rows && mask(y, x)) { const T1 src1_data = src1.ptr(y)[x]; const T2 src2_data = src2.ptr(y)[x]; dst.ptr(y)[x] = op(src1_data, src2_data); } }
static __global__ void transformSmart(const PtrStepSz<T> src_, PtrStep<D> dst_, const Mask mask, const UnOp op) { typedef TransformFunctorTraits<UnOp> ft; typedef typename UnaryReadWriteTraits<T, D, ft::smart_shift>::read_type read_type; typedef typename UnaryReadWriteTraits<T, D, ft::smart_shift>::write_type write_type; const int x = threadIdx.x + blockIdx.x * blockDim.x; const int y = threadIdx.y + blockIdx.y * blockDim.y; const int x_shifted = x * ft::smart_shift; if (y < src_.rows) { const T* src = src_.ptr(y); D* dst = dst_.ptr(y); if (x_shifted + ft::smart_shift - 1 < src_.cols) { const read_type src_n_el = ((const read_type*)src)[x]; write_type dst_n_el = ((const write_type*)dst)[x]; OpUnroller<ft::smart_shift>::unroll(src_n_el, dst_n_el, mask, op, x_shifted, y); ((write_type*)dst)[x] = dst_n_el; } else { for (int real_x = x_shifted; real_x < src_.cols; ++real_x) { if (mask(y, real_x)) dst[real_x] = op(src[real_x]); } } } }
__global__ static void transformSimple(const PtrStepSz<T> src, PtrStep<D> dst, const Mask mask, const UnOp op) { const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; if (x < src.cols && y < src.rows && mask(y, x)) { dst.ptr(y)[x] = op(src.ptr(y)[x]); } }