Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Respect ReadOnly Property in Nanobind Adapter #1809

Merged
merged 94 commits into from
Oct 30, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
94 commits
Select commit Hold shift + click to select a range
eb86cbd
Renamed BlockSizes → ThreadBlockSizes
fthaler Jun 12, 2024
f10f9ad
Add loop blocking to fn GPU backend
fthaler Jun 17, 2024
c616749
Cleanup/refactor
fthaler Jun 18, 2024
44f13fd
Treat compile-time block sizes as compile-time
fthaler Jun 18, 2024
5da7ed2
Merge remote-tracking branch 'origin/master' into loop-blocking
fthaler Jun 18, 2024
b1577c1
Merge remote-tracking branch 'origin/master' into loop-blocking
fthaler Jun 18, 2024
9a5edf1
Minor cleanup
fthaler Jun 18, 2024
90e3a50
Use uint3 instead of dim3
fthaler Jun 18, 2024
f70b885
Silent warning
fthaler Jun 18, 2024
6e77d53
Add __launch_bounds__
fthaler Jun 18, 2024
508852d
Fixed comment
fthaler Jun 18, 2024
bade455
Slightly simpler meta:: calculation
fthaler Jun 18, 2024
1641040
Maybe fix for Clang-CUDA compilation
fthaler Jun 18, 2024
85dbec8
Merge remote-tracking branch 'origin/master' into loop-blocking
fthaler Jun 18, 2024
990555e
Added another ::template
fthaler Jun 18, 2024
ea9e3c9
Make clang-14-cuda-11 happy
fthaler Jun 18, 2024
d3d5228
Fix HIP compilation
fthaler Jun 19, 2024
bba6c78
Some cleanup
fthaler Jun 19, 2024
1b547d0
Fix formatting
fthaler Jun 19, 2024
6bc7bcf
Check types of ThreadBlockSizes and LoopBlockSizes
fthaler Jun 19, 2024
cf22abe
Merge branch 'master' into loop-blocking
fthaler Jun 24, 2024
afb3201
Merge branch 'master' into loop-blocking
fthaler Jun 26, 2024
21e27ea
Faster loop blocking WIP state
fthaler Jul 22, 2024
6be319a
Add const_host_view to const data store
fthaler Jul 22, 2024
ade9e8f
Added sid::make_unrolled_loop
fthaler Jul 22, 2024
0cfff11
Use explicitly unrolled loops in GPU fn backend
fthaler Jul 22, 2024
980933c
Re-enable verification
fthaler Jul 22, 2024
64098e2
Revert some irrelevant changes
fthaler Jul 22, 2024
8a021dc
Revert irrelevant changes to neighbor tables
fthaler Jul 22, 2024
8ab05d1
Revert more irrelevant changes
fthaler Jul 22, 2024
5db3605
Revert compile-time dimensions
fthaler Jul 22, 2024
05361de
Revert irrelevant changes in unstructured backend
fthaler Jul 22, 2024
d8b7b43
Formatting
fthaler Jul 22, 2024
23d2e65
Cleanup nabla stencils
fthaler Jul 22, 2024
2977c39
Fix NVCC warning
fthaler Jul 22, 2024
e2706ad
Re-enable nabla tests
fthaler Jul 22, 2024
2628178
Use vertex_field_id where applicable
fthaler Jul 22, 2024
e43a8b6
Cleanup storage SID adaptor
fthaler Jul 22, 2024
981b30a
Revert irrelevant changes
fthaler Jul 22, 2024
b9d4d34
Increased vertical block size
fthaler Jul 22, 2024
7ea497b
Revert temporary build changes
fthaler Jul 22, 2024
6c4c4f1
Fix missing ;
fthaler Jul 22, 2024
c5b3ad9
Use reasonable unroll factors in loop unrolling tests
fthaler Jul 22, 2024
34fd7e7
Fix missing include
fthaler Jul 23, 2024
e103e15
Merge remote-tracking branch 'upstream/master' into fast-loop-blockin…
fthaler Jul 23, 2024
702fa7c
Only include cuda_type_traits.hpp when required
fthaler Jul 23, 2024
68a2bd9
Enable fast-math on GPUs
fthaler Jul 23, 2024
19af796
Fix capturing of variable
fthaler Jul 23, 2024
8474c66
Disable k-blocking for now
fthaler Jul 23, 2024
69e13e6
Fix accidental use of single precision floats in fn nabla test
fthaler Jul 23, 2024
eb9096c
Updated references
fthaler Jul 23, 2024
e0fcc95
Merge remote-tracking branch 'origin/fix-nabla-float_t' into fast-loo…
fthaler Jul 23, 2024
0996c15
Possible workaround against compiler crash
fthaler Jul 24, 2024
adce6f2
Merge remote-tracking branch 'upstream/master' into fast-loop-blockin…
fthaler Jul 24, 2024
0a29491
Revert "Possible workaround against compiler crash"
fthaler Jul 24, 2024
2e6e3e0
Re-add accidentally lost index check
fthaler Jul 24, 2024
d2a6769
Selectively enable k-blocking
fthaler Jul 24, 2024
bf0ba6c
Merge remote-tracking branch 'upstream/master' into fast-loop-blockin…
fthaler Jul 24, 2024
0b0a306
Fix use of k_blocked_backend_t
fthaler Jul 24, 2024
c591f79
Explicitly use __ldg on pointer derefs
fthaler Jul 24, 2024
7699cb5
Merge remote-tracking branch 'upstream/master' into fast-loop-blockin…
fthaler Sep 2, 2024
07cae69
Fix compilation after merge
fthaler Sep 2, 2024
2a0914f
Revert unrelated changes to master
fthaler Sep 2, 2024
82c35e2
Merge remote-tracking branch 'upstream/master' into loop-blocking
fthaler Sep 25, 2024
c3edf6f
Merge remote-tracking branch 'upstream/master' into fast-loop-blockin…
fthaler Sep 25, 2024
ef4fed7
Merge branch 'fast-loop-blocking-tmp' into loop-blocking
fthaler Sep 25, 2024
ffd070e
Use sid::make_loop in sid::make_unrolled_loop if unroll factor is 1
fthaler Sep 25, 2024
6e4bcbe
Remove leftover --use_fast_math
fthaler Sep 25, 2024
b3bd8e3
Cleanup mp_find
fthaler Sep 25, 2024
304878d
Added a few comments
fthaler Sep 25, 2024
a83027f
Address review comment
fthaler Sep 25, 2024
c896ea1
Unconditional lookup in neighbor table
fthaler Jul 24, 2024
d16c1ea
Merge remote-tracking branch 'upstream/master' into loop-blocking
fthaler Oct 7, 2024
4e115a5
Revert "Unconditional lookup in neighbor table"
fthaler Oct 7, 2024
3e39f82
Revert accidental changes to Daint envs
fthaler Oct 7, 2024
494f8de
Use default loop block sizes in fn_select
fthaler Oct 7, 2024
27fc972
Disable loop unrolling for now
fthaler Oct 22, 2024
05206ac
Updated perftest references
fthaler Oct 22, 2024
071992a
Only 8 loads for neighbors and passing tests
iomaganaris Oct 15, 2024
ac48f24
Reduces loads but errors with "operation not supported on global/shar…
iomaganaris Oct 17, 2024
d69fd86
Only read neighbor tables twice
iomaganaris Oct 17, 2024
b39c679
Fix pointing to pointer of temporary array and m_index == -1
iomaganaris Oct 17, 2024
7884d4d
Keep only the necessary changes (const table ptr)
iomaganaris Oct 17, 2024
2d2a97f
Add const in sid in sid_neighbor_table
iomaganaris Oct 22, 2024
0642864
Create as_const_sid for nanobind bindings to pass input fields as con…
iomaganaris Oct 22, 2024
3aede67
Add const in as_neighbor_table changes
iomaganaris Oct 22, 2024
0a08822
Apply suggestions from code review
iomaganaris Oct 23, 2024
45fcc27
Fix as_sid and as_const_sid variable names
iomaganaris Oct 23, 2024
f341795
Use ndarray.ReadOnly to add_const if necessary in as_sid and remove a…
iomaganaris Oct 23, 2024
d1be29b
Merge remote-tracking branch 'upstream/master' into readonly-nanobind…
fthaler Oct 29, 2024
abd61d3
Exclude changes to SID neighbor table
fthaler Oct 29, 2024
b73a91b
Added test with a read-only nanobind array
fthaler Oct 29, 2024
894e0f7
Revert unnecessary changes
fthaler Oct 29, 2024
2c4a68a
Use at least Nanobind 2.1.0
fthaler Oct 29, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 6 additions & 6 deletions include/gridtools/storage/adapter/nanobind_adapter.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -78,21 +78,21 @@ namespace gridtools {
class Strides = fully_dynamic_strides<sizeof...(Sizes)>,
class StridesKind = sid::unknown_kind>
auto as_sid(nanobind::ndarray<T, nanobind::shape<Sizes...>, Args...> ndarray,
Strides stride_spec_ = {},
Strides stride_spec = {},
StridesKind = {}) {
using sid::property;
const auto ptr = ndarray.data();
constexpr auto ndim = sizeof...(Sizes);
assert(ndim == ndarray.ndim());
gridtools::array<std::size_t, ndim> shape;
std::copy_n(ndarray.shape_ptr(), ndim, shape.begin());
gridtools::array<std::size_t, ndim> strides_;
std::copy_n(ndarray.stride_ptr(), ndim, strides_.begin());
const auto strides = select_static_strides(stride_spec_, strides_.data());
gridtools::array<std::size_t, ndim> strides;
std::copy_n(ndarray.stride_ptr(), ndim, strides.begin());
const auto static_strides = select_static_strides(stride_spec, strides.data());

return sid::synthetic()
.template set<property::origin>(sid::host_device::simple_ptr_holder<T *>{ptr})
.template set<property::strides>(strides)
.template set<property::origin>(sid::host_device::simple_ptr_holder{ptr})
.template set<property::strides>(static_strides)
.template set<property::strides_kind, StridesKind>()
.template set<property::lower_bounds>(gridtools::array<integral_constant<std::size_t, 0>, ndim>())
.template set<property::upper_bounds>(shape);
Expand Down
2 changes: 1 addition & 1 deletion tests/unit_tests/storage/adapter/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@ if (${GT_TESTS_ENABLE_PYTHON_TESTS})
FetchContent_Declare(
nanobind
GIT_REPOSITORY https://github.com/wjakob/nanobind.git
GIT_TAG v2.0.0
GIT_TAG v2.1.0
)
FetchContent_MakeAvailable(nanobind)
nanobind_build_library(nanobind-static)
Expand Down
20 changes: 20 additions & 0 deletions tests/unit_tests/storage/adapter/test_nanobind_adapter.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -42,6 +42,26 @@ TEST_F(python_init_fixture, NanobindAdapterDataDynStrides) {
EXPECT_EQ(strides[1], gridtools::get<1>(s_strides));
}

TEST_F(python_init_fixture, NanobindAdapterReadOnly) {
const auto data = reinterpret_cast<void *>(0xDEADBEEF);
constexpr int ndim = 2;
constexpr std::array<std::size_t, ndim> shape = {3, 4};
constexpr std::array<std::intptr_t, ndim> strides = {1, 3};
nb::ndarray<int, nb::shape<-1, -1>, nb::ro> ndarray{data, ndim, shape.data(), nb::handle{}, strides.data()};

const auto sid = gridtools::nanobind::as_sid(ndarray);
using element_t = gridtools::sid::element_type<decltype(sid)>;
static_assert(std::is_same_v<element_t, int const>);

const auto s_origin = sid_get_origin(sid);
const auto s_strides = sid_get_strides(sid);
const auto s_ptr = s_origin();

EXPECT_EQ(s_ptr, data);
EXPECT_EQ(strides[0], gridtools::get<0>(s_strides));
EXPECT_EQ(strides[1], gridtools::get<1>(s_strides));
}

TEST_F(python_init_fixture, NanobindAdapterStaticStridesMatch) {
const auto data = reinterpret_cast<void *>(0xDEADBEEF);
constexpr int ndim = 2;
Expand Down
Loading