From 8318f2c4cf992a2d5c21485070548b26514dc0b9 Mon Sep 17 00:00:00 2001
From: schnellerhase <56360279+schnellerhase@users.noreply.github.com>
Date: Sun, 20 Oct 2024 22:39:01 +0200
Subject: [PATCH] [WIP] Transfer matrix operation

---
 cpp/dolfinx/CMakeLists.txt                 |   1 +
 cpp/dolfinx/multigrid/CMakeLists.txt       |   4 +
 cpp/dolfinx/multigrid/transfer_matrix.h    | 259 ++++++++++++++
 cpp/test/CMakeLists.txt                    |   1 +
 cpp/test/multigrid/transfer_matrix.cpp     | 386 +++++++++++++++++++++
 python/CMakeLists.txt                      |   1 +
 python/dolfinx/transfer.py                 |  29 ++
 python/dolfinx/wrappers/dolfinx.cpp        |   6 +
 python/dolfinx/wrappers/multigrid.cpp      |  67 ++++
 python/test/unit/transfer/test_transfer.py |  63 ++++
 10 files changed, 817 insertions(+)
 create mode 100644 cpp/dolfinx/multigrid/CMakeLists.txt
 create mode 100644 cpp/dolfinx/multigrid/transfer_matrix.h
 create mode 100644 cpp/test/multigrid/transfer_matrix.cpp
 create mode 100644 python/dolfinx/transfer.py
 create mode 100644 python/dolfinx/wrappers/multigrid.cpp
 create mode 100644 python/test/unit/transfer/test_transfer.py

diff --git a/cpp/dolfinx/CMakeLists.txt b/cpp/dolfinx/CMakeLists.txt
index 7adaaeb73cf..001cb2116be 100644
--- a/cpp/dolfinx/CMakeLists.txt
+++ b/cpp/dolfinx/CMakeLists.txt
@@ -17,6 +17,7 @@ set(DOLFINX_DIRS
     io
     la
     mesh
+    multigrid
     nls
     refinement
 )
