     xarray,
     barray);
 
-  err = cudaThreadSynchronize();CHKERRCUDA(err);
+  err = cudaDeviceSynchronize();CHKERRCUDA(err);
 
   PetscCall(VecCUDARestoreArrayRead(x, &xarray));
   PetscCall(VecCUDARestoreArray(b, &barray));

    (C(data,RIGHT_SUBSPACE)*) ctx->right_subspace_data,
    d_maxs);

  err = cudaThreadSynchronize();CHKERRCUDA(err);

  err = cudaMemcpy(h_maxs, d_maxs, sizeof(PetscReal)*GPU_BLOCK_NUM, cudaMemcpyDeviceToHost);CHKERRCUDA(err);


  PetscCall(C(CopySubspaceData_CUDA,RIGHT_SUBSPACE)(
    (C(data,RIGHT_SUBSPACE)**)&(ctx->right_subspace_data),
    (C(data,RIGHT_SUBSPACE)*)right_subspace_data));

  return 0;
}

PetscErrorCode C(MatDestroyCtx_GPU,C(LEFT_SUBSPACE,RIGHT_SUBSPACE))(Mat A)
