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.

Sign up for GitHub

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

Jump to bottom

Add Const to SID Neighbor Table Element Type #1808

Merged
merged 97 commits into from
Oct 30, 2024
Merged

Add Const to SID Neighbor Table Element Type #1808

merged 97 commits into from
Oct 30, 2024

Conversation

fthaler
Copy link
Contributor

@fthaler fthaler commented Oct 29, 2024 •
edited
Loading

Allows NVCC to use LDG instructions and thus perform better double-load elision. Credit goes to @iomaganaris. Further, fixes sid::as_const for C arrays.

@fthaler
Renamed BlockSizes → ThreadBlockSizes
eb86cbd
@fthaler
Add loop blocking to fn GPU backend
f10f9ad
@fthaler
Cleanup/refactor
c616749
@fthaler
Treat compile-time block sizes as compile-time
44f13fd
@fthaler
Merge remote-tracking branch 'origin/master' into loop-blocking
5da7ed2
@fthaler
Merge remote-tracking branch 'origin/master' into loop-blocking
b1577c1
@fthaler
Minor cleanup
9a5edf1
@fthaler
Use uint3 instead of dim3
90e3a50
@fthaler
Silent warning
f70b885
@fthaler
Add __launch_bounds__
6e77d53
@fthaler
Fixed comment
508852d
@fthaler
Slightly simpler meta:: calculation
bade455
@fthaler
Maybe fix for Clang-CUDA compilation
1641040
@fthaler
Merge remote-tracking branch 'origin/master' into loop-blocking
85dbec8
@fthaler
Added another ::template
990555e
@fthaler
Make clang-14-cuda-11 happy
ea9e3c9
@fthaler
Fix HIP compilation
d3d5228
@fthaler
Some cleanup
bba6c78
@fthaler
Fix formatting
1b547d0
@fthaler
Check types of ThreadBlockSizes and LoopBlockSizes
6bc7bcf
@fthaler
Merge branch 'master' into loop-blocking
cf22abe
@fthaler
Merge branch 'master' into loop-blocking
afb3201
@fthaler
Faster loop blocking WIP state
21e27ea
@fthaler
Add const_host_view to const data store
6be319a
@fthaler
Added sid::make_unrolled_loop
ade9e8f
@fthaler
Use explicitly unrolled loops in GPU fn backend
0cfff11
@fthaler
Re-enable verification
980933c
@fthaler
Revert some irrelevant changes
64098e2
@fthaler
Revert irrelevant changes to neighbor tables
8a021dc
@fthaler
Revert more irrelevant changes
8ab05d1
iomaganaris and others added 13 commits October 22, 2024 14:27
@iomaganaris
Reduces loads but errors with "operation not supported on global/shar…
ac48f24 
…ed address space" because of __ldg on array
@iomaganaris
Only read neighbor tables twice
d69fd86
@iomaganaris
Fix pointing to pointer of temporary array and m_index == -1
b39c679
@iomaganaris
Keep only the necessary changes (const table ptr)
7884d4d
@iomaganaris
Add const in sid in sid_neighbor_table
2d2a97f
@iomaganaris
Create as_const_sid for nanobind bindings to pass input fields as con…
0642864 
…st and avoid reading fields that are independent of the vertical level in each iteration
@iomaganaris
Add const in as_neighbor_table changes
3aede67
@iomaganaris @fthaler
Apply suggestions from code review
0a08822 
Improvements in `sid` handling as `const` in `as_neighbor_table`

Co-authored-by: Felix Thaler <felix.thaler@nummi.ch>
@iomaganaris
Fix as_sid and as_const_sid variable names
45fcc27
@iomaganaris
Use ndarray.ReadOnly to add_const if necessary in as_sid and remove a…
f341795 
…s_const_sid
@fthaler
Merge remote-tracking branch 'upstream/master' into const-sid-neighbo…
c1b8543 
…r-table
@fthaler
Fix minor style inconsistency
cc97916
@fthaler
Exclude nanobind adapter changes
5732d10
@gridtoolsjenkins
Copy link
Collaborator

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

@fthaler
Copy link
Contributor Author

fthaler commented Oct 29, 2024

launch perftest

@fthaler
Copy link
Contributor Author

fthaler commented Oct 29, 2024

launch jenkins

@fthaler
Copy link
Contributor Author

fthaler commented Oct 29, 2024

launch jenkins

@fthaler
Copy link
Contributor Author

fthaler commented Oct 29, 2024

launch perftest

@fthaler
Copy link
Contributor Author

fthaler commented Oct 30, 2024

launch jenkins

@fthaler
Copy link
Contributor Author

fthaler commented Oct 30, 2024

launch jenkins

@fthaler
Copy link
Contributor Author

fthaler commented Oct 30, 2024

launch perftest

iomaganaris
iomaganaris approved these changes Oct 30, 2024
View reviewed changes
Copy link
Contributor

@iomaganaris iomaganaris left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM! Thanks for taking care of this and adding a test 👍

@havogt havogt merged commit 85de9ff into GridTools:master Oct 30, 2024
70 checks passed
havogt pushed a commit that referenced this pull request Oct 30, 2024
@fthaler @iomaganaris @havogt
Add Const to SID Neighbor Table Element Type (#1808)
88af55b 
Allows NVCC to use `LDG` instructions and thus perform better
double-load elision. Credit goes to @iomaganaris. Further, fixes
`sid::as_const` for C arrays.

---------

Co-authored-by: Ioannis Magkanaris <ioannis.magkanaris@cscs.ch>
@iomaganaris iomaganaris mentioned this pull request Oct 31, 2024
Closed
3 tasks
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@iomaganaris iomaganaris iomaganaris approved these changes

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
4 participants
@fthaler @gridtoolsjenkins @iomaganaris @havogt