diff --git a/cpp/dolfinx/multigrid/CMakeLists.txt b/cpp/dolfinx/multigrid/CMakeLists.txt
new file mode 100644
index 00000000000..36d60438135
--- /dev/null
+++ b/cpp/dolfinx/multigrid/CMakeLists.txt
@@ -0,0 +1,4 @@
+set(HEADERS_multigrid
+    ${CMAKE_CURRENT_SOURCE_DIR}/transfer_matrix.h
+    PARENT_SCOPE
+)
diff --git a/cpp/dolfinx/multigrid/transfer_matrix.h b/cpp/dolfinx/multigrid/transfer_matrix.h
new file mode 100644
index 00000000000..bb68fc49fa3
--- /dev/null
+++ b/cpp/dolfinx/multigrid/transfer_matrix.h
@@ -0,0 +1,259 @@
+// Copyright (C) 2024 Paul T. Kühner
+//
+// This file is part of DOLFINX (https://www.fenicsproject.org)
+//
+// SPDX-License-Identifier:    LGPL-3.0-or-later
+
+#pragma once
+
+#include <algorithm>
+#include <concepts>
+#include <cstdint>
+#include <iterator>
+#include <stdexcept>
+#include <vector>
+
+#include <basix/element-families.h>
+
+#include "dolfinx/common/IndexMap.h"
+#include "dolfinx/fem/FunctionSpace.h"
+#include "dolfinx/la/SparsityPattern.h"
+#include "dolfinx/la/utils.h"
+#include "dolfinx/mesh/Mesh.h"
+
+namespace dolfinx::multigrid
+{
+
+template <std::floating_point T>
+std::vector<std::int64_t>
+inclusion_mapping(const dolfinx::mesh::Mesh<T>& mesh_from,
+                  const dolfinx::mesh::Mesh<T>& mesh_to)
+{
+
+  const common::IndexMap& im_from = *mesh_from.topology()->index_map(0);
+  const common::IndexMap& im_to = *mesh_to.topology()->index_map(0);
+
+  std::vector<std::int64_t> map(im_from.size_global(), -1);
+
+  std::span<const T> x_from = mesh_from.geometry().x();
+  std::span<const T> x_to = mesh_to.geometry().x();
+
+  auto build_global_to_local = [&](const auto& im)
+  {
+    return [&](std::int32_t idx)
+    {
+      std::array<std::int64_t, 1> tmp;
+      im.local_to_global(std::vector<std::int32_t>{idx}, tmp);
+      return tmp[0];
+    };
+  };
+
+  auto to_global_to = build_global_to_local(im_to);
+  auto to_global_from = build_global_to_local(im_from);
+
+  for (std::int32_t i = 0; i < im_from.size_local(); i++)
+  {
+    std::ranges::subrange vertex_from(std::next(x_from.begin(), 3 * i),
+                                      std::next(x_from.begin(), 3 * (i + 1)));
+    for (std::int64_t j = 0; j < im_to.size_local() + im_to.num_ghosts(); j++)
+    {
+      std::ranges::subrange vertex_to(std::next(x_to.begin(), 3 * j),
+                                      std::next(x_to.begin(), 3 * (j + 1)));
+
+      if (std::ranges::equal(
+              vertex_from, vertex_to, [](T a, T b)
+              { return std::abs(a - b) <= std::numeric_limits<T>::epsilon(); }))
+      {
+        map[to_global_from(i)] = to_global_to(j);
+        break;
+      }
+    }
+  }
+
+  if (dolfinx::MPI::size(mesh_to.comm()) == 1)
+  {
+    // no communication required
+    assert(std::ranges::all_of(map, [](auto e) { return e >= 0; }));
+    return map;
+  }
+
+  // map holds at this point for every original local index the corresponding
+  // mapped global index. All other entries are still -1, but are available on
+  // other processes.
+  std::vector<std::int64_t> result(map.size(), -1);
+  MPI_Allreduce(map.data(), result.data(), map.size(),
+                dolfinx::MPI::mpi_t<std::int64_t>, MPI_MAX, mesh_from.comm());
+
+  assert(std::ranges::all_of(result, [](auto e) { return e >= 0; }));
+  return result;
+}
+
+template <std::floating_point T>
+la::SparsityPattern
+create_sparsity_pattern(const dolfinx::fem::FunctionSpace<T>& V_from,
+                        const dolfinx::fem::FunctionSpace<T>& V_to,
+                        const std::vector<std::int64_t>& inclusion_map)
+{
+  if (*V_from.element() != *V_to.element())
+    throw std::runtime_error(
+        "Transfer between different element types not supported");
+
+  if (V_from.element()->basix_element().family() != basix::element::family::P)
+    throw std::runtime_error("Only Lagrange elements supported");
+
+  if (V_from.element()->basix_element().degree() != 1)
+    throw std::runtime_error("Only piecewise linear elements supported");
+
+  // TODO: mixed elements? value shapes? DG?
+
+  auto mesh_from = V_from.mesh();
+  auto mesh_to = V_to.mesh();
+  assert(mesh_from);
+  assert(mesh_to);
+
+  MPI_Comm comm = mesh_from->comm();
+  {
+    // Check comms equal
+    int result;
+    MPI_Comm_compare(comm, mesh_to->comm(), &result);
+    assert(result == MPI_CONGRUENT);
+  }
+  assert(mesh_from->topology()->dim() == mesh_to->topology()->dim());
+
+  auto to_v_to_f = mesh_to->topology()->connectivity(0, 1);
+  auto to_f_to_v = mesh_to->topology()->connectivity(1, 0);
+  assert(to_v_to_f);
+  assert(to_f_to_v);
+
+  auto dofmap_from = V_from.dofmap();
+  auto dofmap_to = V_to.dofmap();
+  assert(dofmap_from);
+  assert(dofmap_to);
+
+  assert(mesh_to->topology()->index_map(0));
+  assert(mesh_from->topology()->index_map(0));
+  const common::IndexMap& im_to = *mesh_to->topology()->index_map(0);
+  const common::IndexMap& im_from = *mesh_from->topology()->index_map(0);
+
+  dolfinx::la::SparsityPattern sp(
+      comm, {dofmap_from->index_map, dofmap_to->index_map},
+      {dofmap_from->index_map_bs(), dofmap_to->index_map_bs()});
+
+  assert(inclusion_map.size() == im_from.size_global());
+  for (int dof_from_global = 0; dof_from_global < im_from.size_global();
+       dof_from_global++)
+  {
+    std::int64_t dof_to_global = inclusion_map[dof_from_global];
+
+    std::vector<std::int32_t> local_dof_to_v{0};
+    im_to.global_to_local(std::vector<std::int64_t>{dof_to_global},
+                          local_dof_to_v);
+
+    auto local_dof_to = local_dof_to_v[0];
+
+    bool is_remote = (local_dof_to == -1);
+    bool is_ghost = local_dof_to >= im_to.size_local();
+    if (is_remote || is_ghost)
+      continue;
+
+    std::vector<std::int32_t> dof_from_v{0};
+    im_from.global_to_local(std::vector<std::int64_t>{dof_from_global},
+                            dof_from_v);
+
+    std::ranges::for_each(
+        to_v_to_f->links(local_dof_to),
+        [&](auto e)
+        {
+          sp.insert(
+              std::vector<int32_t>(to_f_to_v->links(e).size(), dof_from_v[0]),
+              to_f_to_v->links(e));
+        });
+  }
+  sp.finalize();
+  return sp;
+}
+
+template <std::floating_point T>
+void assemble_transfer_matrix(la::MatSet<T> auto mat_set,
+                              const dolfinx::fem::FunctionSpace<T>& V_from,
+                              const dolfinx::fem::FunctionSpace<T>& V_to,
+                              const std::vector<std::int64_t>& inclusion_map,
+                              std::function<T(std::int32_t)> weight)
+{
+  if (*V_from.element() != *V_to.element())
+    throw std::runtime_error(
+        "Transfer between different element types not supported");
+
+  if (V_from.element()->basix_element().family() != basix::element::family::P)
+    throw std::runtime_error("Only Lagrange elements supported");
+
+  if (V_from.element()->basix_element().degree() != 1)
+    throw std::runtime_error("Only piecewise linear elements supported");
+
+  // TODO: mixed elements? value shapes? DG?
+
+  auto mesh_from = V_from.mesh();
+  auto mesh_to = V_to.mesh();
+  assert(mesh_from);
+
+  MPI_Comm comm = mesh_from->comm();
+  {
+    // Check comms equal
+    int result;
+    MPI_Comm_compare(comm, mesh_to->comm(), &result);
+    assert(result == MPI_CONGRUENT);
+  }
+  assert(mesh_from->topology()->dim() == mesh_to->topology()->dim());
+
+  auto to_v_to_f = mesh_to->topology()->connectivity(0, 1);
+  auto to_f_to_v = mesh_to->topology()->connectivity(1, 0);
+  assert(to_v_to_f);
+  assert(to_f_to_v);
+
+  auto dofmap_from = V_from.dofmap();
+  auto dofmap_to = V_to.dofmap();
+  assert(dofmap_from);
+  assert(dofmap_to);
+
+  assert(mesh_to->topology()->index_map(0));
+  assert(mesh_from->topology()->index_map(0));
+  const common::IndexMap& im_to = *mesh_to->topology()->index_map(0);
+  const common::IndexMap& im_from = *mesh_from->topology()->index_map(0);
+
+  assert(inclusion_map.size() == im_from.size_global());
+
+  for (int dof_from_global = 0; dof_from_global < im_from.size_global();
+       dof_from_global++)
+  {
+    std::int64_t dof_to_global = inclusion_map[dof_from_global];
+
+    std::vector<std::int32_t> local_dof_to_v{0};
+    im_to.global_to_local(std::vector<std::int64_t>{dof_to_global},
+                          local_dof_to_v);
+
+    auto local_dof_to = local_dof_to_v[0];
+
+    bool is_remote = (local_dof_to == -1);
+    bool is_ghost = local_dof_to >= im_to.size_local();
+    if (is_remote || is_ghost)
+      continue;
+
+    std::vector<std::int32_t> dof_from_v{0};
+    im_from.global_to_local(std::vector<std::int64_t>{dof_from_global},
+                            dof_from_v);
+
+    for (auto e : to_v_to_f->links(local_dof_to))
+    {
+      for (auto n : to_f_to_v->links(e))
+      {
+        // For now we only support distance <= 1 -> this should be somehow
+        // asserted.
+        std::int32_t distance = (n == local_dof_to) ? 0 : 1;
+        mat_set(std::vector<int32_t>{dof_from_v[0]}, std::vector<int32_t>{n},
+                std::vector<T>{weight(distance)});
+      }
+    }
+  }
+}
+
+} // namespace dolfinx::transfer
\ No newline at end of file
diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt
index ada545792a1..e956f3fda46 100644
--- a/cpp/test/CMakeLists.txt
+++ b/cpp/test/CMakeLists.txt
@@ -48,6 +48,7 @@ add_executable(
   mesh/refinement/interval.cpp
   mesh/refinement/option.cpp
   mesh/refinement/rectangle.cpp
+  multigrid/transfer_matrix.cpp
   ${CMAKE_CURRENT_BINARY_DIR}/poisson.c
 )
 target_link_libraries(unittests PRIVATE Catch2::Catch2WithMain dolfinx)
