         }
     }
 
+    goto out;
+err:
+out:
     GHOST_CALL_RETURN(ghost_cu_free(cucolfield));
     GHOST_CALL_RETURN(ghost_cu_free(curowfield));
     GHOST_CALL_RETURN(ghost_cu_free(d_a));
-
-    return GHOST_SUCCESS;
+    GHOST_INSTR_STOP(vscale)
+
+    return ret;
 }
 
 extern "C" ghost_error_t ghost_densemat_cm_cu_fromScalar(ghost_densemat_t *vec, void *a)
 {
+    ghost_error_t ret = GHOST_SUCCESS;
+    
     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));
+    char *cucolfield = NULL, *curowfield = NULL;
+
+    if (!ghost_bitmap_iscompact(vec->ldmask) || 
+            !ghost_bitmap_iscompact(vec->trmask)) {
+        
+        WARNING_LOG("Potentially slow fromScalar operation because some rows or columns are masked out!");
+        GHOST_CALL_GOTO(ghost_cu_malloc((void **)&cucolfield,vec->traits.ncolsorig),err,ret);
+        GHOST_CALL_GOTO(ghost_cu_malloc((void **)&curowfield,vec->traits.nrowsorig),err,ret);
+
+        ghost_densemat_mask2charfield(vec->trmask,vec->traits.ncolsorig,colfield);
+        ghost_densemat_mask2charfield(vec->ldmask,vec->traits.nrowsorig,rowfield);
+
+        GHOST_CALL_GOTO(ghost_cu_upload(cucolfield,colfield,vec->traits.ncolsorig),err,ret);
+        GHOST_CALL_GOTO(ghost_cu_upload(curowfield,rowfield,vec->traits.nrowsorig),err,ret);
+    }
+    
     ghost_densemat_cm_malloc(vec);
     if (vec->traits.datatype & GHOST_DT_COMPLEX)
     {

        {
            if (vec->traits.datatype & GHOST_DT_DOUBLE)
            {
                CUBLAS_CALL_RETURN(cublasZdotc(ghost_cublas_handle,vec->traits.nrows,
                            (const cuDoubleComplex *)v1,1,(const cuDoubleComplex *)v2,1,&((cuDoubleComplex *)res)[v]));
            } 
            else 
            {
                CUBLAS_CALL_RETURN(cublasCdotc(ghost_cublas_handle,vec->traits.nrows,
                            (const cuFloatComplex *)v1,1,(const cuFloatComplex *)v2,1,&((cuFloatComplex *)res)[v]));
            }
        }
        else
        {
            if (vec->traits.datatype & GHOST_DT_DOUBLE)
            {
                CUBLAS_CALL_RETURN(cublasDdot(ghost_cublas_handle,vec->traits.nrows,
                            (const double *)v1,1,(const double *)v2,1,&((double *)res)[v]));
            } 
            else 
            {
                CUBLAS_CALL_RETURN(cublasSdot(ghost_cublas_handle,vec->traits.nrows,
                            (const float *)v1,1,(const float *)v2,1,&((float *)res)[v]));
            }
        }
    }

    if (vec->traits.flags & GHOST_DENSEMAT_SCATTERED) {
        vecclone->destroy(vecclone);
    }
    
    if (vec2->traits.flags & GHOST_DENSEMAT_SCATTERED) {
        vec2clone->destroy(vec2clone);
    }
    return GHOST_SUCCESS;
}

extern "C" ghost_error_t ghost_densemat_cm_cu_axpy(ghost_densemat_t *vec, ghost_densemat_t *vec2, void *a)

                        (float *)vec->cu_val,1));
        }
    }*/
    goto out;
err:
out:
    
    GHOST_CALL_RETURN(ghost_cu_free(cucolfield));
    GHOST_CALL_RETURN(ghost_cu_free(curowfield));
    GHOST_INSTR_STOP(axpy)

    return ret;
}

extern "C" ghost_error_t ghost_densemat_cm_cu_axpby(ghost_densemat_t *v1, ghost_densemat_t *v2, void *a, void *b)
{
    if (v1->traits.datatype != v2->traits.datatype)
    {
        ERROR_LOG("Cannot AXPBY vectors with different data types");
        return GHOST_ERR_NOT_IMPLEMENTED;
    }
    GHOST_INSTR_START(axpby);
    ghost_error_t ret = GHOST_SUCCESS;

    char colfield[v1->traits.ncolsorig];
    char rowfield[v1->traits.nrowsorig];

    char *cucolfield = NULL, *curowfield = NULL;
    
    if (!ghost_bitmap_iscompact(v1->ldmask) || 
            !ghost_bitmap_iscompact(v1->trmask) || 
            !ghost_bitmap_iscompact(v2->ldmask) || 
            !ghost_bitmap_iscompact(v2->trmask)) {
        
        WARNING_LOG("Potentially slow AXPBY operation because some rows or columns are masked out!");
        
        GHOST_CALL_GOTO(ghost_cu_malloc((void **)&cucolfield,v1->traits.ncolsorig),err,ret);
        GHOST_CALL_GOTO(ghost_cu_malloc((void **)&curowfield,v1->traits.nrowsorig),err,ret);

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

        GHOST_CALL_GOTO(ghost_cu_upload(cucolfield,colfield,v1->traits.ncolsorig),err,ret);
        GHOST_CALL_GOTO(ghost_cu_upload(curowfield,rowfield,v1->traits.nrowsorig),err,ret);
    }

    if (v1->traits.datatype & GHOST_DT_COMPLEX)
    {
        if (v1->traits.datatype & GHOST_DT_DOUBLE)
