Repository navigation
UCT/GDA: Use gpunetio submodule #11122
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
Merged
Merged
Changes from 10 commits
Commits
Show all changes
19 commits
Select commit
Hold shift + click to select a range
cd1a43b
UCT/GDA: Use gpunetio submodule
Artemy-Mellanox 5dfed4b
UCT/GDA: Use gpunetio submodule - 2
Artemy-Mellanox bf2100f
UCT/GDA: Use gpunetio submodule - 3
Artemy-Mellanox 682c1e8
UCT/GDA: Use gpunetio submodule - 4
Artemy-Mellanox 432a4ba
UCT/GDA: Use gpunetio submodule - 5
Artemy-Mellanox a3204c7
Merge remote-tracking branch 'origin/master' into topic/gpunetio-use
Artemy-Mellanox b5b2f04
UCT/GDA: Use gpunetio submodule - 6
Artemy-Mellanox e79b87a
UCT/GDA: Use gpunetio submodule - 7
Artemy-Mellanox 9f3b8c4
UCT/GDA: Use gpunetio submodule - 8
Artemy-Mellanox d0e5ea0
UCT/GDA: Use gpunetio submodule - 9
Artemy-Mellanox 1133624
UCT/GDA: Use gpunetio submodule - 10
Artemy-Mellanox 82dbc5e
UCT/GDA: Use gpunetio submodule - 11
Artemy-Mellanox d4b2655
UCT/GDA: Use gpunetio submodule - 12
Artemy-Mellanox 94c18a9
UCT/GDA: Use gpunetio submodule - 13
Artemy-Mellanox 3fe30bc
UCT/GDA: Use gpunetio submodule - 14
Artemy-Mellanox 1dd18a1
UCT/GDA: Use gpunetio submodule - 15
Artemy-Mellanox b8b1885
UCT/GDA: Use gpunetio submodule - 16
Artemy-Mellanox 22762c9
UCT/GDA: Use gpunetio submodule - 17
Artemy-Mellanox af10456
Merge remote-tracking branch 'origin/master' into topic/gpunetio-use
Artemy-Mellanox File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
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
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
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
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
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1 +1,2 @@ | ||
| +usr/lib/ucx/libuct_ib_mlx5_gda.* | ||
|
yosefe marked this conversation as resolved.
Resolved
|
||
| usr/lib/ucx/libuct_ib_mlx5_gda.* | ||
| usr/include/uct/ib/mlx5/gdaki/* | ||
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
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
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
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
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1 @@ | ||
| gpunetio/include/common | ||
|
yosefe marked this conversation as resolved.
Resolved
|
||
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
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1 @@ | ||
| gpunetio/include/device |
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -9,13 +9,20 @@ | |
|
|
||
| #include "gdaki_dev.h" | ||
|
|
||
| #include <doca_gpunetio_dev_verbs_qp.cuh> | ||
| #include <infiniband/mlx5dv.h> /* TODO add to gpunetio */ | ||
|
tvegas1 marked this conversation as resolved.
Resolved
|
||
| #include <cuda.h> /* TODO add to gpunetio */ | ||
| #include <uct/ib/mlx5/gdaki/device/doca_gpunetio_dev_verbs_qp.cuh> | ||
|
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. can we include "device/doca_gpunetio_dev_verbs_qp.cuh" ? |
||
| #include <cooperative_groups.h> | ||
|
|
||
| #define UCT_RC_GDA_RESV_WQE_NO_RESOURCE -1ULL | ||
| #define UCT_RC_GDA_WQE_ERR UCS_BIT(63) | ||
| #define UCT_RC_GDA_WQE_MASK UCS_MASK(63) | ||
|
|
||
| UCS_F_DEVICE uint32_t uct_rc_mlx5_gda_wqe_idx_inc_mask(uint32_t wqe_idx, | ||
|
yosefe marked this conversation as resolved.
Resolved
|
||
| uint32_t increment) | ||
| { | ||
| return (wqe_idx + increment) & DOCA_GPUNETIO_VERBS_WQE_PI_MASK; | ||
| } | ||
|
|
||
| UCS_F_DEVICE uct_rc_gdaki_dev_qp_t * | ||
| uct_rc_mlx5_gda_get_qp(uct_rc_gdaki_dev_ep_t *ep, unsigned cid) | ||
|
|
@@ -30,7 +37,7 @@ UCS_F_DEVICE void *uct_rc_mlx5_gda_get_wqe_ptr(uct_rc_gdaki_dev_ep_t *ep, | |
| const uintptr_t wqe_addr = __ldg((uintptr_t*)&ep->sq_wqe_daddr); | ||
| const uint32_t idx = wqe_idx & (wqe_num - 1); | ||
| const uint32_t full_idx = idx + cid * wqe_num; | ||
| return (void*)(wqe_addr + (full_idx << DOCA_GPUNETIO_MLX5_WQE_SQ_SHIFT)); | ||
| return (void*)(wqe_addr + (full_idx << DOCA_GPUNETIO_IB_MLX5_WQE_SQ_SHIFT)); | ||
| } | ||
|
|
||
| UCS_F_DEVICE void uct_rc_mlx5_gda_ring_db(uct_rc_gdaki_dev_ep_t *ep, | ||
|
|
@@ -320,7 +327,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_single( | |
| if (lane_id == 0) { | ||
| uint16_t wqe_idx = (uint16_t)wqe_base; | ||
| if ((comp != nullptr) || uct_rc_mlx5_gda_fc(ep, wqe_idx)) { | ||
| cflag = DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE; | ||
| cflag = DOCA_GPUNETIO_IB_MLX5_WQE_CTRL_CQ_UPDATE; | ||
| if (comp != nullptr) { | ||
| comp->wqe_idx = wqe_base; | ||
| comp->channel_id = cid; | ||
|
|
@@ -446,7 +453,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi( | |
| return UCS_ERR_NO_RESOURCE; | ||
| } | ||
|
|
||
| wqe_idx = doca_gpu_dev_verbs_wqe_idx_inc_mask(wqe_base, lane_id); | ||
| wqe_idx = uct_rc_mlx5_gda_wqe_idx_inc_mask(wqe_base, lane_id); | ||
| for (uint32_t i = lane_id; i < count; i += num_lanes) { | ||
| if (i == counter_index) { | ||
| atomic = true; | ||
|
|
@@ -468,7 +475,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi( | |
| cflag = 0; | ||
| if (((comp != nullptr) && (i == count - 1)) || | ||
| ((comp == nullptr) && uct_rc_mlx5_gda_fc(ep, wqe_idx))) { | ||
| cflag = DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE; | ||
| cflag = DOCA_GPUNETIO_IB_MLX5_WQE_CTRL_CQ_UPDATE; | ||
| if (comp != nullptr) { | ||
| comp->wqe_idx = wqe_base; | ||
| comp->channel_id = cid; | ||
|
|
@@ -482,7 +489,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi( | |
| ep, wqe_ptr, wqe_idx, opcode, cflag, remote_address, rkey, | ||
| reinterpret_cast<uint64_t>(address), lkey, length, atomic, | ||
| counter_inc_value, cid); | ||
| wqe_idx = doca_gpu_dev_verbs_wqe_idx_inc_mask(wqe_idx, num_lanes); | ||
| wqe_idx = uct_rc_mlx5_gda_wqe_idx_inc_mask(wqe_idx, num_lanes); | ||
| } | ||
|
|
||
| uct_rc_mlx5_gda_sync<level>(); | ||
|
|
@@ -540,7 +547,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi_partial( | |
| return UCS_ERR_NO_RESOURCE; | ||
| } | ||
|
|
||
| wqe_idx = doca_gpu_dev_verbs_wqe_idx_inc_mask(wqe_base, lane_id); | ||
| wqe_idx = uct_rc_mlx5_gda_wqe_idx_inc_mask(wqe_base, lane_id); | ||
| for (uint32_t i = lane_id; i < count; i += num_lanes) { | ||
| if (i == mem_list_count) { | ||
| idx = counter_index; | ||
|
|
@@ -564,7 +571,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi_partial( | |
| cflag = 0; | ||
| if (((comp != nullptr) && (i == count - 1)) || | ||
| ((comp == nullptr) && uct_rc_mlx5_gda_fc(ep, wqe_idx))) { | ||
| cflag = DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE; | ||
| cflag = DOCA_GPUNETIO_IB_MLX5_WQE_CTRL_CQ_UPDATE; | ||
| if (comp != nullptr) { | ||
| comp->wqe_idx = wqe_base; | ||
| comp->channel_id = cid; | ||
|
|
@@ -578,7 +585,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi_partial( | |
| ep, wqe_ptr, wqe_idx, opcode, cflag, remote_address, rkey, | ||
| reinterpret_cast<uint64_t>(address), lkey, length, atomic, | ||
| counter_inc_value, cid); | ||
| wqe_idx = doca_gpu_dev_verbs_wqe_idx_inc_mask(wqe_idx, num_lanes); | ||
| wqe_idx = uct_rc_mlx5_gda_wqe_idx_inc_mask(wqe_idx, num_lanes); | ||
| } | ||
|
|
||
| uct_rc_mlx5_gda_sync<level>(); | ||
|
|
||
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
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
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
Oops, something went wrong.
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.
Uh oh!
There was an error while loading. Please reload this page.