diff --git a/cpp/test/multigrid/transfer_matrix.cpp b/cpp/test/multigrid/transfer_matrix.cpp
new file mode 100644
index 00000000000..60fa6dd3b04
--- /dev/null
+++ b/cpp/test/multigrid/transfer_matrix.cpp
@@ -0,0 +1,386 @@
+// Copyright (C) 2024 Paul T. Kühner
+//
+// This file is part of DOLFINX (https://www.fenicsproject.org)
+//
+// SPDX-License-Identifier:    LGPL-3.0-or-later
+
+#include <basix/cell.h>
+#include <concepts>
+#include <cstdint>
+#include <memory>
+#include <optional>
+#include <vector>
+
+#include <mpi.h>
+
+#include <catch2/catch_template_test_macros.hpp>
+#include <catch2/catch_test_macros.hpp>
+#include <catch2/matchers/catch_matchers.hpp>
+#include <catch2/matchers/catch_matchers_floating_point.hpp>
+#include <catch2/matchers/catch_matchers_range_equals.hpp>
+
+#include <dolfinx/common/IndexMap.h>
+#include <dolfinx/common/MPI.h>
+#include <dolfinx/fem/FunctionSpace.h>
+#include <dolfinx/fem/utils.h>
+#include <dolfinx/la/MatrixCSR.h>
+#include <dolfinx/la/SparsityPattern.h>
+#include <dolfinx/mesh/cell_types.h>
+#include <dolfinx/mesh/generation.h>
+#include <dolfinx/mesh/utils.h>
+#include <dolfinx/refinement/plaza.h>
+#include <dolfinx/refinement/refine.h>
+#include <dolfinx/multigrid/transfer_matrix.h>
+
+using namespace dolfinx;
+using namespace Catch::Matchers;
+
+template <std::floating_point T>
+constexpr auto weight
+    = [](std::int32_t d) -> T { return d == 0 ? T(1) : T(.5); };
+
+template <typename T = double>
+std::shared_ptr<const fem::FiniteElement<T>>
+P1_element(basix::cell::type cell_type)
+{
+  auto element
+      = basix::create_element<T>(basix::element::family::P, cell_type, 1,
+                                 basix::element::lagrange_variant::unset,
+                                 basix::element::dpc_variant::unset, false);
+  return std::make_shared<const fem::FiniteElement<T>>(element);
+}
+
+TEMPLATE_TEST_CASE("Transfer Matrix 1D", "[transfer_matrix]", double, float)
+{
+  using T = TestType;
+
+  auto mesh_coarse
+      = std::make_shared<mesh::Mesh<T>>(dolfinx::mesh::create_interval<T>(
+          MPI_COMM_SELF, 2, {0.0, 1.0}, mesh::GhostMode::none));
+
+  auto [mesh_fine, parent_cell, parent_facet]
+      = refinement::refine(*mesh_coarse, std::nullopt);
+
+  auto element = P1_element<T>(basix::cell::type::interval);
+  auto V_coarse = std::make_shared<fem::FunctionSpace<T>>(
+      fem::create_functionspace<T>(mesh_coarse, element));
+  auto V_fine
+      = std::make_shared<fem::FunctionSpace<T>>(fem::create_functionspace<T>(
+          std::make_shared<mesh::Mesh<T>>(mesh_fine), element));
+
+  mesh_fine.topology()->create_connectivity(1, 0);
+  mesh_fine.topology()->create_connectivity(0, 1);
+
+  std::vector<std::int64_t> inclusion_map
+      = multigrid::inclusion_mapping(*mesh_coarse, mesh_fine);
+
+  CHECK_THAT(inclusion_map, RangeEquals(std::vector<std::int64_t>{0, 2, 4}));
+
+  la::SparsityPattern sp
+      = multigrid::create_sparsity_pattern(*V_coarse, *V_fine, inclusion_map);
+
+  la::MatrixCSR<T> transfer_matrix(std::move(sp), la::BlockMode::compact);
+  multigrid::assemble_transfer_matrix<T>(transfer_matrix.mat_set_values(),
+                                        *V_coarse, *V_fine, inclusion_map,
+                                        weight<T>);
+
+  // clang-format off
+  std::vector<T> expected{1.0, 0.5, 0.0, 0.0, 0.0,
+                          0.0, 0.5, 1.0, 0.5, 0.0,
+                          0.0, 0.0, 0.0, 0.5, 1.0};
+  // clang-format on
+  CHECK_THAT(transfer_matrix.to_dense(),
+             RangeEquals(expected,
+                         [](auto a, auto b) {
+                           return std::abs(a - b)
+                                  <= std::numeric_limits<T>::epsilon();
+                         }));
+}
+
+// TEMPLATE_TEST_CASE("Transfer Matrix 1D (parallel)", "[transfer_matrix]",
+//                    double) // TODO: float
+// {
+//   using T = TestType;
+
+//   int comm_size = dolfinx::MPI::size(MPI_COMM_WORLD);
+//   int rank = dolfinx::MPI::rank(MPI_COMM_WORLD);
+
+//   // TODO: see https://github.com/FEniCS/dolfinx/issues/3358
+//   //   auto part = mesh::create_cell_partitioner(mesh::GhostMode::shared_facet);
+//   mesh::CellPartitionFunction part
+//       = [&](MPI_Comm /* comm */, int /* nparts */,
+//             const std::vector<mesh::CellType>& /* cell_types */,
+//             const std::vector<std::span<const std::int64_t>>& /* cells */)
+//   {
+//     std::vector<std::vector<std::int32_t>> data;
+//     if (comm_size == 1)
+//       data = {{0}, {0}, {0}, {0}};
+//     else if (comm_size == 2)
+//       data = {{0}, {0}, {0}, {0}, {0, 1}, {1, 0}, {1}, {1}, {1}};
+//     else if (comm_size == 3)
+//       data = {{1}, {1},    {1},    {1}, {1, 2}, {2, 1}, {2},
+//               {2}, {2, 0}, {0, 2}, {0}, {0},    {0},    {0}};
+//     else
+//       FAIL("Test only supports <= 3 processes");
+//     return graph::AdjacencyList<std::int32_t>(std::move(data));
+//   };
+
+//   auto mesh_coarse = std::make_shared<mesh::Mesh<T>>(
+//       mesh::create_interval<T>(MPI_COMM_WORLD, 5 * comm_size - 1, {0., 1.},
+//                                mesh::GhostMode::none, part));
+
+//   // TODO: see https://github.com/FEniCS/dolfinx/issues/3358
+//   //   auto part_refine
+//   //       = mesh::create_cell_partitioner(mesh::GhostMode::shared_facet);
+//   mesh::CellPartitionFunction part_refine
+//       = [&](MPI_Comm comm, int /* nparts */,
+//             const std::vector<mesh::CellType>& /* cell_types */,
+//             const std::vector<std::span<const std::int64_t>>& /* cells */)
+//   {
+//     std::vector<std::vector<std::int32_t>> data;
+//     if (dolfinx::MPI::size(comm) == 1)
+//       data = {{0}, {0}, {0}, {0}, {0}, {0}, {0}, {0}};
+//     else if (dolfinx::MPI::size(comm) == 2)
+//     {
+//       if (rank == 0)
+//         data = {{0}, {0}, {0}, {0}, {0}, {0}, {0}, {0}, {0, 1}, {1, 0}};
+//       else
+//         data = {{1}, {1}, {1}, {1}, {1}, {1}, {1}, {1}};
+//     }
+//     else if (dolfinx::MPI::size(comm) == 3)
+//     {
+//       if (rank == 0)
+//         data = {{2, 0}, {0, 2}, {0}, {0}, {0}, {0}, {0}, {0}, {0}, {0}};
+//       else if (rank == 1)
+//         data = {{1}, {1}, {1}, {1}, {1}, {1}, {1}, {1}, {1}, {1, 2}};
+//       else
+//         data = {{2, 1}, {2}, {2}, {2}, {2}, {2}, {2}, {2}};
+//     }
+//     else
+//       FAIL("Test only supports <= 3 processes");
+//     return graph::AdjacencyList<std::int32_t>(std::move(data));
+//   };
+
+//   auto [mesh_fine, parent_cell, parent_facet] = refinement::refine(
+//       *mesh_coarse, std::nullopt, part_refine, refinement::Option::none);
+
+//   auto element = P1_element(basix::cell::type::interval);
+
+//   auto V_coarse = std::make_shared<fem::FunctionSpace<T>>(
+//       fem::create_functionspace<T>(mesh_coarse, element));
+//   auto V_fine
+//       = std::make_shared<fem::FunctionSpace<T>>(fem::create_functionspace<T>(
+//           std::make_shared<mesh::Mesh<T>>(mesh_fine), element));
+
+//   mesh_fine.topology()->create_connectivity(1, 0);
+//   mesh_fine.topology()->create_connectivity(0, 1);
+
+//   std::vector<int64_t> inclusion_map;
+//   switch (comm_size)
+//   {
+//   case 1:
+//     inclusion_map = {0, 2, 4, 6, 8};
+//     break;
+//   case 2:
+//     inclusion_map = {0, 2, 4, 6, 8, 10, 12, 14, 16, 18};
+//     break;
+//   case 3:
+//     inclusion_map = {0, 2, 4, 6, 8, 9, 11, 13, 15, 17, 19, 27, 25, 23, 20};
+//     break;
+//   }
+
+//   CHECK_THAT(multigrid::inclusion_mapping(*mesh_coarse, mesh_fine),
+//              RangeEquals(inclusion_map));
+
+//   la::SparsityPattern sp
+//       = multigrid::create_sparsity_pattern(*V_coarse, *V_fine, inclusion_map);
+
+//   la::MatrixCSR<T> transfer_matrix(std::move(sp), la::BlockMode::compact);
+//   multigrid::assemble_transfer_matrix<T>(transfer_matrix.mat_set_values(),
+//                                         *V_coarse, *V_fine, inclusion_map,
+//                                         weight<T>);
+
+//   std::array<std::vector<T>, 3> expected;
+//   if (comm_size == 1)
+//   {
+//     // clang-format off
+//     expected[0] = {
+//         /* row_0 */ 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_1 */ 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_2 */ 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0,
+//         /* row_3 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0,
+//         /* row_4 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0
+//     };
+//     // clang-format on
+//   }
+//   else if (comm_size == 2)
+//   {
+//     // clang-format off
+//     expected[0] = {
+//         /* row_0 */ 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_1 */ 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_2 */ 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_3 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_4 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_5 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_6 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0
+//     };
+//     expected[1] = {
+//         /* row_0 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_1 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0,
+//         /* row_2 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0,
+//         /* row_3 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0,
+//         /* row_4 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_5 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0
+//     };
+//     // clang-format on
+//   }
+//   else if (comm_size == 3)
+//   {
+//     // clang-format off
+//     expected[0] = {
+//         /* row_0 */ 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_1 */ 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_2 */ 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_3 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_4 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_5 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_6 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//     };
+//     expected[1] = {
+//         /* row_0 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_1 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_2 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_3 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_4 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_5 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5,
+//         /* row_6 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//     };
+//     expected[2] = {
+//         /* row_0 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5,
+//         /* row_1 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0,
+//         /* row_2 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 1.0, 0.5, 0.0, 0.0, 0.0, 0.0,
+//         /* row_3 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.5, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_4 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_5 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+//         /* row_6 */ 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0
+//     };
+//     // clang-format on
+//   }
+//   else
+//   {
+//     CHECK(false);
+//   }
+
+//   CHECK_THAT(transfer_matrix.to_dense(),
+//              RangeEquals(expected[rank],
+//                          [](auto a, auto b) {
+//                            return std::abs(a - b)
+//                                   <= std::numeric_limits<T>::epsilon();
+//                          }));
+// }
+
+TEMPLATE_TEST_CASE("Transfer Matrix 2D", "[transfer_matrix]", double)
+{
+  using T = TestType;
+
+  auto mesh_coarse
+      = std::make_shared<mesh::Mesh<T>>(dolfinx::mesh::create_rectangle<T>(
+          MPI_COMM_SELF, {{{0, 0}, {1, 1}}}, {1, 1}, mesh::CellType::triangle));
+  mesh_coarse->topology()->create_entities(1);
+
+  auto [mesh_fine, parent_cell, parent_facet]
+      = refinement::refine(*mesh_coarse, std::nullopt);
+
+  auto element = P1_element(basix::cell::type::triangle);
+
+  auto V_coarse = std::make_shared<fem::FunctionSpace<T>>(
+      fem::create_functionspace<T>(mesh_coarse, element));
+  auto V_fine
+      = std::make_shared<fem::FunctionSpace<T>>(fem::create_functionspace<T>(
+          std::make_shared<mesh::Mesh<T>>(mesh_fine), element));
+
+  mesh_fine.topology()->create_connectivity(1, 0);
+  mesh_fine.topology()->create_connectivity(0, 1);
+
+  std::vector<std::int64_t> inclusion_map
+      = multigrid::inclusion_mapping(*mesh_coarse, mesh_fine);
+  CHECK_THAT(inclusion_map, RangeEquals(std::vector<std::int64_t>{4, 1, 5, 8}));
+
+  la::SparsityPattern sp
+      = multigrid::create_sparsity_pattern(*V_coarse, *V_fine, inclusion_map);
+  la::MatrixCSR<T> transfer_matrix(std::move(sp), la::BlockMode::compact);
+  multigrid::assemble_transfer_matrix<T>(transfer_matrix.mat_set_values(),
+                                        *V_coarse, *V_fine, inclusion_map,
+                                        weight<T>);
+  transfer_matrix.scatter_rev();
+  std::vector<T> expected{0.5, 0.0, 0.5, 0.0, 1.0, 0.0, 0.5, 0.0, 0.0,
+                          0.5, 1.0, 0.5, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0,
+                          0.5, 0.0, 0.0, 0.5, 0.0, 1.0, 0.0, 0.5, 0.0,
+                          0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.5, 1.0};
+  CHECK_THAT(transfer_matrix.to_dense(),
+             RangeEquals(expected,
+                         [](auto a, auto b) {
+                           return std::abs(a - b)
+                                  <= std::numeric_limits<T>::epsilon();
+                         }));
+}
+
+TEMPLATE_TEST_CASE("Transfer Matrix 3D", "[transfer_matrix]", double)
+{
+  using T = TestType;
+
+  auto mesh_coarse = std::make_shared<mesh::Mesh<T>>(
+      dolfinx::mesh::create_box<T>(MPI_COMM_SELF, {{{0, 0, 0}, {1, 1, 1}}},
+                                   {1, 1, 1}, mesh::CellType::tetrahedron));
+  mesh_coarse->topology()->create_entities(1);
+
+  auto [mesh_fine, parent_cell, parent_facet]
+      = refinement::refine(*mesh_coarse, std::nullopt);
+
+  auto element = P1_element(basix::cell::type::tetrahedron);
+
+  auto V_coarse = std::make_shared<fem::FunctionSpace<T>>(
+      fem::create_functionspace<T>(mesh_coarse, element));
+  auto V_fine
+      = std::make_shared<fem::FunctionSpace<T>>(fem::create_functionspace<T>(
+          std::make_shared<mesh::Mesh<T>>(mesh_fine), element));
+
+  mesh_fine.topology()->create_connectivity(1, 0);
+  mesh_fine.topology()->create_connectivity(0, 1);
+
+  std::vector<std::int64_t> inclusion_map
+      = multigrid::inclusion_mapping(*mesh_coarse, mesh_fine);
+  CHECK_THAT(inclusion_map, RangeEquals(std::vector<std::int64_t>{
+                                0, 6, 15, 25, 17, 9, 11, 22}));
+
+  la::SparsityPattern sp
+      = multigrid::create_sparsity_pattern(*V_coarse, *V_fine, inclusion_map);
+  la::MatrixCSR<T> transfer_matrix(std::move(sp), la::BlockMode::compact);
+  multigrid::assemble_transfer_matrix<T>(transfer_matrix.mat_set_values(),
+                                        *V_coarse, *V_fine, inclusion_map,
+                                        weight<T>);
+
+  std::vector<T> expected{
+      1.0, 0.5, 0.5, 0.5, 0.5, 0.5, 0.0, 0.5, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+      0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.5,
+      0.5, 0.5, 0.0, 1.0, 0.0, 0.0, 0.0, 0.5, 0.0, 0.5, 0.0, 0.5, 0.0, 0.0, 0.0,
+      0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.5, 0.0, 0.0, 0.0,
+      0.0, 0.0, 0.0, 0.0, 0.5, 0.0, 0.0, 0.5, 0.5, 1.0, 0.0, 0.0, 0.0, 0.5, 0.0,
+      0.0, 0.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+      0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.5, 0.0, 0.5, 0.0, 0.5,
+      0.5, 1.0, 0.5, 0.0, 0.5, 0.0, 0.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0,
+      0.5, 0.0, 0.5, 0.0, 0.5, 1.0, 0.0, 0.0, 0.0, 0.5, 0.0, 0.0, 0.5, 0.0, 0.0,
+      0.0, 0.5, 0.5, 0.0, 0.0, 0.5, 0.0, 0.0, 0.5, 1.0, 0.0, 0.0, 0.0, 0.5, 0.0,
+      0.0, 0.0, 0.0, 0.5, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.0,
+      0.0, 0.5, 0.0, 0.0, 0.5, 0.5, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.0,
+      0.0, 0.0, 0.5, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.0, 0.0, 0.0, 0.0,
+      0.0, 0.0, 0.5, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.5, 0.5, 0.5,
+      0.5, 1.0, 0.0, 0.0, 0.0, 0.5};
+  CHECK_THAT(transfer_matrix.to_dense(),
+             RangeEquals(expected,
+                         [](auto a, auto b) {
+                           return std::abs(a - b)
+                                  <= std::numeric_limits<T>::epsilon();
+                         }));
+}
diff --git a/python/CMakeLists.txt b/python/CMakeLists.txt
index 0f440260214..5425d10896f 100644
--- a/python/CMakeLists.txt
+++ b/python/CMakeLists.txt
@@ -44,6 +44,7 @@ nanobind_add_module(
   dolfinx/wrappers/la.cpp
   dolfinx/wrappers/log.cpp
   dolfinx/wrappers/mesh.cpp
+  dolfinx/wrappers/multigrid.cpp
   dolfinx/wrappers/petsc.cpp
   dolfinx/wrappers/refinement.cpp
 )
