         {
             CURAND_CALL_GOTO(curandGenerateUniformDouble(gen,
                         (double *)(valptr),
-                        vec->traits.ncolspadded*vec->traits.nrows*2),err,ret);
+                        compactvec->traits.ncolsorig*compactvec->traits.nrows*2),err,ret);
         } 
         else 
         {
             CURAND_CALL_GOTO(curandGenerateUniform(gen,
                         (float *)(valptr),
-                        vec->traits.ncolspadded*vec->traits.nrows*2),err,ret);
+                        compactvec->traits.ncolsorig*compactvec->traits.nrows*2),err,ret);
         }
     }
     else
     {
         if (vec->traits.datatype & GHOST_DT_DOUBLE)
         {
-            int r;
-                for (r=0; r<vec->traits.nrowsorig; r++) {
-                    if (ghost_bitmap_isset(vec->trmask,r)) {
             CURAND_CALL_GOTO(curandGenerateUniformDouble(gen,
-                        (double *)(valptr)+r*vec->traits.ncolspadded,
-                        /*vec->traits.ncolspadded**/vec->traits.ncols),err,ret);
-                    }
-                }
+                        (double *)(valptr),
+                        compactvec->traits.ncolsorig*compactvec->traits.nrows),err,ret);
         } 
         else 
         {
             CURAND_CALL_GOTO(curandGenerateUniform(gen,
                         (float *)(valptr),
-                        vec->traits.ncolspadded*vec->traits.nrows),err,ret);
-        }
-    }
-    if (vec->traits.datatype & GHOST_DT_DOUBLE) {
-        vec->axpby(vec,onevecview,minusahalf,one);
+                        compactvec->traits.ncolsorig*compactvec->traits.nrows),err,ret);
+        }
+    }
+    if (compactvec->traits.datatype & GHOST_DT_DOUBLE) {
+        compactvec->axpby(compactvec,onevecview,minusahalf,one);
     } else {
-        vec->axpby(vec,onevecview,fminusahalf,fone);
+        compactvec->axpby(compactvec,onevecview,fminusahalf,fone);
+    }
+    if (compactvec != vec) {
+        vec->fromVec(vec,compactvec,0,0);
+        compactvec->destroy(compactvec);
     }
     goto out;
 err:


    one[1] = 0.;
    void *valptr;
    ghost_densemat_cu_valptr(vec,&valptr);
    
    if (vec->traits.datatype & GHOST_DT_COMPLEX)
    {



template<typename T>  
__global__ static void cu_vaxpby_kernel(T *v1, T *v2, T *a, T *b, ghost_lidx_t nrows, ghost_lidx_t ncols, ghost_lidx_t ld1, ghost_lidx_t ld2)
{
    int idx = blockIdx.x*blockDim.x+threadIdx.x;

    for (;idx < nrows; idx+=gridDim.x*blockDim.x) {
        ghost_lidx_t v;
        for (v=0; v<ncols; v++) {
            v1[idx*ld1+v] = axpby<T>(v2[idx*ld2+v],v1[idx*ld1+v],a[v],b[v]);
        }
    }
}

template<typename T>  
__global__ static void cu_axpby_kernel(T *v1, T *v2, T a, T b, ghost_lidx_t nrows, ghost_lidx_t ncols, ghost_lidx_t ld1, ghost_lidx_t ld2) 
{
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    for (;idx < nrows; idx+=gridDim.x*blockDim.x) {
        ghost_lidx_t v;
        for (v=0; v<ncols; v++) {
            v1[idx*ld1+v] = axpby<T>(v2[idx*ld2+v],v1[idx*ld1+v],a,b);
        }
    }
}

template<typename T>  
__global__ static void cu_scale_kernel(T *vec, T a, ghost_lidx_t nrows, ghost_lidx_t ncols, ghost_lidx_t ld)
{
    int idx = blockIdx.x*blockDim.x+threadIdx.x;

    for (;idx < nrows; idx+=gridDim.x*blockDim.x) {
        ghost_lidx_t v;
        for (v=0; v<ncols; v++) {
            vec[idx*ld+v] = scale<T>(a,vec[idx*ld+v]);
        }
    }

}

template<typename T>  
__global__ static void cu_vscale_kernel(T *vec, T *a, ghost_lidx_t nrows, ghost_lidx_t ncols, ghost_lidx_t ld)
{
    int idx = blockIdx.x*blockDim.x+threadIdx.x;

    for (;idx < nrows; idx+=gridDim.x*blockDim.x) {
        ghost_lidx_t v;
        for (v=0; v<ncols; v++) {
            vec[idx*ld+v] = scale<T>(a[v],vec[idx*ld+v]);
        }
    }
}

template<typename T>  
__global__ static void cu_fromscalar_kernel(T *vec, T a, ghost_lidx_t nrows, ghost_lidx_t ncols, ghost_lidx_t ld)
{
    int idx = blockIdx.x*blockDim.x+threadIdx.x;

    for (;idx < nrows; idx+=gridDim.x*blockDim.x) {
        ghost_lidx_t v;
        for (v=0; v<ncols; v++) {
            vec[idx*ld+v] = a;
        }
    }
}


template<typename T>  
__global__ static void cu_communicationassembly_kernel(T *vec, T *work, ghost_lidx_t offs, ghost_lidx_t *duelist, ghost_lidx_t ncols, ghost_lidx_t ndues, ghost_lidx_t ncolspadded, ghost_lidx_t *perm)
