diff --git a/modules/gpu/cuda/mat_operators.cu b/modules/gpu/cuda/mat_operators.cu
index 5d697598d..d1ac7e39a 100644
--- a/modules/gpu/cuda/mat_operators.cu
+++ b/modules/gpu/cuda/mat_operators.cu
@@ -57,28 +57,66 @@ namespace mat_operators
             mat[i] = static_cast<T>(scalar_d[i % channels]);
             unroll<T, channels, count - 1>::unroll_set(mat, i+1);
         }
+
+        __device__ static void unroll_set_with_mask(T * mat, float mask, size_t i)
+        {
+            mat[i] = mask * static_cast<T>(scalar_d[i % channels]);
+            unroll<T, channels, count - 1>::unroll_set_with_mask(mat, mask, i+1);
+        }
     };
 
     template <typename T, int channels>
     struct unroll<T,channels,0>
     {
         __device__ static void unroll_set(T * , size_t){}
+        __device__ static void unroll_set_with_mask(T * , float, size_t){}
     };
 
-
-
     template <typename T, int channels>
     __global__ void kernel_set_to_without_mask(T * mat)
     {
         size_t i = (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(T);
         unroll<T, channels>::unroll_set(mat, i);
     }
+
+    template <typename T, int channels>
+    __global__ void kernel_set_to_with_mask(T * mat, const float * mask)
+    {
+        size_t i = (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(T);
+        unroll<T, channels>::unroll_set_with_mask(mat, i, mask[i]);
+    }
 }
 
 
 extern "C" void cv::gpu::impl::set_to_with_mask(const DevMem2D& mat, const double * scalar, const DevMem2D& mask, int depth, int channels)
 {
+    scalar_d[0] = scalar[0];
+    scalar_d[1] = scalar[1];
+    scalar_d[2] = scalar[2];
+    scalar_d[3] = scalar[3];
 
+    int numBlocks = mat.rows * mat.step / 256;
+
+    dim3 threadsPerBlock(256);
+
+    if (channels == 1)
+    {
+        if (depth == 1) ::mat_operators::kernel_set_to_with_mask<unsigned char,  1><<<numBlocks,threadsPerBlock>>>(mat.ptr, (float *)mask.ptr);
+        if (depth == 2) ::mat_operators::kernel_set_to_with_mask<unsigned short, 1><<<numBlocks,threadsPerBlock>>>((unsigned short *)mat.ptr, (float *)mask.ptr);
+        if (depth == 4) ::mat_operators::kernel_set_to_with_mask<unsigned int,   1><<<numBlocks,threadsPerBlock>>>((unsigned int *)mat.ptr, (float *)mask.ptr);
+    }
+    if (channels == 2)
+    {
+        if (depth == 1) ::mat_operators::kernel_set_to_with_mask<unsigned char,  2><<<numBlocks,threadsPerBlock>>>(mat.ptr, (float *)mask.ptr);
+        if (depth == 2) ::mat_operators::kernel_set_to_with_mask<unsigned short, 2><<<numBlocks,threadsPerBlock>>>((unsigned short *)mat.ptr, (float *)mask.ptr);
+        if (depth == 4) ::mat_operators::kernel_set_to_with_mask<unsigned int,   2><<<numBlocks,threadsPerBlock>>>((unsigned int *)mat.ptr, (float *)mask.ptr);
+    }
+    if (channels == 3)
+    {
+        if (depth == 1) ::mat_operators::kernel_set_to_with_mask<unsigned char,  3><<<numBlocks,threadsPerBlock>>>(mat.ptr, (float *)mask.ptr);
+        if (depth == 2) ::mat_operators::kernel_set_to_with_mask<unsigned short, 3><<<numBlocks,threadsPerBlock>>>((unsigned short *)mat.ptr, (float *)mask.ptr);
+        if (depth == 4) ::mat_operators::kernel_set_to_with_mask<unsigned int,   3><<<numBlocks,threadsPerBlock>>>((unsigned int *)mat.ptr, (float *)mask.ptr);
+    }
 }
 
 extern "C" void cv::gpu::impl::set_to_without_mask(const DevMem2D& mat, const double * scalar, int depth, int channels)
diff --git a/modules/gpu/src/gpumat.cpp b/modules/gpu/src/gpumat.cpp
index 0208d28f7..28498680d 100644
--- a/modules/gpu/src/gpumat.cpp
+++ b/modules/gpu/src/gpumat.cpp
@@ -81,14 +81,13 @@ void GpuMat::convertTo( GpuMat& /*m*/, int /*rtype*/, double /*alpha*/, double /
 
 GpuMat& GpuMat::operator = (const Scalar& s)
 {
-    CV_Assert(!"Not implemented");
     cv::gpu::impl::set_to_without_mask(*this, s.val, this->depth(), this->channels());
     return *this;
 }
 
 GpuMat& GpuMat::setTo(const Scalar& s, const GpuMat& mask)
 {
-    CV_Assert(!"Not implemented");
+    CV_Assert(mask.type() == CV_32F);
 
     CV_DbgAssert(!this->empty());