GDP plugin support patches #11626
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
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,81 @@ | ||
| /** | ||
| * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2026. ALL RIGHTS RESERVED. | ||
| * | ||
| * See file LICENSE for terms. | ||
| */ | ||
|
|
||
| #ifndef UCT_DEV_COMMON_CUH_ | ||
| #define UCT_DEV_COMMON_CUH_ | ||
|
|
||
| #include <ucs/sys/device_code.h> | ||
|
|
||
| #include <infiniband/mlx5dv.h> /* TODO add to gpunetio */ | ||
| #include <cuda.h> /* TODO add to gpunetio */ | ||
| #include "gpunetio/device/doca_gpunetio_dev_verbs_common.cuh" | ||
| #include <cooperative_groups.h> | ||
|
|
||
|
|
||
| template<ucs_device_level_t level> | ||
| UCS_F_DEVICE void uct_dev_exec_init(unsigned &lane_id, unsigned &num_lanes) | ||
|
rakhmets marked this conversation as resolved.
|
||
| { | ||
| switch (level) { | ||
| case UCS_DEVICE_LEVEL_THREAD: | ||
| lane_id = 0; | ||
| num_lanes = 1; | ||
| break; | ||
| case UCS_DEVICE_LEVEL_WARP: | ||
| lane_id = doca_gpu_dev_verbs_get_lane_id(); | ||
| num_lanes = UCS_DEVICE_NUM_THREADS_IN_WARP; | ||
| break; | ||
| case UCS_DEVICE_LEVEL_BLOCK: | ||
| lane_id = threadIdx.x; | ||
| num_lanes = blockDim.x; | ||
| break; | ||
| case UCS_DEVICE_LEVEL_GRID: | ||
| lane_id = threadIdx.x + blockIdx.x * blockDim.x; | ||
| num_lanes = blockDim.x * gridDim.x; | ||
| break; | ||
| } | ||
| } | ||
|
|
||
| template<ucs_device_level_t level> UCS_F_DEVICE void uct_dev_sync(void) | ||
| { | ||
| switch (level) { | ||
| case UCS_DEVICE_LEVEL_WARP: | ||
| __syncwarp(); | ||
| break; | ||
| case UCS_DEVICE_LEVEL_BLOCK: | ||
| __syncthreads(); | ||
| break; | ||
| case UCS_DEVICE_LEVEL_THREAD: | ||
| break; | ||
| case UCS_DEVICE_LEVEL_GRID: | ||
| auto g = cooperative_groups::this_grid(); | ||
| g.sync(); | ||
| } | ||
| } | ||
|
|
||
| template<ucs_device_level_t level> | ||
| UCS_F_DEVICE int uct_dev_bcast(int value, unsigned lane_id) | ||
| { | ||
| switch (level) { | ||
| case UCS_DEVICE_LEVEL_WARP: | ||
| return __shfl_sync(0xffffffff, value, 0); | ||
| case UCS_DEVICE_LEVEL_BLOCK: { | ||
| __shared__ int shared_value; | ||
| if (lane_id == 0) { | ||
| shared_value = value; | ||
| } | ||
| __syncthreads(); | ||
| value = shared_value; | ||
| __syncthreads(); | ||
| return value; | ||
| } | ||
| case UCS_DEVICE_LEVEL_THREAD: | ||
| case UCS_DEVICE_LEVEL_GRID: | ||
| default: | ||
| return value; | ||
| } | ||
| } | ||
|
|
||
| #endif /* UCT_DEV_COMMON_CUH_ */ | ||
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,116 @@ | ||
| /** | ||
| * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2026. ALL RIGHTS RESERVED. | ||
| * | ||
| * See file LICENSE for terms. | ||
| */ | ||
|
|
||
| #ifndef UCT_D2P_CUH_ | ||
| #define UCT_D2P_CUH_ | ||
|
|
||
| #include "d2p.h" | ||
| #include "d2p_proto.h" | ||
| #include "common.cuh" | ||
|
|
||
| #include <uct/api/device/uct_device_types.h> | ||
| #include <ucs/sys/device_code.h> | ||
|
|
||
|
|
||
| template<ucs_device_level_t level> | ||
| UCS_F_DEVICE ucs_status_t uct_ib_d2p_post_desc(uct_ib_d2p_gpu_ep_t *ep, | ||
| uint8_t opcode, uint32_t length, | ||
| uint32_t lkey, uint64_t laddr, | ||
| uint32_t rkey, uint64_t raddr, | ||
| uint64_t add, uint16_t flags) | ||
| { | ||
| ucs_status_t status = UCS_INPROGRESS; | ||
| unsigned lane_id, num_lanes; | ||
|
|
||
| uct_dev_exec_init<level>(lane_id, num_lanes); | ||
|
|
||
| if (lane_id == 0) { | ||
| const long long depth = UCS_BIT(ep->log_depth); | ||
| unsigned long long pi = READ_ONCE(*ep->pi); | ||
|
|
||
| for (;;) { | ||
| unsigned long long ci = READ_ONCE(*ep->ci); | ||
| if (static_cast<long long>(pi - ci) >= depth) { | ||
| status = UCS_ERR_NO_RESOURCE; | ||
| break; | ||
| } | ||
|
|
||
| unsigned long long prev = atomicCAS(ep->pi, pi, pi + 1); | ||
| if (prev == pi) { | ||
| break; | ||
| } | ||
| pi = prev; | ||
| } | ||
|
|
||
| if (status == UCS_INPROGRESS) { | ||
| const uint32_t slot = pi & UCS_MASK(ep->log_depth); | ||
| auto desc = reinterpret_cast<volatile uct_ib_d2p_desc_t*>( | ||
| ep->queue_base) + | ||
| slot; | ||
|
|
||
| desc->opcode = opcode; | ||
| desc->length = length; | ||
| desc->qp_idx = ep->qp_idx; | ||
| desc->lkey = lkey; | ||
| desc->laddr = laddr; | ||
| desc->rkey = rkey; | ||
| desc->raddr = raddr; | ||
| desc->add = add; | ||
| desc->flags = flags; | ||
| desc->owner = (pi >> ep->log_depth) & 0x1; | ||
|
Contributor
There was a problem hiding this comment. Is a __threadfence(); missing here? |
||
| } | ||
| } | ||
|
|
||
| return static_cast<ucs_status_t>( | ||
| uct_dev_bcast<level>(static_cast<int>(status), lane_id)); | ||
| } | ||
|
|
||
| template<ucs_device_level_t level> | ||
| UCS_F_DEVICE ucs_status_t uct_ib_d2p_ep_put( | ||
| uct_device_ep_h tl_ep, const uct_device_mem_elem_t *src_uct_elem, | ||
| const uct_device_mem_elem_t *tl_mem_elem, const void *address, | ||
| uint64_t remote_address, size_t length, uint64_t flags, | ||
| uct_device_completion_t *comp) | ||
| { | ||
| auto ep = reinterpret_cast<uct_ib_d2p_gpu_ep_t*>(tl_ep); | ||
| auto src_ib = reinterpret_cast<const uct_ib_md_device_mem_element_t*>( | ||
| src_uct_elem); | ||
| auto rem_ib = reinterpret_cast<const uct_ib_md_device_mem_element_t*>( | ||
| tl_mem_elem); | ||
|
|
||
| return uct_ib_d2p_post_desc<level>( | ||
| ep, UCT_IB_D2P_OP_RDMA_WRITE, length, src_ib->lkey, | ||
| reinterpret_cast<uint64_t>(address), rem_ib->rkey, remote_address, | ||
| 0, comp == nullptr ? 0 : UCT_IB_D2P_FLAG_CQ_UPDATE); | ||
| } | ||
|
|
||
| template<ucs_device_level_t level> | ||
| UCS_F_DEVICE ucs_status_t uct_ib_d2p_ep_atomic_add( | ||
| uct_device_ep_h tl_ep, const uct_device_mem_elem_t *tl_mem_elem, | ||
| uint64_t inc_value, uint64_t remote_address, uint64_t flags, | ||
| uct_device_completion_t *comp) | ||
| { | ||
| auto ep = reinterpret_cast<uct_ib_d2p_gpu_ep_t*>(tl_ep); | ||
| auto rem_ib = reinterpret_cast<const uct_ib_md_device_mem_element_t*>( | ||
| tl_mem_elem); | ||
|
|
||
| return uct_ib_d2p_post_desc<level>(ep, UCT_IB_D2P_OP_ATOMIC_ADD, | ||
| sizeof(uint64_t), ep->atomic_result_lkey, | ||
| ep->atomic_result_va, rem_ib->rkey, | ||
| remote_address, inc_value, | ||
| comp == nullptr ? | ||
| 0 : | ||
| UCT_IB_D2P_FLAG_CQ_UPDATE); | ||
| } | ||
|
|
||
| template<ucs_device_level_t level> | ||
| UCS_F_DEVICE ucs_status_t uct_ib_d2p_ep_check_completion( | ||
| uct_device_ep_h tl_ep, uct_device_completion_t *comp) | ||
| { | ||
| return UCS_OK; | ||
|
Contributor
There was a problem hiding this comment. That means it's always completed? |
||
| } | ||
|
|
||
| #endif /* UCT_D2P_CUH_ */ | ||
Oops, something went wrong.
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.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
unused?