From 5da5f66706dff4c7037f87fe574966161e5d4c25 Mon Sep 17 00:00:00 2001 From: Gabriel Dos Santos Date: Wed, 30 Jul 2025 15:57:08 +0200 Subject: [PATCH 1/5] fix(benchmarks): attempt to fix 2D stencil MBT(SV) versions --- benchmarks/stencil_2d.cpp | 210 ++++++++++++++++++++++++++++---------- 1 file changed, 156 insertions(+), 54 deletions(-) diff --git a/benchmarks/stencil_2d.cpp b/benchmarks/stencil_2d.cpp index 300a49c..d952f27 100644 --- a/benchmarks/stencil_2d.cpp +++ b/benchmarks/stencil_2d.cpp @@ -39,7 +39,7 @@ namespace { template requires IsMdSpan2D && IsMdSpan2D -auto naive(MdSpanIn in, MdSpanOut out, size_t D = DELTA) -> void { +auto naive(const MdSpanIn in, MdSpanOut out, size_t D = DELTA) -> void { assert(in.extent(0) == out.extent(0)); assert(in.extent(1) == out.extent(1)); assert(D != 0); @@ -61,9 +61,9 @@ auto naive(MdSpanIn in, MdSpanOut out, size_t D = DELTA) -> void { } } -template +template requires IsMdSpan2D && IsMdSpan2D -auto cache_blocked(MdSpanIn in, MdSpanOut out, size_t D = DELTA) -> void { +auto cache_blocked(const MdSpanIn in, MdSpanOut out) -> void { assert(in.extent(0) == out.extent(0)); assert(in.extent(1) == out.extent(1)); assert(D != 0); @@ -91,7 +91,7 @@ auto cache_blocked(MdSpanIn in, MdSpanOut out, size_t D = DELTA) -> void { template requires IsMdSpan2D && IsMdSpan2D -auto cache_blocked(MdSpanIn in, MdSpanOut out, size_t TM, size_t TN, size_t D = DELTA) -> void { +auto cache_blocked(const MdSpanIn in, MdSpanOut out, size_t TM, size_t TN, size_t D = DELTA) -> void { assert(in.extent(0) == out.extent(0)); assert(in.extent(1) == out.extent(1)); assert(D != 0); @@ -117,39 +117,84 @@ auto cache_blocked(MdSpanIn in, MdSpanOut out, size_t TM, size_t TN, size_t D = } } -template +template requires IsMdSpan2D && IsMdSpan2D -auto manual_blocked_tiled(MdSpanIn in, MdSpanOut out, size_t D = DELTA) -> void { +auto manual_blocked_tiled(const MdSpanIn in, MdSpanOut out) -> void { assert(in.extent(0) == out.extent(0)); assert(in.extent(1) == out.extent(1)); assert(D != 0); - const size_t M = out.extent(0); - const size_t N = out.extent(1); + using index_type = typename MdSpanIn::index_type; - auto hin = in.data_handle(); - auto hout = out.data_handle(); + const size_t M = out.extent(0); + const size_t N = out.extent(1); + constexpr index_type tile_size = in.mapping().tile_size(); + const index_type nb_elems_in_tile_row = tile_size * in.mapping().ntiles_in_extent(0); + + const auto hin = in.data_handle(); + auto hout = out.data_handle(); #if defined(LAYOUT_TILED_BENCH_ENABLE_OMP) #pragma omp parallel for schedule(static) #endif - for (size_t x = D; x < M - D; x += TM) { - for (size_t y = D; y < N - D; y += TN) { - const size_t outoff = in.mapping().tile_offset_outer(x, y); - for (size_t i = 0; i < std::min(TM, M - D - x); ++i) { - const size_t ti = TM * i; - const size_t ti_mD = TM * (i - D); - const size_t ti_pD = TM * (i + D); - for (size_t j = 0; j < std::min(TN, N - D - y); ++j) { - const size_t inoff = ti + j; - const size_t baseoff = outoff + inoff; - const size_t inoff_mD = ti_mD + j; - const size_t inoff_pD = ti_pD + j; - const size_t baseoff_mD = outoff + inoff_mD; - const size_t baseoff_pD = outoff + inoff_pD; + for (index_type x = D; x < M - D; x += TM) { + for (index_type y = D; y < N - D; y += TN) { + for (index_type i = x; i < std::min(x + TM, M - D); ++i) { + for (index_type j = y; j < std::min(y + TN, N - D); ++j) { + // Compute local i and j coordinates (used to determine whether we are at a tile boundary or not) + const index_type li = i % TM; + const index_type lj = j % TN; + + const index_type center_toff = in.mapping().tile_offset_outer(i, j); + const index_type center = center_toff + in.mapping().tile_offset_inner(i, j); + + // North neighbor: i-1, j + index_type north = 0; + if (li >= D) { + // Same tile + north = center - TN; + } else { + // Different tile + const index_type north_toff = center_toff - nb_elems_in_tile_row; + north = north_toff + in.mapping().tile_offset_inner(i, j); + } + + // South neighbor: i+1, j + index_type south = 0; + if (li < TM - D && i + D < M) { + // Same tile + south = center + TN; + } else { + // Different tile + const index_type south_toff = center_toff + nb_elems_in_tile_row; + south = south_toff + in.mapping().tile_offset_inner(i, j); + } + + // West neighbor: i, j-D + index_type west = 0; + if (lj >= D) { + // Same tile + west = center - D; + } else { + // Different tile + const index_type west_toff = center_toff - tile_size; + west = west_toff + in.mapping().tile_offset_inner(i, j); + } + + // East neighbor: i, j+D + index_type east; + if (lj < TN - D && j + D < N) { + // Same tile + east = center + D; + } else { + // Different tile + const index_type east_toff = center_toff + tile_size; + east = east_toff + in.mapping().tile_offset_inner(i, j); + } + // clang-format off - hout[baseoff] = hin[baseoff_mD] + - hin[baseoff-D] + hin[baseoff] + hin[baseoff+D] + - hin[baseoff_pD]; + hout[center] = ( hin[north] + + hin[west] + hin[center] + hin[east] + + hin[south] ) / 5.0; // clang-format on } } @@ -157,6 +202,7 @@ auto manual_blocked_tiled(MdSpanIn in, MdSpanOut out, size_t D = DELTA) -> void } } +// TODO(gdossantos): update the runtime-only version template requires IsMdSpan2D && IsMdSpan2D auto manual_blocked_tiled(MdSpanIn in, MdSpanOut out, size_t TM, size_t TN, size_t D = DELTA) -> void { @@ -197,34 +243,90 @@ auto manual_blocked_tiled(MdSpanIn in, MdSpanOut out, size_t TM, size_t TN, size } } -template +template requires IsMdSpan2D && IsMdSpan2D -auto manual_blocked_tiled_subview(MdSpanIn in, MdSpanOut out, size_t D = DELTA) -> void { - assert(in.extent(0) == out.extent(0)); - assert(in.extent(1) == out.extent(1)); - assert(D != 0); - - using InTile = std::mdspan, std::layout_right>; - using OutTile = std::mdspan, std::layout_right>; - - const size_t M = out.extent(0); - const size_t N = out.extent(1); - -#if defined(LAYOUT_TILED_BENCH_ENABLE_OMP) -#pragma omp parallel for schedule(static) -#endif - for (size_t x = 0; x < M; x += TM) { - for (size_t y = 0; y < N; y += TN) { - const size_t outoff = in.mapping().tile_offset_outer(x, y); - const auto t_in = InTile(in.data_handle() + outoff); - auto t_out = OutTile(out.data_handle() + outoff); - for (size_t i = 0; i < std::min(TM, M - x); ++i) { - for (size_t j = 0; j < std::min(TN, N - y); ++j) { - // clang-format off - t_out[i, j] = ( t_in[i-D, j] - + t_in[i, j-D] + t_in[ i, j] + t_in[i, j+D] - + t_in[i+D, j] ) / 5.0; - // clang-format on +auto manual_blocked_tiled_subview(const MdSpanIn in, MdSpanOut out) -> void { + using index_type = typename MdSpanIn::index_type; + using mapping_type = typename MdSpanIn::mapping_type; + using value_type = typename MdSpanIn::value_type; + + constexpr index_type tile_size = mapping_type::tile_size(); + + const index_type M = in.extent(0); + const index_type N = in.extent(1); + + const index_type ntiles_M = in.mapping().ntiles_in_extent(0); + const index_type ntiles_N = in.mapping().ntiles_in_extent(1); + + // Halo tile dimensions + constexpr index_type halo_TM = TM + 2 * D; + constexpr index_type halo_TN = TN + 2 * D; + + using HaloExts = std::extents; + + // Pre-allocate halo buffer + alignas(64) std::array halo_buf{}; + + for (index_type x = 0; x < ntiles_M; x += TM) { + for (index_type y = 0; y < ntiles_N; y += TN) { + for (index_type i = x; i < std::min(x + TM, ntiles_M); ++i) { + const index_type ti = i * TM; + const index_type ti_end = std::min(ti + TM, M); + for (index_type j = y; j < std::min(y + TN, ntiles_N); ++j) { + const index_type tj = j * TN; + const index_type tj_end = std::min(tj + TN, N); + + // Skip tiles without valid stencil region + if (ti + D >= M - D || tj + D >= N - D || ti_end <= D || tj_end <= D) { + continue; + } + + // Fill halo buffer + for (index_type hi = 0; hi < halo_TM; ++hi) { + const index_type global_i = ti + hi - D; + const index_type src_tile_i = global_i / TM; + const index_type src_local_i = global_i % TM; + const index_type halo_i = hi * halo_TN; + for (index_type hj = 0; hj < halo_TN; ++hj) { + const index_type global_j = tj + hj - D; + const index_type src_tile_j = global_j / TN; + const index_type src_local_j = global_j % TN; + const index_type src_tile_base = (src_tile_i * ntiles_N + src_tile_j) * tile_size; + const index_type src_offset = src_tile_base + src_local_i * TN + src_local_j; + const index_type halo_offset = halo_i + hj; + + halo_buf[halo_offset] = in.data_handle()[src_offset]; + } + } + + // Create local mdspan with layout_right over halo buffer + std::mdspan hs(halo_buf.data()); + for (index_type hi = D; hi < std::min(TM + D, ti_end - ti + D); ++hi) { + const index_type global_i = ti + hi - D; + const index_type dst_tile_i = global_i / TM; + const index_type dst_local_i = global_i % TM; + for (index_type hj = D; hj < std::min(TN + D, tj_end - tj + D); ++hj) { + // Global coordinates + const index_type global_j = tj + hj - D; + + // Skip global boundaries + if (global_i < D || global_i >= M - D || global_j < D || global_j >= N - D) { + continue; + } + + // Write to output - compute destination offset manually + const index_type dst_tile_j = global_j / TN; + const index_type dst_local_j = global_j % TN; + const index_type dst_tile_base = (dst_tile_i * ntiles_N + dst_tile_j) * tile_size; + const index_type dst_offset = dst_tile_base + dst_local_i * TN + dst_local_j; + + // clang-format off + out.data_handle()[dst_offset] = ( hs[hi-D, hj] + + hs[hi, hj-D] + hs[ hi, hj] + hs[hi, hj+D] + + hs[hi+D, hj] ) / 5.0; + // clang-format on + } + } } } } From c463280cc7326ac6038921546e8bfd946cf2074c Mon Sep 17 00:00:00 2001 From: Gabriel Dos Santos Date: Fri, 1 Aug 2025 17:28:42 +0200 Subject: [PATCH 2/5] fix(benchmarks): another attempt to fix, with diags Explain the process with additional diagrams in an attempt to fix the 2D stencil MBT impl Signed-off-by: Gabriel Dos Santos --- benchmarks/stencil_2d.cpp | 128 +++++++++++++++++++++++++++++--------- 1 file changed, 100 insertions(+), 28 deletions(-) diff --git a/benchmarks/stencil_2d.cpp b/benchmarks/stencil_2d.cpp index d952f27..7705929 100644 --- a/benchmarks/stencil_2d.cpp +++ b/benchmarks/stencil_2d.cpp @@ -126,8 +126,8 @@ auto manual_blocked_tiled(const MdSpanIn in, MdSpanOut out) -> void { using index_type = typename MdSpanIn::index_type; - const size_t M = out.extent(0); - const size_t N = out.extent(1); + const size_t M = in.extent(0); + const size_t N = in.extent(1); constexpr index_type tile_size = in.mapping().tile_size(); const index_type nb_elems_in_tile_row = tile_size * in.mapping().ntiles_in_extent(0); @@ -138,57 +138,126 @@ auto manual_blocked_tiled(const MdSpanIn in, MdSpanOut out) -> void { #endif for (index_type x = D; x < M - D; x += TM) { for (index_type y = D; y < N - D; y += TN) { + // Compute offset of first element in tile (outside domain if in first tile row/column) + const index_type center_toff = in.mapping().tile_offset_outer(x, y); + for (index_type i = x; i < std::min(x + TM, M - D); ++i) { - for (index_type j = y; j < std::min(y + TN, N - D); ++j) { - // Compute local i and j coordinates (used to determine whether we are at a tile boundary or not) - const index_type li = i % TM; - const index_type lj = j % TN; + const index_type li = i % TM; + const index_type center_row = li * TN; + const index_type north_row = (li - D + TM) * TN; + const index_type south_row = (li + D - TM) * TN; - const index_type center_toff = in.mapping().tile_offset_outer(i, j); - const index_type center = center_toff + in.mapping().tile_offset_inner(i, j); + for (index_type j = y; j < std::min(y + TN, N - D); ++j) { + const index_type lj = j % TN; + const index_type center = center_toff + center_row + lj; - // North neighbor: i-1, j + // North neighbor: i-D, j index_type north = 0; - if (li >= D) { + if (li - D >= 0 || x == 0) { // Same tile north = center - TN; } else { + // ◊ ◌ ¤ ↑ ↗ → ↘ ↓ ↙ ← ↖ ↔ ↕ ■ ▲ ▶ ▼ ◀ △ ▷ ▽ ◁ ╦ ╗ ╔ ═ ╩ ╝ ╚ ║ ╬ ╣ ╠ // Different tile - const index_type north_toff = center_toff - nb_elems_in_tile_row; - north = north_toff + in.mapping().tile_offset_inner(i, j); + // ┌─────────┬─────────┬─────────┬─────────┐ + // │ │△═╗2 │ │ │ If the target element (@) is in the tile above us (*), we do: + // │ │║ ║ 3 │ │ │ + // │ │║1▽══▶@ │ │ │ 1. Subtract $N_elems_per_tile \times N_tiles_in_row$ from the + // ├─────────┼║─────↑──┼─────────┼─────────┤ base tile offset (■); + // │ │■ ← * →│ │ │ 2. Add the pre-computed north_row offset; + // │ │ ↓ │ │ │ 3. Add the local j offset. + // │ │ │ │ │ + // ├─────────┼─────────┼─────────┼─────────┤ + // │ │ │ │ │ + // │ │ │ │ │ + // │ │ │ │ │ + // ├─────────┼─────────┼─────────┼─────────┤ + // │ │ │ │ │ + // │ │ │ │ │ + // │ │ │ │ │ + // └─────────┴─────────┴─────────┴─────────┘ + north = center_toff - nb_elems_in_tile_row + north_row + lj; } - // South neighbor: i+1, j + // South neighbor: i+D, j index_type south = 0; - if (li < TM - D && i + D < M) { + if (li + D < TM && i + D < M) { // Same tile south = center + TN; } else { // Different tile - const index_type south_toff = center_toff + nb_elems_in_tile_row; - south = south_toff + in.mapping().tile_offset_inner(i, j); + // ┌─────────┬─────────┬─────────┬─────────┐ + // │ │ │ │ │ If the target element (@) is in the tile below us (*), we do: + // │ │ │ │ │ + // │ │ │ │ │ 1. Add $N_elems_per_tile \times N_tiles_in_row$ from the + // ├─────────┼─────────┼─────────┼─────────┤ base tile offset (■); + // │ │■ │ │ │ 2. Add the pre-computed south_row offset (may not be 0 if + // │ │║ ↑ │ │ │ D > 1); + // │ │║1 ← * →│ │ │ 3. Add the local j offset. + // ├─────────┼║─────↓──┼─────────┼─────────┤ + // │ │▽════▶@ │ │ │ + // │ │ 2&3 │ │ │ + // │ │ │ │ │ + // ├─────────┼─────────┼─────────┼─────────┤ + // │ │ │ │ │ + // │ │ │ │ │ + // │ │ │ │ │ + // └─────────┴─────────┴─────────┴─────────┘ + south = center_toff + nb_elems_in_tile_row + south_row + lj; } // West neighbor: i, j-D index_type west = 0; - if (lj >= D) { + if (lj - D >= 0 || y == 0) { // Same tile west = center - D; } else { // Different tile - const index_type west_toff = center_toff - tile_size; - west = west_toff + in.mapping().tile_offset_inner(i, j); + // ┌─────────┬─────────┬─────────┬─────────┐ + // │ │ │ │ │ If the target element (@) is in the tile left of us (*), we + // │ │ │ │ │ do: + // │ │ │ │ │ + // ├────1────┼─────────┼─────────┼─────────┤ 1. Subtract $N_elems_per_tile$ from the base tile offset (■); + // │◁═════════■↑ │ │ │ 2. Add the pre-computed center_row offset; + // 2▽══════▶@← * → │ │ │ 3. Add the local j - D + TN offset. + // │ │ ↓ │ │ │ + // ├─────────┼─────────┼─────────┼─────────┤ + // │ │ │ │ │ + // │ │ │ │ │ + // │ │ │ │ │ + // ├─────────┼─────────┼─────────┼─────────┤ + // │ │ │ │ │ + // │ │ │ │ │ + // │ │ │ │ │ + // └─────────┴─────────┴─────────┴─────────┘ + west = center_toff - tile_size + center_row + (lj - D + TN); } // East neighbor: i, j+D index_type east; - if (lj < TN - D && j + D < N) { + if (lj - D < TN && j + D < N) { // Same tile east = center + D; } else { // Different tile - const index_type east_toff = center_toff + tile_size; - east = east_toff + in.mapping().tile_offset_inner(i, j); + // ┌─────────┬─────────┬─────────┬─────────┐ + // │ │ │ │ │ If the target element (@) is in the tile right of us (*), we + // │ │ │ │ │ do: + // │ │ │ │ │ + // ├─────────┼────1────┼2────────┼─────────┤ 1. Add $N_elems_per_tile$ from the base tile offset (■); + // │ │■══════↑══▼ 3 │ │ 2. Add the pre-computed center_row offset; + // │ │ ← * →@ │ │ 3. Add the local j + D - TN offset. + // │ │ ↓ │ │ │ + // ├─────────┼─────────┼─────────┼─────────┤ + // │ │ │ │ │ + // │ │ │ │ │ + // │ │ │ │ │ + // ├─────────┼─────────┼─────────┼─────────┤ + // │ │ │ │ │ + // │ │ │ │ │ + // │ │ │ │ │ + // └─────────┴─────────┴─────────┴─────────┘ + east = center_toff + tile_size + center_row + (lj + D - TN); } // clang-format off @@ -357,11 +426,9 @@ auto run_bench(const Config& cfg) -> int { // Initializations auto rng_naive = ankerl::nanobench::Rng(42); - auto rng_tiled_l1 = rng_naive.copy(); - auto rng_tiled_l2 = rng_naive.copy(); benchmark_utils::fill_2d_uniform01(mlr_in, rng_naive); - benchmark_utils::fill_2d_uniform01(mlrtr_l1_in, rng_tiled_l1); - benchmark_utils::fill_2d_uniform01(mlrtr_l2_in, rng_tiled_l2); + benchmark_utils::fill_2d_copy(mlrtr_l1_in, mlr_in); + benchmark_utils::fill_2d_copy(mlrtr_l2_in, mlr_in); #if defined(LAYOUT_TILED_BENCH_ENABLE_CHECKS) // Checks @@ -369,8 +436,11 @@ auto run_bench(const Config& cfg) -> int { const double big_o = 5.0 * n * n; int res = 0; - naive(mlr_in, mlr_out); - double ref = benchmark_utils::l2_norm(mlr_out); + auto ref_out = MemorySlice(cfg.n_ * cfg.n_); + auto mref_out = std::mdspan(ref_out.get(), cfg.n_, cfg.n_); + naive(mlr_in, mref_out); + double ref = benchmark_utils::l2_norm(mref_out); + benchmark_utils::fill_zeroes(mlr_out); cache_blocked(mlr_in, mlr_out); res = benchmark_utils::check(ref, benchmark_utils::l2_norm(mlr_out), "R_cb-L1", big_o); @@ -384,8 +454,10 @@ auto run_bench(const Config& cfg) -> int { manual_blocked_tiled(mlrtr_l1_in, mlrtr_l1_out); res = benchmark_utils::check(ref, benchmark_utils::l2_norm(mlrtr_l1_out), "RTR-L1_mbt", big_o); + // benchmark_utils::compare_2d_vals(mref_out, mlrtr_l1_out); manual_blocked_tiled(mlrtr_l2_in, mlrtr_l2_out); res = benchmark_utils::check(ref, benchmark_utils::l2_norm(mlrtr_l2_out), "RTR-L2_mbt", big_o); + manual_blocked_tiled_subview(mlrtr_l1_in, mlrtr_l1_out); res = benchmark_utils::check(ref, benchmark_utils::l2_norm(mlrtr_l1_out), "RTR-L1_mbtsv", big_o); manual_blocked_tiled_subview(mlrtr_l2_in, mlrtr_l2_out); From bea6f18afd155d5aa6a99112a3d194b81d7c62b9 Mon Sep 17 00:00:00 2001 From: Gabriel Dos Santos Date: Fri, 1 Aug 2025 17:29:13 +0200 Subject: [PATCH 3/5] fix(benchmarks): fix L2-norm calculation Signed-off-by: Gabriel Dos Santos --- benchmarks/utils.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/benchmarks/utils.hpp b/benchmarks/utils.hpp index 4dc3b85..835f744 100644 --- a/benchmarks/utils.hpp +++ b/benchmarks/utils.hpp @@ -102,7 +102,7 @@ template #pragma omp parallel for reduction(+ : sum) #endif for (size_t e = 0; e < m.size(); ++e) { - sum += ptr[e]; + sum += (ptr[e] * ptr[e]); } return std::sqrt(sum); } From d78240996542f8867cfa0229104fc7b88874aae8 Mon Sep 17 00:00:00 2001 From: Gabriel Dos Santos Date: Fri, 1 Aug 2025 17:29:51 +0200 Subject: [PATCH 4/5] feat(benchmarks): add compare values helper for 2D mdspans Signed-off-by: Gabriel Dos Santos --- benchmarks/utils.hpp | 20 ++++++++++++++++++++ 1 file changed, 20 insertions(+) diff --git a/benchmarks/utils.hpp b/benchmarks/utils.hpp index 835f744..44668b2 100644 --- a/benchmarks/utils.hpp +++ b/benchmarks/utils.hpp @@ -107,6 +107,26 @@ template return std::sqrt(sum); } +template + requires std::is_floating_point_v and std::is_floating_point_v +auto compare_2d_vals(const MdSpanRef ref, const MdSpanGot got) -> void { + assert(ref.extent(0) == got.extent(0)); + assert(ref.extent(1) == got.extent(1)); + + const size_t M = ref.extent(0); + const size_t N = ref.extent(1); + + for (size_t i = 0; i < M; ++i) { + for (size_t j = 0; j < N; ++j) { + auto r = ref[i, j]; + auto g = got[i, j]; + if (std::abs(g - r) > std::numeric_limits::epsilon()) { + fmt::print(stderr, "({:03}, {:03}): expected {:.18f}, got {:.18f}\n", i, j, r, g); + } + } + } +} + /// Fills an `mdspan` with zeroes. template requires std::is_floating_point_v From 314aad8cf716634342f901b8cb5562e6cfe7650a Mon Sep 17 00:00:00 2001 From: Gabriel Dos Santos Date: Fri, 1 Aug 2025 17:35:09 +0200 Subject: [PATCH 5/5] fix(benchmarks): add explanations for control flow Signed-off-by: Gabriel Dos Santos --- benchmarks/stencil_2d.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/benchmarks/stencil_2d.cpp b/benchmarks/stencil_2d.cpp index 7705929..7f585a8 100644 --- a/benchmarks/stencil_2d.cpp +++ b/benchmarks/stencil_2d.cpp @@ -154,10 +154,9 @@ auto manual_blocked_tiled(const MdSpanIn in, MdSpanOut out) -> void { // North neighbor: i-D, j index_type north = 0; if (li - D >= 0 || x == 0) { - // Same tile + // Same tile or first row of tiles north = center - TN; } else { - // ◊ ◌ ¤ ↑ ↗ → ↘ ↓ ↙ ← ↖ ↔ ↕ ■ ▲ ▶ ▼ ◀ △ ▷ ▽ ◁ ╦ ╗ ╔ ═ ╩ ╝ ╚ ║ ╬ ╣ ╠ // Different tile // ┌─────────┬─────────┬─────────┬─────────┐ // │ │△═╗2 │ │ │ If the target element (@) is in the tile above us (*), we do: @@ -209,7 +208,7 @@ auto manual_blocked_tiled(const MdSpanIn in, MdSpanOut out) -> void { // West neighbor: i, j-D index_type west = 0; if (lj - D >= 0 || y == 0) { - // Same tile + // Same tile or first column of tiles west = center - D; } else { // Different tile