     (C(data,RIGHT_SUBSPACE)*) ctx->right_subspace_data,
     d_maxs);
 
-  err = cudaThreadSynchronize();CHKERRCUDA(err);
+  err = cudaDeviceSynchronize();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));
}

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


  PetscCall(VecGetSize(b, &size));

  err = cudaDeviceSynchronize();CHKERRCUDA(err);

  C(device_MatMult,C(LEFT_SUBSPACE,RIGHT_SUBSPACE))<<<GPU_BLOCK_NUM,GPU_BLOCK_SIZE>>>(
    size,
