From 7b663b33ef45dc415956ce299b386bc5afa5a3d0 Mon Sep 17 00:00:00 2001 From: Anton Afanasyev Date: Fri, 16 Oct 2020 14:03:30 +0000 Subject: [PATCH 01/24] naive_nabla is done --- .devcontainer/devcontainer.json | 32 ++ cpputil/unstructured/CMakeLists.txt | 20 - cpputil/unstructured/atlas_field_converter.cc | 63 --- cpputil/unstructured/atlas_field_converter.cu | 61 --- cpputil/unstructured/atlas_mesh_adapter.cc | 51 --- .../include/gridtools/next/atlas_adapter.hpp | 102 +---- .../next/atlas_array_view_adapter.hpp | 22 +- .../gridtools/next/atlas_field_util.hpp | 115 ----- .../include/gridtools/next/domain.hpp | 10 + .../include/gridtools/next/math.hpp | 0 .../include/gridtools/next/mesh.hpp | 93 ---- .../next/test_helper/field_builder.hpp | 34 -- .../next/test_helper/simple_mesh.hpp | 165 ------- .../include/gridtools/next/tmp_storage.hpp | 21 +- .../include/gridtools/next/unstructured.hpp | 23 +- cpputil/unstructured/prototype.cc | 155 ------- .../tests/regression/fvm_nabla/CMakeLists.txt | 19 +- .../tests/regression/fvm_nabla/driver.cc | 421 +++++++----------- .../tests/regression/fvm_nabla/nabla_cuda.hpp | 6 +- .../regression/fvm_nabla/nabla_naive.hpp | 343 +++++--------- .../tests/unit_tests/next/CMakeLists.txt | 4 - .../unit_tests/next/test_atlas_field_util.cc | 62 --- .../next/test_helper/CMakeLists.txt | 2 - .../next/test_helper/test_field_builder.cc | 105 ----- .../next/test_helper/test_simple_mesh.cc | 54 --- .../tests/unit_tests/next/test_mesh.cc | 30 -- 26 files changed, 370 insertions(+), 1643 deletions(-) create mode 100644 .devcontainer/devcontainer.json delete mode 100644 cpputil/unstructured/atlas_field_converter.cc delete mode 100644 cpputil/unstructured/atlas_field_converter.cu delete mode 100644 cpputil/unstructured/atlas_mesh_adapter.cc delete mode 100644 cpputil/unstructured/include/gridtools/next/atlas_field_util.hpp create mode 100644 cpputil/unstructured/include/gridtools/next/domain.hpp delete mode 100644 cpputil/unstructured/include/gridtools/next/math.hpp delete mode 100644 cpputil/unstructured/include/gridtools/next/mesh.hpp delete mode 100644 cpputil/unstructured/include/gridtools/next/test_helper/field_builder.hpp delete mode 100644 cpputil/unstructured/include/gridtools/next/test_helper/simple_mesh.hpp delete mode 100644 cpputil/unstructured/prototype.cc delete mode 100644 cpputil/unstructured/tests/unit_tests/next/test_atlas_field_util.cc delete mode 100644 cpputil/unstructured/tests/unit_tests/next/test_helper/CMakeLists.txt delete mode 100644 cpputil/unstructured/tests/unit_tests/next/test_helper/test_field_builder.cc delete mode 100644 cpputil/unstructured/tests/unit_tests/next/test_helper/test_simple_mesh.cc delete mode 100644 cpputil/unstructured/tests/unit_tests/next/test_mesh.cc diff --git a/.devcontainer/devcontainer.json b/.devcontainer/devcontainer.json new file mode 100644 index 0000000..5453115 --- /dev/null +++ b/.devcontainer/devcontainer.json @@ -0,0 +1,32 @@ +{ + "name": "Existing Dockerfile", + + // Sets the run context to one level up instead of the .devcontainer folder. + "context": "..", + + // Update the 'dockerFile' property if you aren't using the standard 'Dockerfile' filename. + "dockerFile": "../.gitpod.Dockerfile", + + // Set *default* container specific settings.json values on container create. + "settings": { + "terminal.integrated.shell.linux": null + }, + + // Add the IDs of extensions you want installed when the container is created. + "extensions": [] + + // Use 'forwardPorts' to make a list of ports inside the container available locally. + // "forwardPorts": [], + + // Uncomment the next line to run commands after the container is created - for example installing curl. + // "postCreateCommand": "apt-get update && apt-get install -y curl", + + // Uncomment when using a ptrace-based debugger like C++, Go, and Rust + // "runArgs": [ "--cap-add=SYS_PTRACE", "--security-opt", "seccomp=unconfined" ], + + // Uncomment to use the Docker CLI from inside the container. See https://aka.ms/vscode-remote/samples/docker-from-docker. + // "mounts": [ "source=/var/run/docker.sock,target=/var/run/docker.sock,type=bind" ], + + // Uncomment to connect as a non-root user if you've added one. See https://aka.ms/vscode-remote/containers/non-root. + // "remoteUser": "vscode" +} diff --git a/cpputil/unstructured/CMakeLists.txt b/cpputil/unstructured/CMakeLists.txt index 2f6fafa..baea663 100644 --- a/cpputil/unstructured/CMakeLists.txt +++ b/cpputil/unstructured/CMakeLists.txt @@ -27,29 +27,9 @@ FetchContent_MakeAvailable(GridTools) add_library(gtnext INTERFACE) target_include_directories(gtnext INTERFACE ${CMAKE_CURRENT_SOURCE_DIR}/include) -add_executable(prototype prototype.cc) -target_link_libraries(prototype GridTools::gridtools) -target_include_directories(prototype PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include) - find_package(eckit REQUIRED) find_package(Atlas REQUIRED) -add_executable(atlas_mesh_adapter atlas_mesh_adapter.cc) -target_link_libraries(atlas_mesh_adapter atlas GridTools::gridtools) -target_include_directories(atlas_mesh_adapter PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include) - -add_executable(atlas_field_converter atlas_field_converter.cc) -target_link_libraries(atlas_field_converter atlas GridTools::gridtools) -target_include_directories(atlas_field_converter PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include) - -if(CMAKE_CUDA_COMPILER) - add_executable(atlas_field_converter_cuda atlas_field_converter.cu) - target_link_libraries(atlas_field_converter_cuda atlas GridTools::storage_gpu) - target_include_directories(atlas_field_converter_cuda PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include) - target_compile_options(atlas_field_converter_cuda PRIVATE "-std=c++17") - target_compile_options(atlas_field_converter_cuda PRIVATE "-arch=sm_50") -endif() - if(CMAKE_PROJECT_NAME STREQUAL PROJECT_NAME) include(CTest) if(BUILD_TESTING) diff --git a/cpputil/unstructured/atlas_field_converter.cc b/cpputil/unstructured/atlas_field_converter.cc deleted file mode 100644 index 5438c81..0000000 --- a/cpputil/unstructured/atlas_field_converter.cc +++ /dev/null @@ -1,63 +0,0 @@ -#include "atlas/grid.h" -#include "atlas/mesh/actions/BuildCellCentres.h" -#include "atlas/mesh/actions/BuildDualMesh.h" -#include "atlas/mesh/actions/BuildEdges.h" -#include "atlas/meshgenerator.h" -#include "gridtools/common/integral_constant.hpp" -#include -#include -#include -#include -#include -#include -#include -#include - -#include "gridtools/next/atlas_array_view_adapter.hpp" -#include -#include -#include -#include - -#include "gridtools/sid/rename_dimensions.hpp" -#include "tests/include/util/atlas_util.hpp" - -namespace dim { - struct k; -} // namespace dim - -int main() { - auto mesh = atlas_util::make_mesh(); - atlas::mesh::actions::build_edges(mesh); - - int nb_levels = 5; - atlas::functionspace::EdgeColumns fs_edges(mesh, atlas::option::levels(nb_levels) | atlas::option::halo(1)); - - atlas::Field f; - auto my_field = fs_edges.createField(atlas::option::name("my_field")); - - auto view = atlas::array::make_view(my_field); - for (int i = 0; i < fs_edges.size(); ++i) - for (int k = 0; k < nb_levels; ++k) - view(i, k) = i * 10 + k; - - auto my_field_sidified = gridtools::next::atlas_util::as::with_type{}(my_field); - static_assert(gridtools::is_sid{}); - - auto my_field_as_data_store = - gridtools::next::atlas_util::as_data_store::with_type{}(my_field); - static_assert(gridtools::is_sid{}); - - auto strides = gridtools::sid::get_strides(my_field_sidified); - auto strides_ds = gridtools::sid::get_strides(my_field_as_data_store); - for (int i = 0; i < fs_edges.size(); ++i) - for (int k = 0; k < nb_levels; ++k) { - auto ptr = gridtools::sid::get_origin(my_field_sidified)(); - gridtools::sid::shift(ptr, gridtools::at_key(strides), i); - gridtools::sid::shift(ptr, gridtools::at_key(strides), k); - auto ptr2 = gridtools::sid::get_origin(my_field_as_data_store)(); - gridtools::sid::shift(ptr2, gridtools::at_key(strides_ds), i); - gridtools::sid::shift(ptr2, gridtools::at_key(strides_ds), k); - std::cout << view(i, k) << "/" << *ptr << "/" << *ptr2 << std::endl; - } -} diff --git a/cpputil/unstructured/atlas_field_converter.cu b/cpputil/unstructured/atlas_field_converter.cu deleted file mode 100644 index f5ef8f6..0000000 --- a/cpputil/unstructured/atlas_field_converter.cu +++ /dev/null @@ -1,61 +0,0 @@ -#include "atlas/grid.h" -#include "atlas/mesh/actions/BuildCellCentres.h" -#include "atlas/mesh/actions/BuildDualMesh.h" -#include "atlas/mesh/actions/BuildEdges.h" -#include "atlas/meshgenerator.h" -#include "gridtools/common/integral_constant.hpp" -#include -#include -#include -#include -#include -#include -#include -#include - -#include "gridtools/next/atlas_array_view_adapter.hpp" -#include -#include -#include -#include - -#include "tests/include/util/atlas_util.hpp" - -namespace dim { - struct k; -} // namespace dim - -template -__global__ void kernel(Ptr ptr_holder, Strides strides, UpperBounds upper_bounds) { - auto ptr = ptr_holder(); - gridtools::sid::shift(ptr, gridtools::device::at_key(strides), threadIdx.x); - for (int i = 0; i < gridtools::device::at_key(upper_bounds); ++i) { - printf("%f\n", *ptr); - gridtools::sid::shift(ptr, gridtools::device::at_key(strides), 1); - } -} - -int main() { - auto mesh = atlas_util::make_mesh(); - atlas::mesh::actions::build_edges(mesh); - - int nb_levels = 5; - atlas::functionspace::EdgeColumns fs_edges(mesh, atlas::option::levels(nb_levels) | atlas::option::halo(1)); - - atlas::Field f; - auto my_field = fs_edges.createField(atlas::option::name("my_field")); - - auto view = atlas::array::make_view(my_field); - for (int i = 0; i < fs_edges.size(); ++i) - for (int k = 0; k < nb_levels; ++k) - view(i, k) = i * 10 + k; - - auto my_field_as_data_store = - gridtools::next::atlas_util::as_data_store::with_type{}(my_field); - static_assert(gridtools::is_sid{}); - - kernel<<<1, fs_edges.size()>>>(gridtools::sid::get_origin(my_field_as_data_store), - gridtools::sid::get_strides(my_field_as_data_store), - gridtools::sid::get_upper_bounds(my_field_as_data_store)); - cudaDeviceSynchronize(); -} diff --git a/cpputil/unstructured/atlas_mesh_adapter.cc b/cpputil/unstructured/atlas_mesh_adapter.cc deleted file mode 100644 index 9743e54..0000000 --- a/cpputil/unstructured/atlas_mesh_adapter.cc +++ /dev/null @@ -1,51 +0,0 @@ -#include "atlas/grid.h" -#include "atlas/mesh/actions/BuildCellCentres.h" -#include "atlas/mesh/actions/BuildDualMesh.h" -#include "atlas/mesh/actions/BuildEdges.h" -#include "atlas/meshgenerator.h" -#include -#include -#include - -#include -#include -#include - -#include -#include - -#include "tests/include/util/atlas_util.hpp" - -int main() { - auto mesh = atlas_util::make_mesh(); - atlas::mesh::actions::build_edges(mesh); - atlas::mesh::actions::build_node_to_edge_connectivity(mesh); - - auto const &v2e = gridtools::next::mesh::connectivity>(mesh); - - auto n_vertices = gridtools::next::connectivity::size(v2e); - std::cout << n_vertices << std::endl; - std::cout << gridtools::next::connectivity::max_neighbors(v2e) << std::endl; - std::cout << gridtools::next::connectivity::skip_value(v2e) << std::endl; - auto v2e_tbl = gridtools::next::connectivity::neighbor_table(v2e); - - static_assert(gridtools::is_sid{}); - namespace tu = gridtools::tuple_util; - auto composite = tu::make::values>(v2e_tbl); - // auto tmp = sid_get_lower_bounds(composite); - static_assert(gridtools::sid::concept_impl_::is_sid{}); - - auto strides = gridtools::sid::get_strides(v2e_tbl); - - std::cout << gridtools::at_key(strides) << std::endl; - std::cout << gridtools::at_key(strides) << std::endl; - - for (std::size_t v = 0; v < gridtools::next::connectivity::size(v2e); ++v) { - auto ptr = gridtools::sid::get_origin(v2e_tbl)(); - gridtools::sid::shift(ptr, gridtools::at_key(strides), v); - for (std::size_t i = 0; i < gridtools::next::connectivity::max_neighbors(v2e); ++i) { - gridtools::sid::shift(ptr, gridtools::at_key(strides), 1); - } - } -} diff --git a/cpputil/unstructured/include/gridtools/next/atlas_adapter.hpp b/cpputil/unstructured/include/gridtools/next/atlas_adapter.hpp index aaa10c7..2b29cf7 100644 --- a/cpputil/unstructured/include/gridtools/next/atlas_adapter.hpp +++ b/cpputil/unstructured/include/gridtools/next/atlas_adapter.hpp @@ -1,98 +1,18 @@ #pragma once -#include +#include -#include "gridtools/common/layout_map.hpp" -#include -#include -#include #include -#include -#include - -#include "mesh.hpp" -#include "unstructured.hpp" -#include - -#ifdef __CUDACC__ // TODO proper handling -#include -using storage_trait = gridtools::storage::gpu; -#else -#include -using storage_trait = gridtools::storage::cpu_ifirst; -#endif - -namespace gridtools::next::atlas_wrappers { - // not really a connectivity - template - struct primary_connectivity { - std::size_t size_; - - GT_FUNCTION friend std::size_t connectivity_size(primary_connectivity const &conn) { return conn.size_; } - }; - - template - struct regular_connectivity { - struct builder { - auto operator()(std::size_t size) { - return gridtools::storage::builder.template type().template layout<0, 1>().template dimensions( - size, std::integral_constant{}); - } +namespace atlas::mesh { + template + auto make_storage_producer(MaxNeighbors max_neighbors, Connectivity const &src) { + return [&src, max_neighbors](auto traits) { + return gridtools::storage::builder + .template type() + .dimensions(src.rows(), max_neighbors) + .initializer([&src](auto row, auto col) { return col < src.cols(row) ? src(row, col) : -1; }) + .build(); }; - - decltype(builder{}(std::size_t{})()) tbl_; - const atlas::idx_t missing_value_; // TODO Not sure if we can leave the type open - const gridtools::uint_t size_; - - regular_connectivity(atlas::mesh::IrregularConnectivity const &conn) - : tbl_{builder{}(conn.rows()).initializer([&conn](std::size_t row, std::size_t col) { - return col < static_cast(conn.cols(row)) ? conn.row(row)(col) : conn.missing_value(); - })()}, - missing_value_{conn.missing_value()}, size_{tbl_->lengths()[0]} {} - - regular_connectivity(atlas::mesh::MultiBlockConnectivity const &conn) - : tbl_{builder{}(conn.rows()).initializer([&conn](std::size_t row, std::size_t col) { - return col < static_cast(conn.cols(row)) ? conn.row(row)(col) : conn.missing_value(); - })()}, - missing_value_{conn.missing_value()}, size_{tbl_->lengths()[0]} {} - - GT_FUNCTION friend std::size_t connectivity_size(regular_connectivity const &conn) { return conn.size_; } - - GT_FUNCTION friend std::integral_constant connectivity_max_neighbors( - regular_connectivity const &) { - return {}; - } - - GT_FUNCTION friend int connectivity_skip_value(regular_connectivity const &conn) { return conn.missing_value_; } - - friend auto connectivity_neighbor_table(regular_connectivity const &conn) { - return gridtools::sid::rename_numbered_dimensions(conn.tbl_); - } - }; - -} // namespace gridtools::next::atlas_wrappers - -namespace atlas { - template