From a68f59538b2b2a62dd4b27cb22e32fc675399cf5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 17 Aug 2022 19:56:20 +0200 Subject: [PATCH 01/27] added neighbor table and a test --- include/gridtools/fn/sid_neighbor_table.hpp | 42 ++++++++++++++++ tests/unit_tests/fn/CMakeLists.txt | 1 + .../fn/test_fn_sid_neighbor_table.cpp | 49 +++++++++++++++++++ 3 files changed, 92 insertions(+) create mode 100644 include/gridtools/fn/sid_neighbor_table.hpp create mode 100644 tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp new file mode 100644 index 000000000..1b69894b8 --- /dev/null +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -0,0 +1,42 @@ +/* + * GridTools + * + * Copyright (c) 2014-2022, ETH Zurich + * All rights reserved. + * + * Please, refer to the LICENSE file in the root directory. + * SPDX-License-Identifier: BSD-3-Clause + */ +#pragma once + +#include +#include + +namespace gridtools::fn::sid_neighbor_table { + namespace sid_neighbor_table_impl_ { + template + struct sid_neighbor_table { + Sid sid; + }; + + template + auto neighbor_table_neighbors(sid_neighbor_table const &table, size_t index) { + using element_type = sid::element_type; + const auto ptr = sid_get_origin(table.sid); + const auto lower_bounds = sid_get_lower_bounds(table.sid); + const auto upper_bounds = sid_get_upper_bounds(table.sid); + const auto strides = sid_get_strides(table.sid); + std::array neighbors; + std::fill(std::begin(neighbors), std::end(neighbors), element_type{-1}); + return neighbors; + } + + template + auto as_neighbor_table(Sid sid) { + return sid_neighbor_table{std::move(sid)}; + } + } // namespace sid_neighbor_table_impl_ + + using sid_neighbor_table_impl_::as_neighbor_table; + +} // namespace gridtools::fn::sid_neighbor_table \ No newline at end of file diff --git a/tests/unit_tests/fn/CMakeLists.txt b/tests/unit_tests/fn/CMakeLists.txt index 9d920a78f..c9ca3601c 100644 --- a/tests/unit_tests/fn/CMakeLists.txt +++ b/tests/unit_tests/fn/CMakeLists.txt @@ -7,6 +7,7 @@ gridtools_add_unit_test(test_fn_run SOURCES test_fn_run.cpp) gridtools_add_unit_test(test_fn_column_stage SOURCES test_fn_column_stage.cpp) gridtools_add_unit_test(test_fn_stencil_stage SOURCES test_fn_stencil_stage.cpp LABELS fn) gridtools_add_unit_test(test_fn_unstructured SOURCES test_fn_unstructured.cpp LABELS fn) +gridtools_add_unit_test(test_fn_sid_neighbor_table SOURCES test_fn_sid_neighbor_table.cpp LABELS fn) if(TARGET _gridtools_cuda) gridtools_add_unit_test(test_fn_backend_gpu_cuda diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp new file mode 100644 index 000000000..af55eb04d --- /dev/null +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -0,0 +1,49 @@ +/* + * GridTools + * + * Copyright (c) 2014-2022, ETH Zurich + * All rights reserved. + * + * Please, refer to the LICENSE file in the root directory. + * SPDX-License-Identifier: BSD-3-Clause + */ + +#include + +#include +#include +#include +#include + +#include +#include +#include +#include + +using namespace gridtools; +using namespace fn; +using sid_neighbor_table::as_neighbor_table; + +TEST(sid_neighbor_table, correctness) { + constexpr size_t numElements = 3; + constexpr size_t numNeighbors = 2; + std::array data = {0, 1, 10, 11, 20, 21}; + using dim_hymap_t = hymap::keys; + auto contents = sid::synthetic() + .set(sid::host_device::simple_ptr_holder(data.data())) + .set(dim_hymap_t::make_values(0, 0)) + .set(dim_hymap_t::make_values(numElements, numNeighbors)) + .set(dim_hymap_t::make_values(3, 1)); + const auto table = as_neighbor_table(contents); + sid_get_origin(contents); + + auto [n00, n01] = neighbor_table_neighbors(table, 0); + auto [n10, n11] = neighbor_table_neighbors(table, 1); + auto [n20, n21] = neighbor_table_neighbors(table, 2); + EXPECT_EQ(n00, 0); + EXPECT_EQ(n01, 1); + EXPECT_EQ(n10, 10); + EXPECT_EQ(n11, 11); + EXPECT_EQ(n20, 20); + EXPECT_EQ(n21, 21); +} \ No newline at end of file From 9c5a26b2de2784cc68fd18852b88efb69f213ca3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Thu, 18 Aug 2022 11:23:54 +0200 Subject: [PATCH 02/27] proper reading of elements --- include/gridtools/fn/sid_neighbor_table.hpp | 33 ++++++++++++++----- .../fn/test_fn_sid_neighbor_table.cpp | 12 +++---- 2 files changed, 30 insertions(+), 15 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index 1b69894b8..7d14f44a4 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -9,31 +9,46 @@ */ #pragma once +#include +#include +#include #include #include namespace gridtools::fn::sid_neighbor_table { namespace sid_neighbor_table_impl_ { - template + template struct sid_neighbor_table { Sid sid; }; - template - auto neighbor_table_neighbors(sid_neighbor_table const &table, size_t index) { + template + auto neighbor_table_neighbors( + sid_neighbor_table const &table, size_t index) { using element_type = sid::element_type; + const auto ptr = sid_get_origin(table.sid); - const auto lower_bounds = sid_get_lower_bounds(table.sid); - const auto upper_bounds = sid_get_upper_bounds(table.sid); const auto strides = sid_get_strides(table.sid); - std::array neighbors; - std::fill(std::begin(neighbors), std::end(neighbors), element_type{-1}); + + const auto index_stride = at_key(strides); + const auto neighbour_stride = at_key(strides); + + gridtools::array neighbors; + for (int32_t elementIdx = 0; elementIdx < MaxNumNeighbors; ++elementIdx) { + const auto element_ptr = ptr + index * index_stride + elementIdx * neighbour_stride; + neighbors[elementIdx] = *element_ptr(); + } return neighbors; } - template + template auto as_neighbor_table(Sid sid) { - return sid_neighbor_table{std::move(sid)}; + static_assert(gridtools::tuple_util::size()))>::value == 2, + "Neighbor tables must have exactly two dimensions: the index dimension and the neighbor dimension"); + static_assert(!std::is_same_v, + "The index dimension and the neighbor dimension must be different."); + + return sid_neighbor_table{std::move(sid)}; } } // namespace sid_neighbor_table_impl_ diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index af55eb04d..69026e6e8 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -24,18 +24,18 @@ using namespace gridtools; using namespace fn; using sid_neighbor_table::as_neighbor_table; +using edge_dim_t = unstructured::dim::horizontal; +using edge_to_cell_dim_t = struct {}; + TEST(sid_neighbor_table, correctness) { constexpr size_t numElements = 3; constexpr size_t numNeighbors = 2; std::array data = {0, 1, 10, 11, 20, 21}; - using dim_hymap_t = hymap::keys; + using dim_hymap_t = hymap::keys; auto contents = sid::synthetic() .set(sid::host_device::simple_ptr_holder(data.data())) - .set(dim_hymap_t::make_values(0, 0)) - .set(dim_hymap_t::make_values(numElements, numNeighbors)) - .set(dim_hymap_t::make_values(3, 1)); - const auto table = as_neighbor_table(contents); - sid_get_origin(contents); + .set(dim_hymap_t::make_values(numNeighbors, 1)); + const auto table = as_neighbor_table(contents); auto [n00, n01] = neighbor_table_neighbors(table, 0); auto [n10, n11] = neighbor_table_neighbors(table, 1); From 9f2702258303632844f42daaa290f87a4d6b7b6e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 13:16:42 +0200 Subject: [PATCH 03/27] includes and size_t --- include/gridtools/fn/sid_neighbor_table.hpp | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index 7d14f44a4..fae323c3e 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -9,22 +9,23 @@ */ #pragma once -#include -#include -#include -#include +#include #include +#include "../common/array.hpp" +#include "../fn/unstructured.hpp" +#include "../sid/concept.hpp" + namespace gridtools::fn::sid_neighbor_table { namespace sid_neighbor_table_impl_ { - template + template struct sid_neighbor_table { Sid sid; }; - template + template auto neighbor_table_neighbors( - sid_neighbor_table const &table, size_t index) { + sid_neighbor_table const &table, std::size_t index) { using element_type = sid::element_type; const auto ptr = sid_get_origin(table.sid); @@ -34,7 +35,7 @@ namespace gridtools::fn::sid_neighbor_table { const auto neighbour_stride = at_key(strides); gridtools::array neighbors; - for (int32_t elementIdx = 0; elementIdx < MaxNumNeighbors; ++elementIdx) { + for (std::size_t elementIdx = 0; elementIdx < MaxNumNeighbors; ++elementIdx) { const auto element_ptr = ptr + index * index_stride + elementIdx * neighbour_stride; neighbors[elementIdx] = *element_ptr(); } From d93bf1f9b397b46a908da6d66261919e886f6b11 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 13:18:31 +0200 Subject: [PATCH 04/27] more includes and size_t --- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index 69026e6e8..8888b002e 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -10,15 +10,15 @@ #include -#include -#include -#include -#include - #include #include #include + #include +#include +#include +#include +#include using namespace gridtools; using namespace fn; @@ -28,8 +28,8 @@ using edge_dim_t = unstructured::dim::horizontal; using edge_to_cell_dim_t = struct {}; TEST(sid_neighbor_table, correctness) { - constexpr size_t numElements = 3; - constexpr size_t numNeighbors = 2; + constexpr std::size_t numElements = 3; + constexpr std::size_t numNeighbors = 2; std::array data = {0, 1, 10, 11, 20, 21}; using dim_hymap_t = hymap::keys; auto contents = sid::synthetic() From 813a41f6911be050c91b8d03268a7a6160fd0ca9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 13:20:38 +0200 Subject: [PATCH 05/27] includes --- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index 8888b002e..1f3b223c1 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -8,12 +8,12 @@ * SPDX-License-Identifier: BSD-3-Clause */ -#include - #include #include #include +#include + #include #include #include From ae473c28bc040168606c0e64b6d6ef94114730d7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 13:21:31 +0200 Subject: [PATCH 06/27] includes --- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index 1f3b223c1..761f4fa1b 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -8,13 +8,14 @@ * SPDX-License-Identifier: BSD-3-Clause */ +#include + #include #include #include #include -#include #include #include #include From 47c839378086b7da4f97ca594106b91c20f9e00c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 13:24:01 +0200 Subject: [PATCH 07/27] test namespaces --- .../fn/test_fn_sid_neighbor_table.cpp | 58 ++++++++++--------- 1 file changed, 31 insertions(+), 27 deletions(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index 761f4fa1b..0ff0b6349 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -21,30 +21,34 @@ #include #include -using namespace gridtools; -using namespace fn; -using sid_neighbor_table::as_neighbor_table; - -using edge_dim_t = unstructured::dim::horizontal; -using edge_to_cell_dim_t = struct {}; - -TEST(sid_neighbor_table, correctness) { - constexpr std::size_t numElements = 3; - constexpr std::size_t numNeighbors = 2; - std::array data = {0, 1, 10, 11, 20, 21}; - using dim_hymap_t = hymap::keys; - auto contents = sid::synthetic() - .set(sid::host_device::simple_ptr_holder(data.data())) - .set(dim_hymap_t::make_values(numNeighbors, 1)); - const auto table = as_neighbor_table(contents); - - auto [n00, n01] = neighbor_table_neighbors(table, 0); - auto [n10, n11] = neighbor_table_neighbors(table, 1); - auto [n20, n21] = neighbor_table_neighbors(table, 2); - EXPECT_EQ(n00, 0); - EXPECT_EQ(n01, 1); - EXPECT_EQ(n10, 10); - EXPECT_EQ(n11, 11); - EXPECT_EQ(n20, 20); - EXPECT_EQ(n21, 21); -} \ No newline at end of file +namespace gridtools::fn { + namespace { + + using sid_neighbor_table::as_neighbor_table; + + using edge_dim_t = unstructured::dim::horizontal; + using edge_to_cell_dim_t = struct {}; + + TEST(sid_neighbor_table, correctness) { + constexpr std::size_t numElements = 3; + constexpr std::size_t numNeighbors = 2; + std::array data = {0, 1, 10, 11, 20, 21}; + using dim_hymap_t = hymap::keys; + auto contents = sid::synthetic() + .set(sid::host_device::simple_ptr_holder(data.data())) + .set(dim_hymap_t::make_values(numNeighbors, 1)); + const auto table = as_neighbor_table(contents); + + auto [n00, n01] = neighbor_table_neighbors(table, 0); + auto [n10, n11] = neighbor_table_neighbors(table, 1); + auto [n20, n21] = neighbor_table_neighbors(table, 2); + EXPECT_EQ(n00, 0); + EXPECT_EQ(n01, 1); + EXPECT_EQ(n10, 10); + EXPECT_EQ(n11, 11); + EXPECT_EQ(n20, 20); + EXPECT_EQ(n21, 21); + } + + } // namespace +} // namespace gridtools::fn \ No newline at end of file From d1f99ef30bf622b4b8fc0d97d8274e064b477dae Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 13:57:56 +0200 Subject: [PATCH 08/27] use sid::shift --- include/gridtools/fn/sid_neighbor_table.hpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index fae323c3e..fa8b25469 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -25,19 +25,19 @@ namespace gridtools::fn::sid_neighbor_table { template auto neighbor_table_neighbors( - sid_neighbor_table const &table, std::size_t index) { + sid_neighbor_table const &table, + std::size_t index) { using element_type = sid::element_type; - const auto ptr = sid_get_origin(table.sid); + auto ptr = sid_get_origin(table.sid)(); const auto strides = sid_get_strides(table.sid); - const auto index_stride = at_key(strides); - const auto neighbour_stride = at_key(strides); - gridtools::array neighbors; + + sid::shift(ptr, sid::get_stride(strides), index); for (std::size_t elementIdx = 0; elementIdx < MaxNumNeighbors; ++elementIdx) { - const auto element_ptr = ptr + index * index_stride + elementIdx * neighbour_stride; - neighbors[elementIdx] = *element_ptr(); + neighbors[elementIdx] = *ptr; + sid::shift(ptr, sid::get_stride(strides), 1); } return neighbors; } From dd7406991d73dd6af043065acd5cfbd0ca599093 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 16:08:24 +0200 Subject: [PATCH 09/27] use C arrays instead of std::array --- include/gridtools/fn/sid_neighbor_table.hpp | 10 +++++----- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 10 +++------- 2 files changed, 8 insertions(+), 12 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index fa8b25469..157131733 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -29,8 +29,8 @@ namespace gridtools::fn::sid_neighbor_table { std::size_t index) { using element_type = sid::element_type; - auto ptr = sid_get_origin(table.sid)(); - const auto strides = sid_get_strides(table.sid); + auto ptr = sid::get_origin(table.sid)(); + const auto strides = sid::get_strides(table.sid); gridtools::array neighbors; @@ -43,13 +43,13 @@ namespace gridtools::fn::sid_neighbor_table { } template - auto as_neighbor_table(Sid sid) { - static_assert(gridtools::tuple_util::size()))>::value == 2, + auto as_neighbor_table(Sid &&sid) { + static_assert(gridtools::tuple_util::size()))>::value == 2, "Neighbor tables must have exactly two dimensions: the index dimension and the neighbor dimension"); static_assert(!std::is_same_v, "The index dimension and the neighbor dimension must be different."); - return sid_neighbor_table{std::move(sid)}; + return sid_neighbor_table{std::forward(sid)}; } } // namespace sid_neighbor_table_impl_ diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index 0ff0b6349..2e5538c19 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -26,17 +26,13 @@ namespace gridtools::fn { using sid_neighbor_table::as_neighbor_table; - using edge_dim_t = unstructured::dim::horizontal; - using edge_to_cell_dim_t = struct {}; + using edge_dim_t = integral_constant; + using edge_to_cell_dim_t = integral_constant; TEST(sid_neighbor_table, correctness) { constexpr std::size_t numElements = 3; constexpr std::size_t numNeighbors = 2; - std::array data = {0, 1, 10, 11, 20, 21}; - using dim_hymap_t = hymap::keys; - auto contents = sid::synthetic() - .set(sid::host_device::simple_ptr_holder(data.data())) - .set(dim_hymap_t::make_values(numNeighbors, 1)); + std::int32_t contents[numElements][numNeighbors] = {{0, 1}, {10, 11}, {20, 21}}; const auto table = as_neighbor_table(contents); auto [n00, n01] = neighbor_table_neighbors(table, 0); From fa40976d6d83f2849446acbf86770c9a700363c8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 16:25:51 +0200 Subject: [PATCH 10/27] code style --- include/gridtools/fn/sid_neighbor_table.hpp | 4 ++-- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 8 ++++---- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index 157131733..f1c7451a4 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -35,8 +35,8 @@ namespace gridtools::fn::sid_neighbor_table { gridtools::array neighbors; sid::shift(ptr, sid::get_stride(strides), index); - for (std::size_t elementIdx = 0; elementIdx < MaxNumNeighbors; ++elementIdx) { - neighbors[elementIdx] = *ptr; + for (std::size_t element_idx = 0; element_idx < MaxNumNeighbors; ++element_idx) { + neighbors[element_idx] = *ptr; sid::shift(ptr, sid::get_stride(strides), 1); } return neighbors; diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index 2e5538c19..22036cdec 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -30,10 +30,10 @@ namespace gridtools::fn { using edge_to_cell_dim_t = integral_constant; TEST(sid_neighbor_table, correctness) { - constexpr std::size_t numElements = 3; - constexpr std::size_t numNeighbors = 2; - std::int32_t contents[numElements][numNeighbors] = {{0, 1}, {10, 11}, {20, 21}}; - const auto table = as_neighbor_table(contents); + constexpr std::size_t num_elements = 3; + constexpr std::size_t num_neighbors = 2; + std::int32_t contents[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; + const auto table = as_neighbor_table(contents); auto [n00, n01] = neighbor_table_neighbors(table, 0); auto [n10, n11] = neighbor_table_neighbors(table, 1); From 6fb6a86a143fbc8c5cce589a1e4b73a93e08cc9c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Wed, 24 Aug 2022 16:40:39 +0200 Subject: [PATCH 11/27] changed implementation to use ptr holder and strides instead of whole sid --- include/gridtools/fn/sid_neighbor_table.hpp | 35 +++++++++++++++------ 1 file changed, 25 insertions(+), 10 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index f1c7451a4..a9458bc7c 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -18,26 +18,34 @@ namespace gridtools::fn::sid_neighbor_table { namespace sid_neighbor_table_impl_ { - template + template struct sid_neighbor_table { - Sid sid; + PtrHolder origin; + Strides strides; }; - template + template auto neighbor_table_neighbors( - sid_neighbor_table const &table, + sid_neighbor_table const &table, std::size_t index) { - using element_type = sid::element_type; - auto ptr = sid::get_origin(table.sid)(); - const auto strides = sid::get_strides(table.sid); + auto ptr = table.origin(); + using element_type = std::remove_reference_t; gridtools::array neighbors; - sid::shift(ptr, sid::get_stride(strides), index); + sid::shift(ptr, sid::get_stride(table.strides), index); for (std::size_t element_idx = 0; element_idx < MaxNumNeighbors; ++element_idx) { neighbors[element_idx] = *ptr; - sid::shift(ptr, sid::get_stride(strides), 1); + sid::shift(ptr, sid::get_stride(table.strides), 1); } return neighbors; } @@ -49,7 +57,14 @@ namespace gridtools::fn::sid_neighbor_table { static_assert(!std::is_same_v, "The index dimension and the neighbor dimension must be different."); - return sid_neighbor_table{std::forward(sid)}; + const auto origin = sid::get_origin(sid); + const auto strides = sid::get_strides(sid); + + return sid_neighbor_table, + sid::strides_type>{origin, strides}; } } // namespace sid_neighbor_table_impl_ From 0bbbdc737bc648bd3374330e3fc210ed4a66d7c6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Thu, 25 Aug 2022 10:28:23 +0200 Subject: [PATCH 12/27] make ADL function a friend --- include/gridtools/fn/sid_neighbor_table.hpp | 31 +++++++++------------ 1 file changed, 13 insertions(+), 18 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index a9458bc7c..28078b906 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -26,29 +26,24 @@ namespace gridtools::fn::sid_neighbor_table { struct sid_neighbor_table { PtrHolder origin; Strides strides; - }; - template - auto neighbor_table_neighbors( - sid_neighbor_table const &table, - std::size_t index) { + friend auto neighbor_table_neighbors( + sid_neighbor_table const &table, + std::size_t index) { - auto ptr = table.origin(); - using element_type = std::remove_reference_t; + auto ptr = table.origin(); + using element_type = std::remove_reference_t; - gridtools::array neighbors; + gridtools::array neighbors; - sid::shift(ptr, sid::get_stride(table.strides), index); - for (std::size_t element_idx = 0; element_idx < MaxNumNeighbors; ++element_idx) { - neighbors[element_idx] = *ptr; - sid::shift(ptr, sid::get_stride(table.strides), 1); + sid::shift(ptr, sid::get_stride(table.strides), index); + for (std::size_t element_idx = 0; element_idx < MaxNumNeighbors; ++element_idx) { + neighbors[element_idx] = *ptr; + sid::shift(ptr, sid::get_stride(table.strides), 1); + } + return neighbors; } - return neighbors; - } + }; template auto as_neighbor_table(Sid &&sid) { From ab55a320795dc774a125add6efafc3614d66cbf9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Thu, 25 Aug 2022 10:56:36 +0200 Subject: [PATCH 13/27] test on cuda --- include/gridtools/fn/sid_neighbor_table.hpp | 2 +- tests/unit_tests/fn/CMakeLists.txt | 4 ++ .../fn/test_fn_sid_neighbor_table.cu | 53 +++++++++++++++++++ 3 files changed, 58 insertions(+), 1 deletion(-) create mode 100644 tests/unit_tests/fn/test_fn_sid_neighbor_table.cu diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index 28078b906..a04b2c84d 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -27,7 +27,7 @@ namespace gridtools::fn::sid_neighbor_table { PtrHolder origin; Strides strides; - friend auto neighbor_table_neighbors( + GT_FUNCTION friend auto neighbor_table_neighbors( sid_neighbor_table const &table, std::size_t index) { diff --git a/tests/unit_tests/fn/CMakeLists.txt b/tests/unit_tests/fn/CMakeLists.txt index c9ca3601c..bd2d59c84 100644 --- a/tests/unit_tests/fn/CMakeLists.txt +++ b/tests/unit_tests/fn/CMakeLists.txt @@ -30,4 +30,8 @@ if(TARGET _gridtools_cuda) SOURCES test_extents.cu LIBRARIES _gridtools_cuda LABELS cuda fn) + gridtools_add_unit_test(test_fn_sid_neighbor_table_cuda + SOURCES test_fn_sid_neighbor_table.cu + LIBRARIES _gridtools_cuda + LABELS cuda fn) endif() diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu new file mode 100644 index 000000000..70802fa6f --- /dev/null +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu @@ -0,0 +1,53 @@ +/* + * GridTools + * + * Copyright (c) 2014-2021, ETH Zurich + * All rights reserved. + * + * Please, refer to the LICENSE file in the root directory. + * SPDX-License-Identifier: BSD-3-Clause + */ + +#include + +#include +#include +#include +#include + +#include + +#include + +namespace gridtools::fn { + namespace { + using sid_neighbor_table::as_neighbor_table; + + using edge_dim_t = integral_constant; + using edge_to_cell_dim_t = integral_constant; + + template + __device__ auto neighbor_table_neighbors_device(const Table &table, size_t index) -> array { + return neighbor_table_neighbors(table, index); + } + constexpr std::size_t num_elements = 3; + constexpr std::size_t num_neighbors = 2; + __device__ std::int32_t contents[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; + + TEST(sid_neighbor_table, correctness_cuda) { + const auto table = as_neighbor_table(contents); + + const auto instantiation = &neighbor_table_neighbors_device>; + + auto [n00, n01] = on_device::exec(GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(instantiation), table, 0); + auto [n10, n11] = on_device::exec(GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(instantiation), table, 1); + auto [n20, n21] = on_device::exec(GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(instantiation), table, 2); + EXPECT_EQ(n00, 0); + EXPECT_EQ(n01, 1); + EXPECT_EQ(n10, 10); + EXPECT_EQ(n11, 11); + EXPECT_EQ(n20, 20); + EXPECT_EQ(n21, 21); + } + } // namespace +} // namespace gridtools::fn \ No newline at end of file From fa522e318ea54ccf5e2cbf91d9d79144addaf8a2 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Thu, 25 Aug 2022 11:02:57 +0200 Subject: [PATCH 14/27] remove unused includes --- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index 22036cdec..684ed9024 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -16,11 +16,6 @@ #include -#include -#include -#include -#include - namespace gridtools::fn { namespace { From d9569af9cc10d92f2b0eecf1a04ef3d32d160f02 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Thu, 25 Aug 2022 11:21:53 +0200 Subject: [PATCH 15/27] cuda playing funny games --- .../fn/test_fn_sid_neighbor_table.cu | 24 +++++++++++-------- 1 file changed, 14 insertions(+), 10 deletions(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu index 70802fa6f..2f46bdbe0 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu @@ -26,22 +26,26 @@ namespace gridtools::fn { using edge_dim_t = integral_constant; using edge_to_cell_dim_t = integral_constant; - template - __device__ auto neighbor_table_neighbors_device(const Table &table, size_t index) -> array { - return neighbor_table_neighbors(table, index); - } constexpr std::size_t num_elements = 3; constexpr std::size_t num_neighbors = 2; __device__ std::int32_t contents[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; + template + __device__ auto neighbor_table_neighbors_device(const Table &table, std::size_t index) + -> array { + return neighbor_table_neighbors(table, index); + } + TEST(sid_neighbor_table, correctness_cuda) { const auto table = as_neighbor_table(contents); - - const auto instantiation = &neighbor_table_neighbors_device>; - - auto [n00, n01] = on_device::exec(GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(instantiation), table, 0); - auto [n10, n11] = on_device::exec(GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(instantiation), table, 1); - auto [n20, n21] = on_device::exec(GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(instantiation), table, 2); + using table_t = std::decay_t; + + auto [n00, n01] = on_device::exec( + GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(&neighbor_table_neighbors_device), table, 0); + auto [n10, n11] = on_device::exec( + GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(&neighbor_table_neighbors_device), table, 1); + auto [n20, n21] = on_device::exec( + GT_MAKE_INTEGRAL_CONSTANT_FROM_VALUE(&neighbor_table_neighbors_device), table, 2); EXPECT_EQ(n00, 0); EXPECT_EQ(n01, 1); EXPECT_EQ(n10, 10); From b67bdf83c2b489427346ce4751abeaadd07d8c77 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Thu, 25 Aug 2022 12:14:52 +0200 Subject: [PATCH 16/27] use dynamic device memory --- .../unit_tests/fn/test_fn_sid_neighbor_table.cu | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu index 2f46bdbe0..bf3434fb1 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu @@ -18,6 +18,7 @@ #include #include +#include namespace gridtools::fn { namespace { @@ -26,10 +27,6 @@ namespace gridtools::fn { using edge_dim_t = integral_constant; using edge_to_cell_dim_t = integral_constant; - constexpr std::size_t num_elements = 3; - constexpr std::size_t num_neighbors = 2; - __device__ std::int32_t contents[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; - template __device__ auto neighbor_table_neighbors_device(const Table &table, std::size_t index) -> array { @@ -37,6 +34,17 @@ namespace gridtools::fn { } TEST(sid_neighbor_table, correctness_cuda) { + constexpr std::size_t num_elements = 3; + constexpr std::size_t num_neighbors = 2; + + const std::int32_t data[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; + const auto device_data = cuda_util::cuda_malloc(num_elements * num_neighbors); + GT_CUDA_CHECK(cudaMemcpy(device_data.get(), &data, sizeof data, cudaMemcpyHostToDevice)); + using dim_hymap_t = hymap::keys; + auto contents = sid::synthetic() + .set(sid::host_device::simple_ptr_holder(device_data.get())) + .set(dim_hymap_t::make_values(num_neighbors, 1)); + const auto table = as_neighbor_table(contents); using table_t = std::decay_t; From 9afdaae317f8012eb630a00faf305ad8bd096cec Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Thu, 25 Aug 2022 14:06:08 +0200 Subject: [PATCH 17/27] to trailing return type --- include/gridtools/fn/sid_neighbor_table.hpp | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index a04b2c84d..47ebf886f 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -46,7 +46,12 @@ namespace gridtools::fn::sid_neighbor_table { }; template - auto as_neighbor_table(Sid &&sid) { + auto as_neighbor_table(Sid &&sid) -> sid_neighbor_table, + sid::strides_type> { + static_assert(gridtools::tuple_util::size()))>::value == 2, "Neighbor tables must have exactly two dimensions: the index dimension and the neighbor dimension"); static_assert(!std::is_same_v, @@ -55,11 +60,7 @@ namespace gridtools::fn::sid_neighbor_table { const auto origin = sid::get_origin(sid); const auto strides = sid::get_strides(sid); - return sid_neighbor_table, - sid::strides_type>{origin, strides}; + return {origin, strides}; } } // namespace sid_neighbor_table_impl_ From f47a4eceec7284e2ded3e8f54a221bfc32f35d3c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Thu, 25 Aug 2022 14:48:39 +0200 Subject: [PATCH 18/27] using neighbor_table::neighbors wrapper fun --- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 6 +++--- tests/unit_tests/fn/test_fn_sid_neighbor_table.cu | 2 +- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index 684ed9024..990f7c032 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -30,9 +30,9 @@ namespace gridtools::fn { std::int32_t contents[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; const auto table = as_neighbor_table(contents); - auto [n00, n01] = neighbor_table_neighbors(table, 0); - auto [n10, n11] = neighbor_table_neighbors(table, 1); - auto [n20, n21] = neighbor_table_neighbors(table, 2); + auto [n00, n01] = neighbor_table::neighbors(table, 0); + auto [n10, n11] = neighbor_table::neighbors(table, 1); + auto [n20, n21] = neighbor_table::neighbors(table, 2); EXPECT_EQ(n00, 0); EXPECT_EQ(n01, 1); EXPECT_EQ(n10, 10); diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu index bf3434fb1..12ba9d12e 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu @@ -30,7 +30,7 @@ namespace gridtools::fn { template __device__ auto neighbor_table_neighbors_device(const Table &table, std::size_t index) -> array { - return neighbor_table_neighbors(table, index); + return neighbor_table::neighbors(table, index); } TEST(sid_neighbor_table, correctness_cuda) { From af631682e6947d239d6947b15f924e26d7774de3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Fri, 26 Aug 2022 11:16:50 +0200 Subject: [PATCH 19/27] removed superfluous template params --- include/gridtools/fn/sid_neighbor_table.hpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index 47ebf886f..8528df39f 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -27,10 +27,7 @@ namespace gridtools::fn::sid_neighbor_table { PtrHolder origin; Strides strides; - GT_FUNCTION friend auto neighbor_table_neighbors( - sid_neighbor_table const &table, - std::size_t index) { - + GT_FUNCTION friend auto neighbor_table_neighbors(sid_neighbor_table const &table, std::size_t index) { auto ptr = table.origin(); using element_type = std::remove_reference_t; From fbde9423c167c194df910022564580ad1b007c0c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Fri, 26 Aug 2022 11:17:42 +0200 Subject: [PATCH 20/27] consistent const placement --- tests/unit_tests/fn/test_fn_sid_neighbor_table.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu index 12ba9d12e..5bade402b 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu @@ -28,7 +28,7 @@ namespace gridtools::fn { using edge_to_cell_dim_t = integral_constant; template - __device__ auto neighbor_table_neighbors_device(const Table &table, std::size_t index) + __device__ auto neighbor_table_neighbors_device(Table const &table, std::size_t index) -> array { return neighbor_table::neighbors(table, index); } From 63eb8329343936e65db0e68ec590cce416bc8b51 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Fri, 26 Aug 2022 12:55:42 +0200 Subject: [PATCH 21/27] i don't know, just trying to get it compile... --- include/gridtools/fn/sid_neighbor_table.hpp | 2 +- tests/unit_tests/fn/test_fn_sid_neighbor_table.cu | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index 8528df39f..b91f367c8 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -27,7 +27,7 @@ namespace gridtools::fn::sid_neighbor_table { PtrHolder origin; Strides strides; - GT_FUNCTION friend auto neighbor_table_neighbors(sid_neighbor_table const &table, std::size_t index) { + GT_FUNCTION friend auto neighbor_table_neighbors(sid_neighbor_table const &table, int index) { auto ptr = table.origin(); using element_type = std::remove_reference_t; diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu index 5bade402b..bdacb046c 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu @@ -28,7 +28,7 @@ namespace gridtools::fn { using edge_to_cell_dim_t = integral_constant; template - __device__ auto neighbor_table_neighbors_device(Table const &table, std::size_t index) + __device__ auto neighbor_table_neighbors_device(Table const &table, int index) -> array { return neighbor_table::neighbors(table, index); } From 37742bdff801646ae90f1541fd0f284d0edcdc02 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Fri, 26 Aug 2022 15:07:11 +0200 Subject: [PATCH 22/27] ints --- include/gridtools/fn/sid_neighbor_table.hpp | 2 +- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 2 +- tests/unit_tests/fn/test_fn_sid_neighbor_table.cu | 6 +++--- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index b91f367c8..ed16bd1e9 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -42,7 +42,7 @@ namespace gridtools::fn::sid_neighbor_table { } }; - template + template auto as_neighbor_table(Sid &&sid) -> sid_neighbor_table(contents); auto [n00, n01] = neighbor_table::neighbors(table, 0); diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu index bdacb046c..14c393496 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu @@ -29,7 +29,7 @@ namespace gridtools::fn { template __device__ auto neighbor_table_neighbors_device(Table const &table, int index) - -> array { + -> array { return neighbor_table::neighbors(table, index); } @@ -37,8 +37,8 @@ namespace gridtools::fn { constexpr std::size_t num_elements = 3; constexpr std::size_t num_neighbors = 2; - const std::int32_t data[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; - const auto device_data = cuda_util::cuda_malloc(num_elements * num_neighbors); + const int data[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; + const auto device_data = cuda_util::cuda_malloc(num_elements * num_neighbors); GT_CUDA_CHECK(cudaMemcpy(device_data.get(), &data, sizeof data, cudaMemcpyHostToDevice)); using dim_hymap_t = hymap::keys; auto contents = sid::synthetic() From 33bbd766ada9f2f1978a481f5fcbf7a13b0dd3e5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Fri, 26 Aug 2022 17:16:02 +0200 Subject: [PATCH 23/27] removed friend function to fix old gcc builds --- include/gridtools/fn/sid_neighbor_table.hpp | 30 +++++++++++++-------- 1 file changed, 19 insertions(+), 11 deletions(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index ed16bd1e9..e263bda50 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -26,21 +26,29 @@ namespace gridtools::fn::sid_neighbor_table { struct sid_neighbor_table { PtrHolder origin; Strides strides; + }; - GT_FUNCTION friend auto neighbor_table_neighbors(sid_neighbor_table const &table, int index) { - auto ptr = table.origin(); - using element_type = std::remove_reference_t; + template + GT_FUNCTION auto neighbor_table_neighbors( + sid_neighbor_table const &table, + int index) { + + auto ptr = table.origin(); + using element_type = std::remove_reference_t; - gridtools::array neighbors; + gridtools::array neighbors; - sid::shift(ptr, sid::get_stride(table.strides), index); - for (std::size_t element_idx = 0; element_idx < MaxNumNeighbors; ++element_idx) { - neighbors[element_idx] = *ptr; - sid::shift(ptr, sid::get_stride(table.strides), 1); - } - return neighbors; + sid::shift(ptr, sid::get_stride(table.strides), index); + for (std::size_t element_idx = 0; element_idx < MaxNumNeighbors; ++element_idx) { + neighbors[element_idx] = *ptr; + sid::shift(ptr, sid::get_stride(table.strides), 1); } - }; + return neighbors; + } template auto as_neighbor_table(Sid &&sid) -> sid_neighbor_table Date: Mon, 29 Aug 2022 12:14:54 +0200 Subject: [PATCH 24/27] use integral constant expressions --- include/gridtools/fn/sid_neighbor_table.hpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index e263bda50..dafabce4a 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -37,6 +37,8 @@ namespace gridtools::fn::sid_neighbor_table { sid_neighbor_table const &table, int index) { + using namespace gridtools::literals; + auto ptr = table.origin(); using element_type = std::remove_reference_t; @@ -45,7 +47,7 @@ namespace gridtools::fn::sid_neighbor_table { sid::shift(ptr, sid::get_stride(table.strides), index); for (std::size_t element_idx = 0; element_idx < MaxNumNeighbors; ++element_idx) { neighbors[element_idx] = *ptr; - sid::shift(ptr, sid::get_stride(table.strides), 1); + sid::shift(ptr, sid::get_stride(table.strides), 1_c); } return neighbors; } From 95f088ac4e34ee3918fbbac66fbc3e4d2c297a87 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Tue, 30 Aug 2022 15:09:58 +0200 Subject: [PATCH 25/27] use decay_t to support const ptr holders --- include/gridtools/fn/sid_neighbor_table.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp index dafabce4a..f61733bf3 100644 --- a/include/gridtools/fn/sid_neighbor_table.hpp +++ b/include/gridtools/fn/sid_neighbor_table.hpp @@ -40,7 +40,7 @@ namespace gridtools::fn::sid_neighbor_table { using namespace gridtools::literals; auto ptr = table.origin(); - using element_type = std::remove_reference_t; + using element_type = std::decay_t; gridtools::array neighbors; From 4497a48b5ae25fa61c449cf4c72c6f8c57bfd1bb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?P=C3=A9ter=20Kardos?= Date: Tue, 30 Aug 2022 15:11:16 +0200 Subject: [PATCH 26/27] use const array in tests --- tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp index e5bc44e3a..2f43ab7ee 100644 --- a/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp +++ b/tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp @@ -27,7 +27,7 @@ namespace gridtools::fn { TEST(sid_neighbor_table, correctness) { constexpr std::size_t num_elements = 3; constexpr std::size_t num_neighbors = 2; - int contents[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; + const int contents[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}}; const auto table = as_neighbor_table(contents); auto [n00, n01] = neighbor_table::neighbors(table, 0); From 907d43cb8784dfd067a43bb7ad17f5e3d23d28c2 Mon Sep 17 00:00:00 2001 From: Hannes Vogt Date: Tue, 4 Oct 2022 11:51:05 +0200 Subject: [PATCH 27/27] Update AUTHORS --- AUTHORS | 1 + 1 file changed, 1 insertion(+) diff --git a/AUTHORS b/AUTHORS index 20ca08997..45790d056 100644 --- a/AUTHORS +++ b/AUTHORS @@ -15,3 +15,4 @@ Willem Deconinck (wdeconinck), ECMWF Auriane Reverdell (aurianer), ETH Zurich (CSCS) Mikael Simberg (msimberg), ETH Zurich (CSCS) Till Ehrengruber (tehrengruber), ETH Zurich (CSCS) +Péter Kardos (petiaccja), ETH Zurich (EXCLAIM)