Using SIDs as neighbour tables in functional API #1730

petiaccja · 2022-08-23T11:57:38Z

Adds a simple class that wraps an SID and implements the neighbour table concept. This makes it possible to use Python buffers as neighbour tables by first wrapping them into an SID, but any SID is suitable as a neighbour table.

gridtoolsjenkins · 2022-08-23T11:59:11Z

Hi there, this is jenkins continuous integration...
Do you want me to verify this patch?

havogt

Looks good. Mostly style comments.

include/gridtools/fn/sid_neighbor_table.hpp

tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp

include/gridtools/fn/sid_neighbor_table.hpp

tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp

include/gridtools/fn/sid_neighbor_table.hpp

tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp

include/gridtools/fn/sid_neighbor_table.hpp

… sid

petiaccja · 2022-08-25T09:28:00Z

launch jenkins

petiaccja · 2022-08-25T10:16:57Z

launch jenkins

petiaccja · 2022-08-25T11:00:40Z

launch perftest

include/gridtools/fn/sid_neighbor_table.hpp

tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp

tests/unit_tests/fn/test_fn_sid_neighbor_table.cu

petiaccja · 2022-08-25T13:54:34Z

launch jenkins

petiaccja · 2022-08-26T08:22:30Z

launch perftest

petiaccja · 2022-08-26T08:22:39Z

launch jenkins

petiaccja · 2022-08-26T09:17:57Z

launch jenkins

petiaccja · 2022-08-26T10:56:04Z

launch jenkins

include/gridtools/fn/sid_neighbor_table.hpp

havogt · 2022-08-26T12:25:24Z

tests/unit_tests/fn/test_fn_sid_neighbor_table.cu

