 {
     int idx = blockIdx.x*blockDim.x+threadIdx.x;
 
-    for (;idx < nrows; idx+=gridDim.x*blockDim.x) {
-        if (rowmask[idx]) {
+    if (rowmask || colmask) {
+        for (;idx < nrows; idx+=gridDim.x*blockDim.x) {
+            if (rowmask[idx]) {
+                ghost_idx_t v;
+                for (v=0; v<ncols; v++) {
+                    if (colmask[v]) {
+                        v1[v*nrowspadded+idx] = axpby<T>(v2[v*nrowspadded+idx],v1[v*nrowspadded+idx],a,b);
+                    }
+                }
+            }
+        }
+    } else {
+        for (;idx < nrows; idx+=gridDim.x*blockDim.x) {
             ghost_idx_t v;
             for (v=0; v<ncols; v++) {
-                if (colmask[v]) {
-                    v1[v*nrowspadded+idx] = axpby<T>(v2[v*nrowspadded+idx],v1[v*nrowspadded+idx],a,b);
-                }
+                v1[v*nrowspadded+idx] = axpby<T>(v2[v*nrowspadded+idx],v1[v*nrowspadded+idx],a,b);
             }
         }
     }

                    vec->traits.nrowsorig,curowfield,vec->traits.ncolsorig,cucolfield,vec->traits.nrowspadded);
        }
    }

    GHOST_CALL_RETURN(ghost_cu_free(cucolfield));
    GHOST_CALL_RETURN(ghost_cu_free(curowfield));

    return GHOST_SUCCESS;
}

extern "C" ghost_error_t ghost_densemat_cm_cu_vscale(ghost_densemat_t *vec, void *a)
{
    void *d_a;
    size_t sizeofdt;
    ghost_datatype_size(&sizeofdt,vec->traits.datatype);
    GHOST_CALL_RETURN(ghost_cu_malloc(&d_a,vec->traits.ncols*sizeofdt));
    ghost_cu_upload(d_a,a,vec->traits.ncols*sizeofdt);
    char colfield[vec->traits.ncolsorig];
    char rowfield[vec->traits.nrowsorig];

    char *cucolfield, *curowfield;
    GHOST_CALL_RETURN(ghost_cu_malloc((void **)&cucolfield,vec->traits.ncolsorig));
    GHOST_CALL_RETURN(ghost_cu_malloc((void **)&curowfield,vec->traits.nrowsorig));

    ghost_densemat_mask2charfield(vec->trmask,vec->traits.ncolsorig,colfield);
    ghost_densemat_mask2charfield(vec->ldmask,vec->traits.nrowsorig,rowfield);

    GHOST_CALL_RETURN(ghost_cu_upload(cucolfield,colfield,vec->traits.ncolsorig));
    GHOST_CALL_RETURN(ghost_cu_upload(curowfield,rowfield,vec->traits.nrowsorig));

    if (vec->traits.datatype & GHOST_DT_COMPLEX)
    {

        {
            if (vec->traits.datatype & GHOST_DT_DOUBLE)
            {
                CURAND_CALL_GOTO(curandGenerateUniformDouble(gen,
                            &((double *)(vec->cu_val))[v*vec->traits.nrowspadded],
                            vec->traits.nrows*2),err,ret);
            } 
            else 
            {
                CURAND_CALL_GOTO(curandGenerateUniform(gen,
                            &((float *)(vec->cu_val))[v*vec->traits.nrowspadded],
                            vec->traits.nrows*2),err,ret);
            }
        }
        else
        {
            if (vec->traits.datatype & GHOST_DT_DOUBLE)
            {
                CURAND_CALL_GOTO(curandGenerateUniformDouble(gen,
                            &((double *)(vec->cu_val))[v*vec->traits.nrowspadded],
                            vec->traits.nrows),err,ret);
            } 
            else 
            {
                CURAND_CALL_GOTO(curandGenerateUniform(gen,
                            &((float *)(vec->cu_val))[v*vec->traits.nrowspadded],
                            vec->traits.nrows),err,ret);
            }
        }
    }
