     template <typename T>
     void set_to_gpu(PtrStepSzb mat, const T* scalar, int channels, cudaStream_t stream)
     {
-        writeScalar(scalar);
-
-        dim3 threadsPerBlock(32, 8, 1);
-        dim3 numBlocks (mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1);
-
-        set_to_without_mask<T><<<numBlocks, threadsPerBlock, 0, stream>>>((T*)mat.data, mat.cols, mat.rows, mat.step, channels);
+        typedef typename TypeVec<T, 4>::vec_type vec_type;
+
+        dim3 block(32, 8);
+        dim3 grid(divUp(mat.rows, block.y));
+
+        set_to_without_mask<T><<<grid, block, 0, stream>>>(PtrStepSz<T>(mat), VecTraits<vec_type>::make(scalar), channels);
         cudaSafeCall( cudaGetLastError() );
 
         if (stream == 0)

#include "opencv2/gpu/device/transform.hpp"
#include "opencv2/gpu/device/functional.hpp"
#include "opencv2/gpu/device/type_traits.hpp"

namespace cv { namespace gpu { namespace device
{

    ////////////////////////////////// SetTo //////////////////////////////////
    ///////////////////////////////////////////////////////////////////////////

    template<typename T>
    __global__ void set_to_without_mask(PtrStepSz<T> mat, typename TypeVec<T, 4>::vec_type val, int channels)
    {
        const int y = blockIdx.x * blockDim.y + threadIdx.y;

        if (y < mat.rows)
        {
            const T vals[] = {
                val.x, val.y, val.z, val.w
            };

            T* row = mat.ptr(y);

            for (int x = threadIdx.x; x < mat.cols * channels; x += 32)
            {
                row[x] = vals[x % channels];
            }
        }
    }

    template<typename T>
    __global__ void set_to_with_mask(PtrStepSz<T> mat, const PtrStepb mask, typename TypeVec<T, 4>::vec_type val, int channels)
    {
        const int y = blockIdx.x * blockDim.y + threadIdx.y;

        if (y < mat.rows)
        {
            const T vals[] = {
                val.x, val.y, val.z, val.w
            };

            T* row = mat.ptr(y);
            const uchar* mask_row = mask.ptr(y);

            for (int x = threadIdx.x; x < mat.cols * channels; x += 32)
            {
                if (mask_row[x / channels])
                {
                    row[x] = vals[x % channels];
                }
            }
        }
    }

    template <typename T>
    void set_to_gpu(PtrStepSzb mat, const T* scalar, PtrStepSzb mask, int channels, cudaStream_t stream)
    {
        typedef typename TypeVec<T, 4>::vec_type vec_type;

        dim3 block(32, 8);
        dim3 grid(divUp(mat.rows, block.y));

        set_to_with_mask<T><<<grid, block, 0, stream>>>(PtrStepSz<T>(mat), mask, VecTraits<vec_type>::make(scalar), channels);
        cudaSafeCall( cudaGetLastError() );

        if (stream == 0)