diff --git a/python/dolfinx/transfer.py b/python/dolfinx/transfer.py
new file mode 100644
index 00000000000..7b9e8b59e7c
--- /dev/null
+++ b/python/dolfinx/transfer.py
@@ -0,0 +1,29 @@
+# Copyright (C) 2024 Paul T. Kühner
+#
+# This file is part of DOLFINx (https://www.fenicsproject.org)
+#
+# SPDX-License-Identifier:    LGPL-3.0-or-later
+
+import numpy as np
+from numpy.typing import NDArray
+
+from dolfinx.cpp.la import SparsityPattern
+from dolfinx.cpp.transfer import (
+    assemble_transfer_matrix,
+)
+from dolfinx.cpp.transfer import create_sparsity_pattern as _create_sparsity_pattern
+from dolfinx.cpp.transfer import inclusion_mapping as _inclusion_mapping
+from dolfinx.fem import FunctionSpace
+from dolfinx.mesh import Mesh
+
+__all__ = ["assemble_transfer_matrix", "create_sparsity_pattern", "inclusion_mapping"]
+
+
+def inclusion_mapping(mesh_from: Mesh, mesh_to: Mesh) -> NDArray[np.int64]:
+    return _inclusion_mapping(mesh_from._cpp_object, mesh_to._cpp_object)
+
+
+def create_sparsity_pattern(
+    V_from: FunctionSpace, V_to: FunctionSpace, inclusion_map: NDArray[np.int64]
+) -> SparsityPattern:
+    return _create_sparsity_pattern(V_from._cpp_object, V_to._cpp_object, inclusion_map)
diff --git a/python/dolfinx/wrappers/dolfinx.cpp b/python/dolfinx/wrappers/dolfinx.cpp
index 9c0864f9ef1..9d2ca3e716d 100644
--- a/python/dolfinx/wrappers/dolfinx.cpp
+++ b/python/dolfinx/wrappers/dolfinx.cpp
@@ -25,6 +25,8 @@ void la(nb::module_& m);
 void mesh(nb::module_& m);
 void nls(nb::module_& m);
 void refinement(nb::module_& m);
