     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)

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

    __constant__ uchar scalar_8u[4];
    __constant__ schar scalar_8s[4];
    __constant__ ushort scalar_16u[4];
    __constant__ short scalar_16s[4];
    __constant__ int scalar_32s[4];
    __constant__ float scalar_32f[4];
    __constant__ double scalar_64f[4];

    template <typename T> __device__ __forceinline__ T readScalar(int i);
    template <> __device__ __forceinline__ uchar readScalar<uchar>(int i) {return scalar_8u[i];}
    template <> __device__ __forceinline__ schar readScalar<schar>(int i) {return scalar_8s[i];}
    template <> __device__ __forceinline__ ushort readScalar<ushort>(int i) {return scalar_16u[i];}
    template <> __device__ __forceinline__ short readScalar<short>(int i) {return scalar_16s[i];}
    template <> __device__ __forceinline__ int readScalar<int>(int i) {return scalar_32s[i];}
    template <> __device__ __forceinline__ float readScalar<float>(int i) {return scalar_32f[i];}
    template <> __device__ __forceinline__ double readScalar<double>(int i) {return scalar_64f[i];}

    void writeScalar(const uchar* vals)
    {
        cudaSafeCall( cudaMemcpyToSymbol(scalar_8u, vals, sizeof(uchar) * 4) );
    }
    void writeScalar(const schar* vals)
    {
        cudaSafeCall( cudaMemcpyToSymbol(scalar_8s, vals, sizeof(schar) * 4) );
    }
    void writeScalar(const ushort* vals)
    {
        cudaSafeCall( cudaMemcpyToSymbol(scalar_16u, vals, sizeof(ushort) * 4) );
    }
    void writeScalar(const short* vals)
    {
        cudaSafeCall( cudaMemcpyToSymbol(scalar_16s, vals, sizeof(short) * 4) );
    }
    void writeScalar(const int* vals)
    {
        cudaSafeCall( cudaMemcpyToSymbol(scalar_32s, vals, sizeof(int) * 4) );
    }
    void writeScalar(const float* vals)
    {
        cudaSafeCall( cudaMemcpyToSymbol(scalar_32f, vals, sizeof(float) * 4) );
    }
    void writeScalar(const double* vals)
    {
        cudaSafeCall( cudaMemcpyToSymbol(scalar_64f, vals, sizeof(double) * 4) );
    }

    template<typename T>
    __global__ void set_to_without_mask(T* mat, int cols, int rows, size_t step, int channels)
    {
        size_t x = blockIdx.x * blockDim.x + threadIdx.x;
        size_t y = blockIdx.y * blockDim.y + threadIdx.y;

        if ((x < cols * channels ) && (y < rows))
        {
            size_t idx = y * ( step >> shift_and_sizeof<T>::shift ) + x;
            mat[idx] = readScalar<T>(x % channels);
        }
    }

    template<typename T>
    __global__ void set_to_with_mask(T* mat, const uchar* mask, int cols, int rows, size_t step, int channels, size_t step_mask)
    {
        size_t x = blockIdx.x * blockDim.x + threadIdx.x;
        size_t y = blockIdx.y * blockDim.y + threadIdx.y;

        if ((x < cols * channels ) && (y < rows))
            if (mask[y * step_mask + x / channels] != 0)
            {
                size_t idx = y * ( step >> shift_and_sizeof<T>::shift ) + x;
                mat[idx] = readScalar<T>(x % channels);
            }
    }
    template <typename T>
    void set_to_gpu(PtrStepSzb mat, const T* scalar, PtrStepSzb mask, 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_with_mask<T><<<numBlocks, threadsPerBlock, 0, stream>>>((T*)mat.data, (uchar*)mask.data, mat.cols, mat.rows, mat.step, channels, mask.step);
        cudaSafeCall( cudaGetLastError() );

        if (stream == 0)

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

namespace cv { namespace gpu { namespace device
{
