Closed
Conversation
msimberg
commented
Oct 27, 2025
Comment on lines
+98
to
+104
| # --------------------------------------------------------------------- | ||
| # nccl setup | ||
| # --------------------------------------------------------------------- | ||
| if(GHEX_USE_NCCL) | ||
| link_libraries("-lnccl") | ||
| # include_directories("") | ||
| endif() |
Contributor
Author
There was a problem hiding this comment.
This is a hack. Add proper FindNCCL.cmake module.
This can be tested using the icon uenv manually setting export LIBRARY_PATH=/user-environment/env/default/lib64:/user-environment/env/default/lib.
msimberg
commented
Oct 27, 2025
Comment on lines
+22
to
+47
| struct cuda_event { | ||
| cudaEvent_t m_event; | ||
| ghex::util::moved_bit m_moved; | ||
|
|
||
| cuda_event() { | ||
| GHEX_CHECK_CUDA_RESULT(cudaEventCreateWithFlags(&m_event, cudaEventDisableTiming)) | ||
| } | ||
| cuda_event(const cuda_event&) = delete; | ||
| cuda_event& operator=(const cuda_event&) = delete; | ||
| cuda_event(cuda_event&& other) = default; | ||
| cuda_event& operator=(cuda_event&&) = default; | ||
|
|
||
| ~cuda_event() | ||
| { | ||
| if (!m_moved) | ||
| { | ||
| GHEX_CHECK_CUDA_RESULT_NO_THROW(cudaEventDestroy(m_event)) | ||
| } | ||
| } | ||
|
|
||
| operator bool() const noexcept { return m_moved; } | ||
| operator cudaEvent_t() const noexcept { return m_event; } | ||
| cudaEvent_t& get() noexcept { return m_event; } | ||
| const cudaEvent_t& get() const noexcept { return m_event; } | ||
| }; | ||
|
|
msimberg
commented
Oct 27, 2025
Comment on lines
+215
to
+220
| ghex::util::moved_bit m_moved; | ||
| bool m_valid; | ||
| communicator_type m_comm; | ||
| #ifdef GHEX_USE_NCCL | ||
| ncclComm_t m_nccl_comm; | ||
| #endif |
Contributor
Author
There was a problem hiding this comment.
Move this to oomph.
msimberg
commented
Oct 27, 2025
Comment on lines
+255
to
+266
| private: | ||
| template<typename... Archs, typename... Fields> | ||
| void nccl_exchange_impl(buffer_info_type<Archs, Fields>... buffer_infos) { | ||
| pack_nccl(); | ||
|
|
||
| ncclGroupStart(); | ||
| post_sends_nccl(); | ||
| post_recvs_nccl(); | ||
| ncclGroupEnd(); | ||
|
|
||
| unpack_nccl(); | ||
| } |
Contributor
Author
There was a problem hiding this comment.
Add customization point or similar to allow doing this with NCCL?
msimberg
commented
Oct 27, 2025
| // application data | ||
| auto& d = local_domains[0]; | ||
| ghex::test::util::memory<int> field(d.size()*levels, 0); | ||
| #ifndef GHEX_USE_NCCL |
Contributor
Author
There was a problem hiding this comment.
To do: tests etc. don't like NCCL when using host memory. Copy to device when using NCCL or disallow it completely?
msimberg
commented
Oct 27, 2025
Comment on lines
297
to
301
| auto h_gpu = co.exchange(patterns(data_gpu)); | ||
| #ifdef GHEX_USE_NCCL | ||
| cudaDeviceSynchronize(); | ||
| #endif | ||
| h_gpu.wait(); |
Contributor
Author
There was a problem hiding this comment.
Introduce another API that allows skipping a blocking wait.
Contributor
Author
|
I'm keeping this branch around for use with icon4py, but closing this PR by replacing it with a cleaner version. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Opening this just for reference. This first version was just a proof of concept to see what improvement we might see.
This version hacks in NCCL support in
communication_object. Performance is very good intra-node in icon-exclaim. I haven't tested inter-node performance (which would use libfabric; the icon uenv is still using an old version of NCCL and libfabric which may not perform so well).I plan to work on updating this to a cleaner version with an oomph NCCL backend. However, this may need some new customization points etc. as the NCCL communication has to/can be set up a bit differently compared to the MPI/libfabric implementations (pack + send/recv + unpack can all be scheduled in one go).
Based on #184.