+void transfer(nb::module_& m);
+
 } // namespace dolfinx_wrappers
 
 NB_MODULE(cpp, m)
@@ -74,6 +76,10 @@ NB_MODULE(cpp, m)
   nb::module_ refinement = m.def_submodule("refinement", "Refinement module");
   dolfinx_wrappers::refinement(refinement);
 
+  // Create transfer submodule
+  nb::module_ transfer = m.def_submodule("transfer", "Transfer module");
+  dolfinx_wrappers::transfer(transfer);
+
 #if defined(HAS_PETSC) && defined(HAS_PETSC4PY)
   // PETSc-specific wrappers
   nb::module_ nls = m.def_submodule("nls", "Nonlinear solver module");
diff --git a/python/dolfinx/wrappers/multigrid.cpp b/python/dolfinx/wrappers/multigrid.cpp
new file mode 100644
index 00000000000..cabc8e53726
--- /dev/null
+++ b/python/dolfinx/wrappers/multigrid.cpp
@@ -0,0 +1,67 @@
+// Copyright (C) 2024 Paul T. Kühner
+//
+// This file is part of DOLFINX (https://www.fenicsproject.org)
+//
+// SPDX-License-Identifier:    LGPL-3.0-or-later
+
+#include <cstdint>
+#include <vector>
+
+#include <nanobind/nanobind.h>
+#include <nanobind/stl/vector.h>
+
+#include <dolfinx/la/MatrixCSR.h>
+#include <dolfinx/multigrid/transfer_matrix.h>
+
+#include "array.h"
+
+namespace nb = nanobind;
+
+namespace dolfinx_wrappers
+{
+
+void transfer(nb::module_& m)
+{
+  m.def(
+      "inclusion_mapping",
+      [](const dolfinx::mesh::Mesh<double>& mesh_from,
+         const dolfinx::mesh::Mesh<double>& mesh_to)
+      {
+        std::vector<std::int64_t> map
+            = dolfinx::multigrid::inclusion_mapping<double>(mesh_from, mesh_to);
+        return dolfinx_wrappers::as_nbarray(std::move(map));
+      },
+      nb::arg("mesh_from"), nb::arg("mesh_to"),
+      "Computes inclusion mapping between two meshes");
+
+  m.def(
+      "create_sparsity_pattern",
+      [](const dolfinx::fem::FunctionSpace<double>& V_from,
+         const dolfinx::fem::FunctionSpace<double>& V_to,
+         nb::ndarray<const std::int64_t, nb::numpy>& inclusion_map)
+      {
+        // TODO: cahnge to accepting range;
+        auto vec = std::vector(inclusion_map.data(),
+                               inclusion_map.data() + inclusion_map.size());
+        return dolfinx::multigrid::create_sparsity_pattern<double>(V_from, V_to,
+                                                                  vec);
+      },
+      nb::arg("V_from"), nb::arg("V_to"), nb::arg("inclusion_map"));
+
+  m.def(
+      "assemble_transfer_matrix",
+      [](dolfinx::la::MatrixCSR<double>& A,
+         const dolfinx::fem::FunctionSpace<double>& V_from,
+         const dolfinx::fem::FunctionSpace<double>& V_to,
+         const std::vector<std::int64_t>& inclusion_map,
+         std::function<double(std::int32_t)> weight)
+      {
+        dolfinx::multigrid::assemble_transfer_matrix(
+            A.mat_set_values(), V_from, V_to, inclusion_map, weight);
+      },
+      nb::arg("A"), nb::arg("V_from"), nb::arg("V_to"),
+      nb::arg("inclusion_map"), nb::arg("weight"),
+      "Assembles a transfer matrix between two function spaces.");
+}
+
+} // namespace dolfinx_wrappers
diff --git a/python/test/unit/transfer/test_transfer.py b/python/test/unit/transfer/test_transfer.py
new file mode 100644
index 00000000000..8d60dfb12e2
--- /dev/null
+++ b/python/test/unit/transfer/test_transfer.py
@@ -0,0 +1,63 @@
+# Copyright (C) 2024 Paul T. Kühner
+#
+# This file is part of DOLFINx (https://www.fenicsproject.org)
+#
+# SPDX-License-Identifier:    LGPL-3.0-or-later
+
+from mpi4py import MPI
+
+import pytest
+
+from dolfinx.fem import functionspace
+from dolfinx.mesh import GhostMode, create_interval, refine
+from dolfinx.transfer import create_sparsity_pattern, inclusion_mapping
+
+
+@pytest.mark.parametrize(
+    "ghost_mode", [GhostMode.none, GhostMode.shared_vertex, GhostMode.shared_facet]
+)
+def test_1d(ghost_mode):
+    mesh = create_interval(MPI.COMM_WORLD, 10, (0, 1), ghost_mode=ghost_mode)
+    mesh_fine, _, _ = refine(mesh)
+    inclusion_map = inclusion_mapping(mesh, mesh_fine)
+    V = functionspace(mesh, ("Lagrange", 1, (1,)))
+    V_fine = functionspace(mesh_fine, ("Lagrange", 1, (1,)))
+    assert V.element == V_fine.element
+    V_fine.mesh.topology.create_connectivity(1, 0)
+    V_fine.mesh.topology.create_connectivity(0, 1)
+    create_sparsity_pattern(V, V_fine, inclusion_map)
+    # T = matrix_csr(sp)
+    # assemble_transfer_matrix(
+    #     T._cpp_object,
+    #     V._cpp_object,
+    #     V_fine._cpp_object,
+    #     inclusion_map,
+    #     lambda i: 1.0 if i == 0 else 0.5,
+    # )
+    # continue with assembly of matrix
+
+
+@pytest.mark.parametrize("degree", [2, 3, 4])
+@pytest.mark.parametrize(
+    "ghost_mode", [GhostMode.none, GhostMode.shared_vertex, GhostMode.shared_facet]
+)
+def test_1d_higher_order(degree, ghost_mode):
+    mesh = create_interval(MPI.COMM_WORLD, 10, (0, 1), ghost_mode=ghost_mode)
+    mesh_fine, _, _ = refine(mesh)
+
+    # this is a strictly geometric operation, and thus should pass
+    inclusion_map = inclusion_mapping(mesh, mesh_fine)
+
+    V = functionspace(mesh, ("Lagrange", degree, (1,)))
+    V_fine = functionspace(mesh_fine, ("Lagrange", degree, (1,)))
+
+    V_fine.mesh.topology.create_connectivity(1, 0)
+    V_fine.mesh.topology.create_connectivity(0, 1)
+
+    # Check not supported throws
+    with pytest.raises(Exception):
+        create_sparsity_pattern(V, V_fine, inclusion_map)
+
+
+if __name__ == "__main__":
+    pytest.main([__file__])