a wierd result of a project with GpuMat and cuda

edit

The code like following

#include <opencv2/core/core.hpp>
#include <opencv2/gpu/gpu.hpp>
#include <opencv2/core/cuda_devptrs.hpp>
#include <opencv2/gpu/stream_accessor.hpp>
#include <opencv2/gpu/device/common.hpp>  
#include <opencv2/highgui/highgui.hpp>
#include <iostream>
#include <cuda_runtime.h>
using namespace std;
using namespace cv;
using namespace cv::gpu;

__global__ void AddColor_Kernel(const PtrStepSz<uchar3> src,const PtrStepSz<uchar> mask, 
PtrStep<uchar3> dst, PtrStep<uchar> dst_mask)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.x;

    if(y < mask.rows && x < mask.cols)
    {
        if(mask(y,x))
            dst.ptr(y )[ x] = src.ptr(y)[x];
        dst_mask.ptr(y )[ x] |= mask.ptr(y)[x];
    }
}




void AddColorCaller(const PtrStepSz<uchar3> &src,const PtrStepSz<uchar> &mask, PtrStep<uchar3> dst,
PtrStep<uchar> dst_mask,cudaStream_t stream)
    {

    dim3 block(32,8);
    //  dim3 grid(divUp(src.cols ,block.x),divUp(src.rows ,block.y));
    dim3 grid((mask.cols + block.x - 1)/block.x,(mask.rows + block.y - 1)/block.y);
    AddColor_Kernel<<<grid,block,0,stream>>>(src,mask,dst,dst_mask);
    cudaSafeCall(cudaGetLastError());
    if(stream == 0)
        cudaSafeCall(cudaDeviceSynchronize());
}

void AddColor(const GpuMat &src,const GpuMat &mask,  GpuMat &dst,  GpuMat &dst_mask,
Stream& stream = Stream::Null())
    {   
        dst.create(src.size(),src.type());
        dst.setTo(Scalar::all(0));
        dst_mask.create(mask.size(),mask.type());
        dst_mask.setTo(Scalar::all(0));
        cudaStream_t st = StreamAccessor::getStream(stream);
        AddColorCaller(src,mask,dst,dst_mask,st);
    }

int main()
{
    Mat image = imread("1.jpg");
    Mat mask  = imread("mask.jpg",-1);
    cout<<mask.channels()<<endl;
    resize(mask,mask,image.size());

    imshow("src",image);
    imshow("mask",mask);
    GpuMat gpuMat,output,mask_gpu,dst,dst_mask;

    gpuMat.upload(image);
    mask_gpu.upload(mask);
    AddColor(gpuMat,mask_gpu,dst,dst_mask);
    namedWindow("1",CV_WINDOW_OPENGL);
    imshow("1",dst);
    waitKey(0); 
    return 0;
}

According to code,the final_mask should be the same as origin_mask.The pic should be cut according to the mask.But result is disappointed,So I hope somebody can point out error of my code.Thanks!