     onevec->fromScalar(onevec,one);
 
     one[1] = 0.;
-
-    ghost_lidx_t r;
-    for (r=0; r<vec->traits.nrows; r++)
-    {
-        if (vec->traits.datatype & GHOST_DT_COMPLEX)
-        {
-            if (vec->traits.datatype & GHOST_DT_DOUBLE)
-            {
-                CURAND_CALL_GOTO(curandGenerateUniformDouble(gen,
-                            &((double *)(vec->cu_val))[r*vec->traits.ncolspadded],
-                            vec->traits.ncols*2),err,ret);
-            } 
-            else 
-            {
-                CURAND_CALL_GOTO(curandGenerateUniform(gen,
-                            &((float *)(vec->cu_val))[r*vec->traits.ncolspadded],
-                            vec->traits.ncols*2),err,ret);
-            }
-        }
-        else
-        {
-            if (vec->traits.datatype & GHOST_DT_DOUBLE)
-            {
-                CURAND_CALL_GOTO(curandGenerateUniformDouble(gen,
-                            &((double *)(vec->cu_val))[r*vec->traits.ncolspadded],
-                            vec->traits.ncols),err,ret);
-            } 
-            else 
-            {
-                CURAND_CALL_GOTO(curandGenerateUniform(gen,
-                            &((float *)(vec->cu_val))[r*vec->traits.ncolspadded],
-                            vec->traits.ncols),err,ret);
-            }
+    
+    if (vec->traits.datatype & GHOST_DT_COMPLEX)
+    {
+        if (vec->traits.datatype & GHOST_DT_DOUBLE)
+        {
+            CURAND_CALL_GOTO(curandGenerateUniformDouble(gen,
+                        (double *)(vec->cu_val),
+                        vec->traits.ncolspadded*vec->traits.nrows*2),err,ret);
+        } 
+        else 
+        {
+            CURAND_CALL_GOTO(curandGenerateUniform(gen,
+                        (float *)(vec->cu_val),
+                        vec->traits.ncolspadded*vec->traits.nrows*2),err,ret);
+        }
+    }
+    else
+    {
+        if (vec->traits.datatype & GHOST_DT_DOUBLE)
+        {
+            CURAND_CALL_GOTO(curandGenerateUniformDouble(gen,
+                        (double *)(vec->cu_val),
+                        vec->traits.ncolspadded*vec->traits.nrows),err,ret);
+        } 
+        else 
+        {
+            CURAND_CALL_GOTO(curandGenerateUniform(gen,
+                        (float *)(vec->cu_val),
+                        vec->traits.ncolspadded*vec->traits.nrows),err,ret);
         }
     }
     if (vec->traits.datatype & GHOST_DT_DOUBLE) {

    char *cucolfield = NULL, *curowfield = NULL;
    int grid = (int)ceil((double)v1->traits.nrows/THREADSPERBLOCK);
    dim3 block (THREADSPERBLOCK/v1->traits.ncolsorig,v1->traits.ncolsorig); 
    
    if (ghost_bitmap_weight(v1->ldmask) != v1->traits.ncolsorig || 
            ghost_bitmap_weight(v1->trmask) != v1->traits.nrowsorig ||
            ghost_bitmap_weight(v2->ldmask) != v2->traits.ncolsorig ||



    int nrank, proc;
    ghost_context_t *ctx = vec->context;
    
    ghost_nrank(&nrank,ctx->mpicomm); 
            
    for (proc=0 ; proc<nrank ; proc++){
        dim3 block((int)ceil((double)THREADSPERBLOCK/vec->traits.ncols),vec->traits.ncols);
        dim3 grid((int)ceil((double)ctx->dues[proc]/block.x));
        INFO_LOG("communication assembly with grid %d block %dx%d",grid.x,block.x,block.y);
        if (vec->traits.datatype & GHOST_DT_COMPLEX)
        {
            if (vec->traits.datatype & GHOST_DT_DOUBLE)
            {
                if (ctx->dues[proc]) {
                    cu_communicationassembly_kernel<cuDoubleComplex><<< grid,block >>>((cuDoubleComplex *)vec->cu_val, ((cuDoubleComplex *)work),dueptr[proc],ctx->cu_duelist[proc],vec->traits.ncols,ctx->dues[proc],vec->traits.ncolspadded,perm);
                }
            } 
            else 
            {
                if (ctx->dues[proc]) {
                    cu_communicationassembly_kernel<cuFloatComplex><<< grid,block >>>((cuFloatComplex *)vec->cu_val, ((cuFloatComplex *)work),dueptr[proc],ctx->cu_duelist[proc],vec->traits.ncols,ctx->dues[proc],vec->traits.ncolspadded,perm);
                }
            }
        }