+        __device__ auto neighbor_table_neighbors_device(Table const &table, int index)
+            -> array<std::int32_t, 2> {


Suggested change

__device__ auto neighbor_table_neighbors_device(Table const &table, int index)

-> array<std::int32_t, 2> {

__device__ auto neighbor_table_neighbors_device(Table const &table, int index) {

I tried, but:

[build] <PATH>/tests/unit_tests/fn/test_fn_sid_neighbor_table.cu(54): error: A non-constexpr __device__ function ("neighbor_table_neighbors_device") with "auto" deduced return type cannot be directly referenced outside the bodies of device functions, except if the reference is absent when __CUDA_ARCH__ is undefined

What's weird is that it first builds it, and complains on rebuild.

tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp

havogt · 2022-08-26T12:37:27Z

tests/unit_tests/fn/test_fn_sid_neighbor_table.cu

+            const std::int32_t data[num_elements][num_neighbors] = {{0, 1}, {10, 11}, {20, 21}};
+            const auto device_data = cuda_util::cuda_malloc<std::int32_t>(num_elements * num_neighbors);
+            GT_CUDA_CHECK(cudaMemcpy(device_data.get(), &data, sizeof data, cudaMemcpyHostToDevice));
+            using dim_hymap_t = hymap::keys<edge_dim_t, edge_to_cell_dim_t>;
+            auto contents = sid::synthetic()
+                                .set<sid::property::origin>(sid::host_device::simple_ptr_holder(device_data.get()))
+                                .set<sid::property::strides>(dim_hymap_t::make_values(num_neighbors, 1));


probably good, unless we want to introduce a dependency on storage which would simplify this.

Is there a way to use a global statically allocated array on the device? I.e. __constant__? Then we might not need this malloc, memcpy and synthetic stuff.

petiaccja · 2022-08-26T13:18:35Z

launch tests

petiaccja · 2022-08-29T10:15:32Z

launch jenkins

petiaccja · 2022-08-29T11:02:28Z

launch jenkins

include/gridtools/fn/sid_neighbor_table.hpp

havogt · 2022-08-29T11:50:37Z

include/gridtools/fn/sid_neighbor_table.hpp

+            class Strides>
+        struct sid_neighbor_table {
+            PtrHolder origin;
+            Strides strides;
+        };
+
+        template <class IndexDimension,
+            class NeighborDimension,
+            std::size_t MaxNumNeighbors,
+            class PtrHolder,
+            class Strides>
+        GT_FUNCTION auto neighbor_table_neighbors(
+            sid_neighbor_table<IndexDimension, NeighborDimension, MaxNumNeighbors, PtrHolder, Strides> const &table,
+            int index) {
+
+            using namespace gridtools::literals;
+
+            auto ptr = table.origin();
+            using element_type = std::remove_reference_t<decltype(*ptr)>;
+
+            gridtools::array<element_type, MaxNumNeighbors> neighbors;
+
+            sid::shift(ptr, sid::get_stride<IndexDimension>(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<NeighborDimension>(table.strides), 1_c);
+            }
+            return neighbors;
+        }


With explicit return type, friend works for the case that failed in CI, but we can keep the current version.

Suggested change

class Strides>

struct sid_neighbor_table {

PtrHolder origin;

Strides strides;

};

template <class IndexDimension,

class NeighborDimension,

std::size_t MaxNumNeighbors,

class PtrHolder,

class Strides>

GT_FUNCTION auto neighbor_table_neighbors(

sid_neighbor_table<IndexDimension, NeighborDimension, MaxNumNeighbors, PtrHolder, Strides> const &table,

int index) {

using namespace gridtools::literals;

auto ptr = table.origin();

using element_type = std::remove_reference_t<decltype(*ptr)>;

gridtools::array<element_type, MaxNumNeighbors> neighbors;

sid::shift(ptr, sid::get_stride<IndexDimension>(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<NeighborDimension>(table.strides), 1_c);

}

return neighbors;

}

class Strides,

class ElementType = std::remove_reference_t<decltype(*std::declval<PtrHolder>()())>>

struct sid_neighbor_table {

PtrHolder origin;

Strides strides;

friend GT_FUNCTION gridtools::array<ElementType, MaxNumNeighbors> neighbor_table_neighbors(sid_neighbor_table const &table, int index) {

auto ptr = table.origin();

gridtools::array<ElementType, MaxNumNeighbors> neighbors;

sid::shift(ptr, sid::get_stride<IndexDimension>(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<NeighborDimension>(table.strides), 1);

}

return neighbors;

}

};

I kinda got used to the friend method so it would be nice, but I think I'll just leave it and we can move it once we drop support for older GCC.

havogt · 2022-08-29T11:51:49Z

tests/unit_tests/fn/test_fn_sid_neighbor_table.cu

+        __device__ auto neighbor_table_neighbors_device(Table const &table, int index)
+            -> array<int, 2> {


Suggested change

__device__ auto neighbor_table_neighbors_device(Table const &table, int index)

-> array<int, 2> {

constexpr __device__ auto neighbor_table_neighbors_device(Table const &table, int index) {

does this work?

Well, surprisingly, yes, but I don't think it should though, because it's calling non-constexpr functions.

There is some special handling of constexpr functions in nvcc, even if not called in constexpr context it affects how nvcc handles the functions in good and in bad ways...

include/gridtools/fn/sid_neighbor_table.hpp

petiaccja · 2022-08-30T13:46:57Z

launch jenkins

petiaccja · 2022-08-30T14:47:53Z

launch perftest

havogt

Looks good.

petiaccja · 2022-10-07T07:17:35Z

launch perftest

petiaccja · 2022-10-07T07:17:42Z

launch jenkins

Adds a simple class that wraps an SID and implements the neighbour table concept. This makes it possible to use Python buffers as neighbour tables by first wrapping them into an SID, but any SID is suitable as a neighbour table.

petiaccja added 2 commits August 17, 2022 19:56

added neighbor table and a test

a68f595

proper reading of elements

9c5a26b

petiaccja requested review from fthaler and havogt August 24, 2022 08:22

havogt requested changes Aug 24, 2022

View reviewed changes

petiaccja added 7 commits August 24, 2022 13:16

includes and size_t

9f27022

more includes and size_t

d93bf1f

includes

813a41f

includes

ae473c2

test namespaces

47c8393

use sid::shift

d1f99ef

use C arrays instead of std::array

dd74069

fthaler requested changes Aug 24, 2022

View reviewed changes

include/gridtools/fn/sid_neighbor_table.hpp Outdated Show resolved Hide resolved

tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp Outdated Show resolved Hide resolved

include/gridtools/fn/sid_neighbor_table.hpp Outdated Show resolved Hide resolved

fthaler reviewed Aug 24, 2022

View reviewed changes

include/gridtools/fn/sid_neighbor_table.hpp Outdated Show resolved Hide resolved

petiaccja added 6 commits August 24, 2022 16:25

code style

fa40976

changed implementation to use ptr holder and strides instead of whole…

6fb6a86

… sid

make ADL function a friend

0bbbdc7

test on cuda

ab55a32

remove unused includes

fa522e3

cuda playing funny games

d9569af

use dynamic device memory

b67bdf8

fthaler reviewed Aug 25, 2022

View reviewed changes

include/gridtools/fn/sid_neighbor_table.hpp Outdated Show resolved Hide resolved

to trailing return type

9afdaae

fthaler reviewed Aug 25, 2022

View reviewed changes

tests/unit_tests/fn/test_fn_sid_neighbor_table.cpp Outdated Show resolved Hide resolved

fthaler reviewed Aug 25, 2022

View reviewed changes

tests/unit_tests/fn/test_fn_sid_neighbor_table.cu Outdated Show resolved Hide resolved

using neighbor_table::neighbors wrapper fun

f47a4ec

petiaccja added 2 commits August 26, 2022 11:16

removed superfluous template params

af63168

consistent const placement

fbde942

i don't know, just trying to get it compile...

63eb832

havogt requested changes Aug 26, 2022

View reviewed changes

ints

37742bd

petiaccja added 2 commits August 26, 2022 17:16

removed friend function to fix old gcc builds

33bbd76

use integral constant expressions

9e58c84

havogt requested changes Aug 29, 2022

View reviewed changes

petiaccja added 2 commits August 30, 2022 15:09

use decay_t to support const ptr holders

95f088a

use const array in tests

4497a48

havogt mentioned this pull request Sep 1, 2022

Prototype: unstructured in compiled backends GridTools/gt4py#875

Closed

tehrengruber mentioned this pull request Sep 19, 2022

Functional: Compiled backends integration unstructured GridTools/gt4py#939

Merged

havogt mentioned this pull request Sep 22, 2022

Sketch of Python neighbor table #1719

Closed

havogt approved these changes Oct 4, 2022

View reviewed changes

Update AUTHORS

907d43c

fthaler approved these changes Oct 11, 2022

View reviewed changes

havogt merged commit 27119b5 into GridTools:master Oct 11, 2022

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Using SIDs as neighbour tables in functional API #1730

Using SIDs as neighbour tables in functional API #1730

petiaccja commented Aug 23, 2022

gridtoolsjenkins commented Aug 23, 2022

havogt left a comment

petiaccja commented Aug 25, 2022

petiaccja commented Aug 25, 2022

petiaccja commented Aug 25, 2022

petiaccja commented Aug 25, 2022

petiaccja commented Aug 26, 2022

petiaccja commented Aug 26, 2022

petiaccja commented Aug 26, 2022

petiaccja commented Aug 26, 2022

havogt Aug 26, 2022

petiaccja Aug 26, 2022

havogt Aug 26, 2022

petiaccja Aug 26, 2022

petiaccja commented Aug 26, 2022

petiaccja commented Aug 29, 2022

petiaccja commented Aug 29, 2022

havogt Aug 29, 2022

petiaccja Aug 30, 2022

havogt Aug 29, 2022

havogt Aug 29, 2022

petiaccja Aug 30, 2022

havogt Aug 31, 2022

petiaccja commented Aug 30, 2022

petiaccja commented Aug 30, 2022

havogt left a comment

petiaccja commented Oct 7, 2022

petiaccja commented Oct 7, 2022

		__device__ auto neighbor_table_neighbors_device(Table const &table, int index)
		-> array<std::int32_t, 2> {

		__device__ auto neighbor_table_neighbors_device(Table const &table, int index)
		-> array<int, 2> {

	__device__ auto neighbor_table_neighbors_device(Table const &table, int index)
	-> array<int, 2> {
	constexpr __device__ auto neighbor_table_neighbors_device(Table const &table, int index) {

Using SIDs as neighbour tables in functional API #1730

Using SIDs as neighbour tables in functional API #1730

Conversation

petiaccja commented Aug 23, 2022

gridtoolsjenkins commented Aug 23, 2022

havogt left a comment

Choose a reason for hiding this comment

petiaccja commented Aug 25, 2022

petiaccja commented Aug 25, 2022

petiaccja commented Aug 25, 2022

petiaccja commented Aug 25, 2022

petiaccja commented Aug 26, 2022

petiaccja commented Aug 26, 2022

petiaccja commented Aug 26, 2022

petiaccja commented Aug 26, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

petiaccja commented Aug 26, 2022

petiaccja commented Aug 29, 2022

petiaccja commented Aug 29, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

petiaccja commented Aug 30, 2022

petiaccja commented Aug 30, 2022

havogt left a comment

Choose a reason for hiding this comment

petiaccja commented Oct 7, 2022

petiaccja commented Oct 7, 2022