From 42f2a1ed4b8996b18dc11fd7fdc559b3d66be7cc Mon Sep 17 00:00:00 2001 From: foraxe Date: Fri, 28 Aug 2026 13:54:07 +0800 Subject: [PATCH 01/25] fix(gpunetio): make multi-chunk requests safe Signed-off-by: foraxe (cherry picked from commit 8a61ec6c00c3c41c9704abadf328fdd788b959ab) --- src/plugins/gpunetio/gpunetio_backend.cpp | 176 +++++++++++++------- src/plugins/gpunetio/gpunetio_backend.h | 6 +- src/plugins/gpunetio/gpunetio_backend_aux.h | 1 + 3 files changed, 118 insertions(+), 65 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index f8d2fe008f..da432bdd1b 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -38,6 +38,10 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) int ret; union ibv_gid rgid; + for (auto &reserved : xferReqReserved) { + reserved.store(false, std::memory_order_relaxed); + } + result = doca_log_backend_create_standard(); if (result != DOCA_SUCCESS) throw std::invalid_argument("Can't initialize doca log"); @@ -1106,20 +1110,38 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, nixlBackendReqH *&handle, const nixl_opt_b_args_t *opt_args) const { uint32_t pos; - nixlDocaBckndReq *treq = new nixlDocaBckndReq; + nixlDocaBckndReq *treq; nixlDocaPrivateMetadata *lmd; nixlDocaPublicMetadata *rmd; - uint32_t lcnt = (uint32_t)local.descCount(); - uint32_t rcnt = (uint32_t)remote.descCount(); - uint32_t stream_id; + const uint32_t lcnt = (uint32_t)local.descCount(); + const uint32_t rcnt = (uint32_t)remote.descCount(); + uint32_t stream_id = DOCA_POST_STREAM_NUM; struct nixlDocaRdmaQp *rdma_qp; uintptr_t notif_addr; + if (operation != NIXL_READ && operation != NIXL_WRITE) { + return NIXL_ERR_INVALID_PARAM; + } + if (lcnt != rcnt || lcnt == 0) { + return NIXL_ERR_INVALID_PARAM; + } + if ((lcnt + DOCA_XFER_REQ_SIZE - 1) / DOCA_XFER_REQ_SIZE > DOCA_XFER_REQ_MAX) { + return NIXL_ERR_INVALID_PARAM; + } + + for (uint32_t idx = 0; idx < lcnt; idx++) { + if (local[idx].len != remote[idx].len) { + return NIXL_ERR_INVALID_PARAM; + } + } + // TODO: check device id from local dlist mr that should be all the same and same of // the engine for (uint32_t idx = 0; idx < lcnt; idx++) { lmd = (nixlDocaPrivateMetadata *)local[idx].metadataP; - if (lmd->devId != gdevs[0].first) return NIXL_ERR_INVALID_PARAM; + if (lmd->devId != gdevs[0].first) { + return NIXL_ERR_INVALID_PARAM; + } } auto search = qpMap.find(remote_agent); @@ -1130,52 +1152,75 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, rdma_qp = search->second; - if (lcnt != rcnt) return NIXL_ERR_INVALID_PARAM; - - if (lcnt == 0) return NIXL_ERR_INVALID_PARAM; + treq = new nixlDocaBckndReq; + auto abandon_request = [&]() { + for (uint32_t reserved_pos : treq->positions) { + xferReqReserved[reserved_pos].store(false, std::memory_order_release); + } + delete treq; + }; - if (opt_args->customParam.empty()) { + if (opt_args == nullptr || opt_args->customParam.empty()) { stream_id = (xferStream.fetch_add(1) & (nstreams - 1)); treq->stream = post_stream[stream_id]; } else { treq->stream = (cudaStream_t) * ((uintptr_t *)opt_args->customParam.data()); } - treq->start_pos = (xferRingPos.fetch_add(1) & (DOCA_XFER_REQ_MAX - 1)); - pos = treq->start_pos; + auto reserve_position = [&]() -> bool { + for (uint32_t attempt = 0; attempt < DOCA_XFER_REQ_MAX; ++attempt) { + const uint32_t candidate = xferRingPos.fetch_add(1) & DOCA_XFER_REQ_MASK; + bool expected = false; + if (xferReqReserved[candidate].compare_exchange_strong( + expected, true, std::memory_order_acq_rel)) { + pos = candidate; + treq->positions.push_back(candidate); + return true; + } + } + NIXL_ERROR << "GPUNETIO transfer ring exhausted"; + return false; + }; + + treq->positions.reserve((lcnt + DOCA_XFER_REQ_SIZE - 1) / DOCA_XFER_REQ_SIZE); + if (!reserve_position()) { + abandon_request(); + return NIXL_ERR_BACKEND; + } + uint32_t desc_offset = 0; do { - for (uint32_t idx = 0; idx < lcnt && idx < DOCA_XFER_REQ_SIZE; idx++) { - size_t lsize = local[idx].len; - size_t rsize = remote[idx].len; - if (lsize != rsize) return NIXL_ERR_INVALID_PARAM; - - lmd = (nixlDocaPrivateMetadata *)local[idx].metadataP; - rmd = (nixlDocaPublicMetadata *)remote[idx].metadataP; - - xferReqRingCpu[pos].lbuf[idx] = (uintptr_t)lmd->mr->get_addr(); - xferReqRingCpu[pos].lkey[idx] = (uintptr_t)lmd->mr->get_lkey(); - xferReqRingCpu[pos].rbuf[idx] = (uintptr_t)rmd->mr->get_addr(); - xferReqRingCpu[pos].rkey[idx] = (uintptr_t)rmd->mr->get_rkey(); - xferReqRingCpu[pos].size[idx] = lsize; - xferReqRingCpu[pos].num++; + docaXferReqGpu staged_req{}; + staged_req.has_notif_msg_idx = DOCA_NOTIF_NULL; + + while (desc_offset < lcnt && staged_req.num < DOCA_XFER_REQ_SIZE) { + const uint32_t idx = staged_req.num; + const uint32_t desc_idx = desc_offset++; + + lmd = (nixlDocaPrivateMetadata *)local[desc_idx].metadataP; + rmd = (nixlDocaPublicMetadata *)remote[desc_idx].metadataP; + + staged_req.lbuf[idx] = local[desc_idx].addr; + staged_req.lkey[idx] = lmd->mr->get_lkey(); + staged_req.rbuf[idx] = remote[desc_idx].addr; + staged_req.rkey[idx] = rmd->mr->get_rkey(); + staged_req.size[idx] = local[desc_idx].len; + staged_req.num++; } - xferReqRingCpu[pos].last_rsvd = last_rsvd_flags; - xferReqRingCpu[pos].last_posted = last_posted_flags; - - xferReqRingCpu[pos].qp_data = rdma_qp->qp_data->get_qp_gpu_dev(); - xferReqRingCpu[pos].qp_notif = rdma_qp->qp_notif->get_qp_gpu_dev(); + staged_req.last_rsvd = last_rsvd_flags; + staged_req.last_posted = last_posted_flags; + staged_req.qp_data = rdma_qp->qp_data->get_qp_gpu_dev(); + staged_req.qp_notif = rdma_qp->qp_notif->get_qp_gpu_dev(); + memcpy(&xferReqRingCpu[pos], &staged_req, sizeof(staged_req)); - if (lcnt > DOCA_XFER_REQ_SIZE) { - lcnt -= DOCA_XFER_REQ_SIZE; - pos = (xferRingPos.fetch_add(1) & (DOCA_XFER_REQ_MAX - 1)); - } else { - lcnt = 0; + if (desc_offset < lcnt && !reserve_position()) { + abandon_request(); + return NIXL_ERR_BACKEND; } - } while (lcnt > 0); + } while (desc_offset < lcnt); - treq->end_pos = xferRingPos; + const uint32_t final_pos = treq->positions.back(); if (opt_args && opt_args->hasNotif) { struct nixlDocaNotif *notif; @@ -1183,6 +1228,7 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, auto search = notifMap.find(remote_agent); if (search == notifMap.end()) { NIXL_ERROR << "Can't find notif for remote_agent " << remote_agent; + abandon_request(); return NIXL_ERR_INVALID_PARAM; } @@ -1192,27 +1238,26 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, std::string newMsg = msg_tag_start + std::to_string(opt_args->notifMsg.size()) + msg_tag_end + opt_args->notifMsg; + auto &final_request = xferReqRingCpu[final_pos]; + final_request.has_notif_msg_idx = (notif->send_pi.fetch_add(1) & (notif->elems_num - 1)); notif_addr = - (uintptr_t)(notif->send_addr + - (xferReqRingCpu[treq->end_pos - 1].has_notif_msg_idx * notif->elems_size)); - xferReqRingCpu[treq->end_pos - 1].has_notif_msg_idx = - (notif->send_pi.fetch_add(1) & (notif->elems_num - 1)); - xferReqRingCpu[treq->end_pos - 1].msg_sz = newMsg.size(); - xferReqRingCpu[treq->end_pos - 1].lbuf_notif = notif_addr; - xferReqRingCpu[treq->end_pos - 1].lkey_notif = notif->send_mr->get_lkey(); + (uintptr_t)(notif->send_addr + (final_request.has_notif_msg_idx * notif->elems_size)); + final_request.msg_sz = newMsg.size(); + final_request.lbuf_notif = notif_addr; + final_request.lkey_notif = notif->send_mr->get_lkey(); memcpy((void *)notif_addr, newMsg.c_str(), newMsg.size()); NIXL_INFO << "DOCA prepXfer with notif to " << remote_agent << " at " - << xferReqRingCpu[treq->end_pos - 1].has_notif_msg_idx << " msg " << newMsg - << " to " << remote_agent; + << final_request.has_notif_msg_idx << " msg " << newMsg << " to " << remote_agent; } else { - xferReqRingCpu[treq->end_pos - 1].has_notif_msg_idx = DOCA_NOTIF_NULL; + xferReqRingCpu[final_pos].has_notif_msg_idx = DOCA_NOTIF_NULL; } - NIXL_INFO << "DOCA REQUEST from " << treq->start_pos << " to " << treq->end_pos - 1 - << " stream " << stream_id << std::endl; + NIXL_INFO << "DOCA REQUEST with " << treq->positions.size() << " ring positions, first " + << treq->positions.front() << ", last " << final_pos << ", stream " << stream_id + << std::endl; treq->backendHandleGpu = 0; @@ -1230,7 +1275,7 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, const nixl_opt_b_args_t *opt_args) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; - for (uint32_t idx = treq->start_pos; idx < treq->end_pos; idx++) { + for (uint32_t idx : treq->positions) { xferReqRingCpu[idx].id = (lastPostedReq.fetch_add(1) & (DOCA_MAX_COMPLETION_INFLIGHT_MASK)); completion_list_cpu[xferReqRingCpu[idx].id].xferReqRingGpu = xferReqRingGpu + idx; completion_list_cpu[xferReqRingCpu[idx].id].completed = 0; @@ -1255,16 +1300,17 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; uint32_t completion_index; - for (uint32_t idx = treq->start_pos; idx < treq->end_pos; idx++) { + for (uint32_t idx : treq->positions) { completion_index = xferReqRingCpu[idx].id & (DOCA_MAX_COMPLETION_INFLIGHT_MASK); - if (((volatile docaXferCompletion *)completion_list_cpu)[completion_index].completed == 1) { - *((volatile uint8_t *)&xferReqRingCpu[idx].in_use) = 0; - NIXL_INFO << "DOCA checkXfer pos " << idx << " compl_idx " << completion_index - << " COMPLETED!\n"; - return NIXL_SUCCESS; - } else + if (((volatile docaXferCompletion *)completion_list_cpu)[completion_index].completed != 1) { return NIXL_IN_PROG; + } + } + + for (uint32_t idx : treq->positions) { + *((volatile uint8_t *)&xferReqRingCpu[idx].in_use) = 0; + NIXL_INFO << "DOCA checkXfer pos " << idx << " COMPLETED!"; } return NIXL_SUCCESS; @@ -1272,11 +1318,17 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixl_status_t nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { - uint32_t tmp = xferRingPos.load() & (DOCA_XFER_REQ_MAX - 1); - if (((volatile docaXferCompletion *)completion_list_cpu)[tmp].completed > 0) - return NIXL_SUCCESS; - else - return NIXL_IN_PROG; + nixl_status_t status = checkXfer(handle); + if (status == NIXL_IN_PROG) { + return status; + } + + auto *treq = static_cast(handle); + for (uint32_t idx : treq->positions) { + xferReqReserved[idx].store(false, std::memory_order_release); + } + delete treq; + return status; } nixl_status_t diff --git a/src/plugins/gpunetio/gpunetio_backend.h b/src/plugins/gpunetio/gpunetio_backend.h index 9a17efc9f2..8cc1904b61 100644 --- a/src/plugins/gpunetio/gpunetio_backend.h +++ b/src/plugins/gpunetio/gpunetio_backend.h @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -166,6 +166,7 @@ class nixlDocaEngine : public nixlBackendEngine { struct docaXferReqGpu *xferReqRingGpu; struct docaXferReqGpu *xferReqRingCpu; mutable std::atomic xferRingPos; + mutable std::array xferReqReserved; struct docaXferCompletion *completion_list_gpu; struct docaXferCompletion *completion_list_cpu; @@ -192,8 +193,7 @@ class nixlDocaEngine : public nixlBackendEngine { public: cudaStream_t stream; uint32_t devId; - uint32_t start_pos; - uint32_t end_pos; + std::vector positions; uintptr_t backendHandleGpu; nixlDocaBckndReq() : nixlBackendReqH() {} diff --git a/src/plugins/gpunetio/gpunetio_backend_aux.h b/src/plugins/gpunetio/gpunetio_backend_aux.h index b3637663ed..0efd9cbb59 100644 --- a/src/plugins/gpunetio/gpunetio_backend_aux.h +++ b/src/plugins/gpunetio/gpunetio_backend_aux.h @@ -18,6 +18,7 @@ #ifndef GPUNETIO_BACKEND_AUX_H #define GPUNETIO_BACKEND_AUX_H +#include #include #include #include From b819b0a981c9c3d8bba795c823879bcab0f1e466 Mon Sep 17 00:00:00 2001 From: foraxe Date: Tue, 1 Sep 2026 02:11:06 +0800 Subject: [PATCH 02/25] fix(gpunetio): harden request release lifecycle Signed-off-by: foraxe (cherry picked from commit c68db923c9693886fe6756f4a9dc06f8230eec3e) --- src/core/nixl_agent.cpp | 9 +-- src/plugins/gpunetio/gpunetio_backend.cpp | 74 ++++++++++++++++------- src/plugins/gpunetio/gpunetio_backend.h | 5 +- test/gtest/unit/agent/agent.cpp | 39 ++++++++++++ 4 files changed, 99 insertions(+), 28 deletions(-) diff --git a/src/core/nixl_agent.cpp b/src/core/nixl_agent.cpp index 9f3b398000..d0c430afe6 100644 --- a/src/core/nixl_agent.cpp +++ b/src/core/nixl_agent.cpp @@ -1298,15 +1298,16 @@ nixlAgent::releaseXferReq(nixlXferReqH *req_hndl) const { if(req_hndl->status == NIXL_IN_PROG) { - req_hndl->status = req_hndl->engine->releaseReqH( - req_hndl->backendHandle); + const nixl_status_t release_status = + req_hndl->engine->releaseReqH(req_hndl->backendHandle); - if (req_hndl->status < 0) { + if (release_status < 0) { NIXL_ERROR_FUNC << "backend '" << req_hndl->engine->getType() << "' could not release transfer request and returned error status " - << req_hndl->status; + << release_status; return NIXL_ERR_REPOST_ACTIVE; // Might need renaming } + req_hndl->status = release_status; // just in case the backend doesn't set to NULL on success // this will prevent calling releaseReqH again in destructor req_hndl->backendHandle = nullptr; diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index da432bdd1b..2b430f6836 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -38,7 +38,7 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) int ret; union ibv_gid rgid; - for (auto &reserved : xferReqReserved) { + for (auto &reserved : xferReqReserved_) { reserved.store(false, std::memory_order_relaxed); } @@ -1155,7 +1155,7 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, treq = new nixlDocaBckndReq; auto abandon_request = [&]() { for (uint32_t reserved_pos : treq->positions) { - xferReqReserved[reserved_pos].store(false, std::memory_order_release); + xferReqReserved_[reserved_pos].store(false, std::memory_order_release); } delete treq; }; @@ -1164,14 +1164,18 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, stream_id = (xferStream.fetch_add(1) & (nstreams - 1)); treq->stream = post_stream[stream_id]; } else { - treq->stream = (cudaStream_t) * ((uintptr_t *)opt_args->customParam.data()); + if (opt_args->customParam.size() != sizeof(cudaStream_t)) { + abandon_request(); + return NIXL_ERR_INVALID_PARAM; + } + std::memcpy(&treq->stream, opt_args->customParam.data(), sizeof(treq->stream)); } auto reserve_position = [&]() -> bool { for (uint32_t attempt = 0; attempt < DOCA_XFER_REQ_MAX; ++attempt) { const uint32_t candidate = xferRingPos.fetch_add(1) & DOCA_XFER_REQ_MASK; bool expected = false; - if (xferReqReserved[candidate].compare_exchange_strong( + if (xferReqReserved_[candidate].compare_exchange_strong( expected, true, std::memory_order_acq_rel)) { pos = candidate; treq->positions.push_back(candidate); @@ -1237,6 +1241,10 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, // Check notifMsg size std::string newMsg = msg_tag_start + std::to_string(opt_args->notifMsg.size()) + msg_tag_end + opt_args->notifMsg; + if (newMsg.size() >= notif->elems_size) { + abandon_request(); + return NIXL_ERR_INVALID_PARAM; + } auto &final_request = xferReqRingCpu[final_pos]; final_request.has_notif_msg_idx = (notif->send_pi.fetch_add(1) & (notif->elems_num - 1)); @@ -1247,6 +1255,7 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, final_request.lkey_notif = notif->send_mr->get_lkey(); memcpy((void *)notif_addr, newMsg.c_str(), newMsg.size()); + reinterpret_cast(notif_addr)[newMsg.size()] = '\0'; NIXL_INFO << "DOCA prepXfer with notif to " << remote_agent << " at " << final_request.has_notif_msg_idx << " msg " << newMsg << " to " << remote_agent; @@ -1275,24 +1284,33 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, const nixl_opt_b_args_t *opt_args) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; - for (uint32_t idx : treq->positions) { - xferReqRingCpu[idx].id = (lastPostedReq.fetch_add(1) & (DOCA_MAX_COMPLETION_INFLIGHT_MASK)); - completion_list_cpu[xferReqRingCpu[idx].id].xferReqRingGpu = xferReqRingGpu + idx; - completion_list_cpu[xferReqRingCpu[idx].id].completed = 0; + if (operation != NIXL_READ && operation != NIXL_WRITE) { + return NIXL_ERR_INVALID_PARAM; + } - switch (operation) { - case NIXL_READ: - doca_kernel_read(treq->stream, xferReqRingCpu[idx].qp_data, xferReqRingGpu, idx); - break; - case NIXL_WRITE: + treq->postedCount = 0; + treq->postStatus = NIXL_SUCCESS; + for (uint32_t idx : treq->positions) { + std::lock_guard lock(postLock_); + const uint32_t completion_index = + lastPostedReq.load(std::memory_order_relaxed) & DOCA_MAX_COMPLETION_INFLIGHT_MASK; + xferReqRingCpu[idx].id = completion_index; + completion_list_cpu[completion_index].xferReqRingGpu = nullptr; + completion_list_cpu[completion_index].completed = 0; + + const doca_error_t result = operation == NIXL_READ ? + doca_kernel_read(treq->stream, xferReqRingCpu[idx].qp_data, xferReqRingGpu, idx) : doca_kernel_write(treq->stream, xferReqRingCpu[idx].qp_data, xferReqRingGpu, idx); + if (result != DOCA_SUCCESS) { + treq->postStatus = NIXL_ERR_BACKEND; break; - default: - return NIXL_ERR_INVALID_PARAM; } + completion_list_cpu[completion_index].xferReqRingGpu = xferReqRingGpu + idx; + lastPostedReq.fetch_add(1, std::memory_order_relaxed); + ++treq->postedCount; } - return NIXL_IN_PROG; + return treq->postedCount == 0 ? treq->postStatus : NIXL_IN_PROG; } nixl_status_t @@ -1300,7 +1318,8 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; uint32_t completion_index; - for (uint32_t idx : treq->positions) { + for (size_t i = 0; i < treq->postedCount; ++i) { + const uint32_t idx = treq->positions[i]; completion_index = xferReqRingCpu[idx].id & (DOCA_MAX_COMPLETION_INFLIGHT_MASK); if (((volatile docaXferCompletion *)completion_list_cpu)[completion_index].completed != 1) { @@ -1308,27 +1327,36 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { } } - for (uint32_t idx : treq->positions) { + for (size_t i = 0; i < treq->postedCount; ++i) { + const uint32_t idx = treq->positions[i]; *((volatile uint8_t *)&xferReqRingCpu[idx].in_use) = 0; NIXL_INFO << "DOCA checkXfer pos " << idx << " COMPLETED!"; } - return NIXL_SUCCESS; + return treq->postStatus; } nixl_status_t nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { + auto *treq = static_cast(handle); + if (treq->postedCount == 0) { + for (uint32_t idx : treq->positions) { + xferReqReserved_[idx].store(false, std::memory_order_release); + } + delete treq; + return NIXL_SUCCESS; + } + nixl_status_t status = checkXfer(handle); if (status == NIXL_IN_PROG) { - return status; + return NIXL_ERR_BACKEND; } - auto *treq = static_cast(handle); for (uint32_t idx : treq->positions) { - xferReqReserved[idx].store(false, std::memory_order_release); + xferReqReserved_[idx].store(false, std::memory_order_release); } delete treq; - return status; + return NIXL_SUCCESS; } nixl_status_t diff --git a/src/plugins/gpunetio/gpunetio_backend.h b/src/plugins/gpunetio/gpunetio_backend.h index 8cc1904b61..bceab6bb51 100644 --- a/src/plugins/gpunetio/gpunetio_backend.h +++ b/src/plugins/gpunetio/gpunetio_backend.h @@ -162,11 +162,12 @@ class nixlDocaEngine : public nixlBackendEngine { cudaStream_t wait_stream; mutable std::atomic xferStream; mutable std::atomic lastPostedReq; + mutable std::mutex postLock_; struct docaXferReqGpu *xferReqRingGpu; struct docaXferReqGpu *xferReqRingCpu; mutable std::atomic xferRingPos; - mutable std::array xferReqReserved; + mutable std::array xferReqReserved_; struct docaXferCompletion *completion_list_gpu; struct docaXferCompletion *completion_list_cpu; @@ -195,6 +196,8 @@ class nixlDocaEngine : public nixlBackendEngine { uint32_t devId; std::vector positions; uintptr_t backendHandleGpu; + size_t postedCount = 0; + nixl_status_t postStatus = NIXL_SUCCESS; nixlDocaBckndReq() : nixlBackendReqH() {} diff --git a/test/gtest/unit/agent/agent.cpp b/test/gtest/unit/agent/agent.cpp index a0eb0dd245..24a16ba4dd 100644 --- a/test/gtest/unit/agent/agent.cpp +++ b/test/gtest/unit/agent/agent.cpp @@ -433,6 +433,45 @@ namespace agent { EXPECT_EQ(local_agent_->releaseXferReq(xfer_req), NIXL_SUCCESS); } + TEST_F(dualAgentBridgeFixture, FailedActiveReleaseKeepsRequestPollable) { + DualAgentSetup s(DRAM_SEG); + setupDualAgent(s); + + nixl_xfer_dlist_t local_xfer_dlist(DRAM_SEG), remote_xfer_dlist(DRAM_SEG); + local_xfer_dlist.addDesc(s.local_blob.getDesc()); + remote_xfer_dlist.addDesc(s.remote_blob.getDesc()); + + nixlBackendReqH backend_request; + auto &engine = local_agent_helper_->getGMockEngine(); + EXPECT_CALL(engine, prepXfer) + .WillOnce([&](const auto &, + const auto &, + const auto &, + const auto &, + nixlBackendReqH *&request, + const auto *) { + request = &backend_request; + return NIXL_SUCCESS; + }); + EXPECT_CALL(engine, postXfer).WillOnce(testing::Return(NIXL_IN_PROG)); + EXPECT_CALL(engine, checkXfer) + .WillOnce(testing::Return(NIXL_IN_PROG)) + .WillOnce(testing::Return(NIXL_SUCCESS)); + EXPECT_CALL(engine, releaseReqH) + .WillOnce(testing::Return(NIXL_ERR_BACKEND)) + .WillOnce(testing::Return(NIXL_SUCCESS)); + + nixlXferReqH *xfer_req; + EXPECT_EQ( + local_agent_->createXferReq( + NIXL_WRITE, local_xfer_dlist, remote_xfer_dlist, s.remote_agent_name, xfer_req), + NIXL_SUCCESS); + EXPECT_EQ(local_agent_->postXferReq(xfer_req), NIXL_IN_PROG); + EXPECT_EQ(local_agent_->releaseXferReq(xfer_req), NIXL_ERR_REPOST_ACTIVE); + EXPECT_EQ(local_agent_->getXferStatus(xfer_req), NIXL_SUCCESS); + EXPECT_EQ(local_agent_->releaseXferReq(xfer_req), NIXL_SUCCESS); + } + TEST_F(dualAgentBridgeFixture, PrepMemViewRemoteDRAM) { DualAgentSetup s(DRAM_SEG); setupDualAgent(s, /*register_local=*/false); From c7471528635960ed2a27ee33389565bbb3ad25d2 Mon Sep 17 00:00:00 2001 From: foraxe Date: Tue, 1 Sep 2026 02:37:16 +0800 Subject: [PATCH 03/25] fix(gpunetio): transmit notification terminator Signed-off-by: foraxe (cherry picked from commit d2f1557d91b3dbb35eabc86e79803ef65fc78061) --- src/plugins/gpunetio/gpunetio_backend.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index 2b430f6836..831f5324c3 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -1250,12 +1250,11 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, final_request.has_notif_msg_idx = (notif->send_pi.fetch_add(1) & (notif->elems_num - 1)); notif_addr = (uintptr_t)(notif->send_addr + (final_request.has_notif_msg_idx * notif->elems_size)); - final_request.msg_sz = newMsg.size(); + final_request.msg_sz = newMsg.size() + 1; final_request.lbuf_notif = notif_addr; final_request.lkey_notif = notif->send_mr->get_lkey(); - memcpy((void *)notif_addr, newMsg.c_str(), newMsg.size()); - reinterpret_cast(notif_addr)[newMsg.size()] = '\0'; + memcpy((void *)notif_addr, newMsg.c_str(), final_request.msg_sz); NIXL_INFO << "DOCA prepXfer with notif to " << remote_agent << " at " << final_request.has_notif_msg_idx << " msg " << newMsg << " to " << remote_agent; From fb26d82195d8c746276bb3fca879c69aacff047f Mon Sep 17 00:00:00 2001 From: foraxe Date: Tue, 1 Sep 2026 03:17:07 +0800 Subject: [PATCH 04/25] fix(gpunetio): latch request completion state Signed-off-by: foraxe (cherry picked from commit 8b0883b70a6a009a031d8333562d1a958e56ba90) --- src/plugins/gpunetio/gpunetio_backend.cpp | 14 +++++++++++--- src/plugins/gpunetio/gpunetio_backend.h | 1 + 2 files changed, 12 insertions(+), 3 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index 831f5324c3..2769c51fca 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -1289,6 +1289,7 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, treq->postedCount = 0; treq->postStatus = NIXL_SUCCESS; + treq->completed = false; for (uint32_t idx : treq->positions) { std::lock_guard lock(postLock_); const uint32_t completion_index = @@ -1317,6 +1318,10 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; uint32_t completion_index; + if (treq->completed) { + return treq->postStatus; + } + for (size_t i = 0; i < treq->postedCount; ++i) { const uint32_t idx = treq->positions[i]; completion_index = xferReqRingCpu[idx].id & (DOCA_MAX_COMPLETION_INFLIGHT_MASK); @@ -1332,6 +1337,7 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { NIXL_INFO << "DOCA checkXfer pos " << idx << " COMPLETED!"; } + treq->completed = true; return treq->postStatus; } @@ -1346,9 +1352,11 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { return NIXL_SUCCESS; } - nixl_status_t status = checkXfer(handle); - if (status == NIXL_IN_PROG) { - return NIXL_ERR_BACKEND; + if (!treq->completed) { + nixl_status_t status = checkXfer(handle); + if (status == NIXL_IN_PROG) { + return NIXL_ERR_BACKEND; + } } for (uint32_t idx : treq->positions) { diff --git a/src/plugins/gpunetio/gpunetio_backend.h b/src/plugins/gpunetio/gpunetio_backend.h index bceab6bb51..c156456910 100644 --- a/src/plugins/gpunetio/gpunetio_backend.h +++ b/src/plugins/gpunetio/gpunetio_backend.h @@ -198,6 +198,7 @@ class nixlDocaEngine : public nixlBackendEngine { uintptr_t backendHandleGpu; size_t postedCount = 0; nixl_status_t postStatus = NIXL_SUCCESS; + bool completed = false; nixlDocaBckndReq() : nixlBackendReqH() {} From b2500882d65e889016dd7782f307eb296dfa2695 Mon Sep 17 00:00:00 2001 From: foraxe Date: Tue, 1 Sep 2026 04:14:27 +0800 Subject: [PATCH 05/25] fix(gpunetio): reject repost after launch failure Signed-off-by: foraxe (cherry picked from commit dd8d094f3deb5d209741ff4ab29953050a41e9dc) --- src/plugins/gpunetio/gpunetio_backend.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index 2769c51fca..eadf68f31b 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -1286,6 +1286,9 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, if (operation != NIXL_READ && operation != NIXL_WRITE) { return NIXL_ERR_INVALID_PARAM; } + if (treq->postStatus != NIXL_SUCCESS) { + return treq->postStatus; + } treq->postedCount = 0; treq->postStatus = NIXL_SUCCESS; From b659c2c30d9e227f37161e7156c336625f2e740c Mon Sep 17 00:00:00 2001 From: foraxe Date: Tue, 1 Sep 2026 04:51:46 +0800 Subject: [PATCH 06/25] fix(gpunetio): serialize completion retirement Signed-off-by: foraxe (cherry picked from commit 4e6d61c930cbc679e74ae62c3b6fe48736415237) --- src/plugins/gpunetio/gpunetio_backend.cpp | 33 +++++++++++++++++------ src/plugins/gpunetio/gpunetio_backend.h | 4 ++- 2 files changed, 28 insertions(+), 9 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index eadf68f31b..ee849ba364 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -1292,7 +1292,8 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, treq->postedCount = 0; treq->postStatus = NIXL_SUCCESS; - treq->completed = false; + treq->completionState.store(nixlDocaBckndReq::CompletionState::IN_PROGRESS, + std::memory_order_release); for (uint32_t idx : treq->positions) { std::lock_guard lock(postLock_); const uint32_t completion_index = @@ -1321,7 +1322,12 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; uint32_t completion_index; - if (treq->completed) { + auto state = treq->completionState.load(std::memory_order_acquire); + if (state == nixlDocaBckndReq::CompletionState::COMPLETE) { + return treq->postStatus; + } + if (state == nixlDocaBckndReq::CompletionState::COMPLETING) { + treq->completionState.wait(state, std::memory_order_acquire); return treq->postStatus; } @@ -1334,13 +1340,23 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { } } - for (size_t i = 0; i < treq->postedCount; ++i) { - const uint32_t idx = treq->positions[i]; - *((volatile uint8_t *)&xferReqRingCpu[idx].in_use) = 0; - NIXL_INFO << "DOCA checkXfer pos " << idx << " COMPLETED!"; + auto expected = nixlDocaBckndReq::CompletionState::IN_PROGRESS; + if (treq->completionState.compare_exchange_strong( + expected, + nixlDocaBckndReq::CompletionState::COMPLETING, + std::memory_order_acq_rel)) { + for (size_t i = 0; i < treq->postedCount; ++i) { + const uint32_t idx = treq->positions[i]; + *((volatile uint8_t *)&xferReqRingCpu[idx].in_use) = 0; + NIXL_INFO << "DOCA checkXfer pos " << idx << " COMPLETED!"; + } + treq->completionState.store(nixlDocaBckndReq::CompletionState::COMPLETE, + std::memory_order_release); + treq->completionState.notify_all(); + } else if (expected == nixlDocaBckndReq::CompletionState::COMPLETING) { + treq->completionState.wait(expected, std::memory_order_acquire); } - treq->completed = true; return treq->postStatus; } @@ -1355,7 +1371,8 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { return NIXL_SUCCESS; } - if (!treq->completed) { + if (treq->completionState.load(std::memory_order_acquire) != + nixlDocaBckndReq::CompletionState::COMPLETE) { nixl_status_t status = checkXfer(handle); if (status == NIXL_IN_PROG) { return NIXL_ERR_BACKEND; diff --git a/src/plugins/gpunetio/gpunetio_backend.h b/src/plugins/gpunetio/gpunetio_backend.h index c156456910..4471040fc6 100644 --- a/src/plugins/gpunetio/gpunetio_backend.h +++ b/src/plugins/gpunetio/gpunetio_backend.h @@ -192,13 +192,15 @@ class nixlDocaEngine : public nixlBackendEngine { class nixlDocaBckndReq : public nixlBackendReqH { private: public: + enum class CompletionState : uint8_t { IN_PROGRESS, COMPLETING, COMPLETE }; + cudaStream_t stream; uint32_t devId; std::vector positions; uintptr_t backendHandleGpu; size_t postedCount = 0; nixl_status_t postStatus = NIXL_SUCCESS; - bool completed = false; + std::atomic completionState{CompletionState::IN_PROGRESS}; nixlDocaBckndReq() : nixlBackendReqH() {} From a5e844136d20366d299286730a17d512c8a88bea Mon Sep 17 00:00:00 2001 From: foraxe Date: Tue, 1 Sep 2026 04:52:13 +0800 Subject: [PATCH 07/25] style(gpunetio): format completion transition Signed-off-by: foraxe (cherry picked from commit 6846c18d502841c646e49bbf4ee5637981cc0474) --- src/plugins/gpunetio/gpunetio_backend.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index ee849ba364..83e7ced9a3 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -1342,9 +1342,7 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { auto expected = nixlDocaBckndReq::CompletionState::IN_PROGRESS; if (treq->completionState.compare_exchange_strong( - expected, - nixlDocaBckndReq::CompletionState::COMPLETING, - std::memory_order_acq_rel)) { + expected, nixlDocaBckndReq::CompletionState::COMPLETING, std::memory_order_acq_rel)) { for (size_t i = 0; i < treq->postedCount; ++i) { const uint32_t idx = treq->positions[i]; *((volatile uint8_t *)&xferReqRingCpu[idx].in_use) = 0; From 1ca6efcc9206f87f1428395fa0eeb127ba91170c Mon Sep 17 00:00:00 2001 From: foraxe Date: Tue, 1 Sep 2026 05:15:46 +0800 Subject: [PATCH 08/25] refactor(gpunetio): name completion state consistently Signed-off-by: foraxe (cherry picked from commit 57dfc9485eafcd51b8563cf40c7afab8fc93bce1) --- src/plugins/gpunetio/gpunetio_backend.cpp | 47 +++++++++++++++-------- src/plugins/gpunetio/gpunetio_backend.h | 7 +++- 2 files changed, 35 insertions(+), 19 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index 83e7ced9a3..de802be166 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -1292,7 +1292,7 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, treq->postedCount = 0; treq->postStatus = NIXL_SUCCESS; - treq->completionState.store(nixlDocaBckndReq::CompletionState::IN_PROGRESS, + treq->completionState.store(nixlDocaBckndReq::completion_state::IN_PROGRESS, std::memory_order_release); for (uint32_t idx : treq->positions) { std::lock_guard lock(postLock_); @@ -1323,10 +1323,10 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { uint32_t completion_index; auto state = treq->completionState.load(std::memory_order_acquire); - if (state == nixlDocaBckndReq::CompletionState::COMPLETE) { + if (state == nixlDocaBckndReq::completion_state::COMPLETE) { return treq->postStatus; } - if (state == nixlDocaBckndReq::CompletionState::COMPLETING) { + if (state == nixlDocaBckndReq::completion_state::COMPLETING) { treq->completionState.wait(state, std::memory_order_acquire); return treq->postStatus; } @@ -1340,22 +1340,35 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { } } - auto expected = nixlDocaBckndReq::CompletionState::IN_PROGRESS; - if (treq->completionState.compare_exchange_strong( - expected, nixlDocaBckndReq::CompletionState::COMPLETING, std::memory_order_acq_rel)) { - for (size_t i = 0; i < treq->postedCount; ++i) { - const uint32_t idx = treq->positions[i]; + retireRequest(treq); + return treq->postStatus; +} + +void +nixlDocaEngine::retireRequest(nixlDocaBckndReq *request) const { + auto state = request->completionState.load(std::memory_order_acquire); + if (state == nixlDocaBckndReq::completion_state::COMPLETE) { + return; + } + if (state == nixlDocaBckndReq::completion_state::COMPLETING) { + request->completionState.wait(state, std::memory_order_acquire); + return; + } + + auto expected = nixlDocaBckndReq::completion_state::IN_PROGRESS; + if (request->completionState.compare_exchange_strong( + expected, nixlDocaBckndReq::completion_state::COMPLETING, std::memory_order_acq_rel)) { + for (size_t i = 0; i < request->postedCount; ++i) { + const uint32_t idx = request->positions[i]; *((volatile uint8_t *)&xferReqRingCpu[idx].in_use) = 0; - NIXL_INFO << "DOCA checkXfer pos " << idx << " COMPLETED!"; + NIXL_INFO << "DOCA retireRequest pos " << idx << " COMPLETED!"; } - treq->completionState.store(nixlDocaBckndReq::CompletionState::COMPLETE, - std::memory_order_release); - treq->completionState.notify_all(); - } else if (expected == nixlDocaBckndReq::CompletionState::COMPLETING) { - treq->completionState.wait(expected, std::memory_order_acquire); + request->completionState.store(nixlDocaBckndReq::completion_state::COMPLETE, + std::memory_order_release); + request->completionState.notify_all(); + } else if (expected == nixlDocaBckndReq::completion_state::COMPLETING) { + request->completionState.wait(expected, std::memory_order_acquire); } - - return treq->postStatus; } nixl_status_t @@ -1370,7 +1383,7 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { } if (treq->completionState.load(std::memory_order_acquire) != - nixlDocaBckndReq::CompletionState::COMPLETE) { + nixlDocaBckndReq::completion_state::COMPLETE) { nixl_status_t status = checkXfer(handle); if (status == NIXL_IN_PROG) { return NIXL_ERR_BACKEND; diff --git a/src/plugins/gpunetio/gpunetio_backend.h b/src/plugins/gpunetio/gpunetio_backend.h index 4471040fc6..0081ffc1f0 100644 --- a/src/plugins/gpunetio/gpunetio_backend.h +++ b/src/plugins/gpunetio/gpunetio_backend.h @@ -192,7 +192,7 @@ class nixlDocaEngine : public nixlBackendEngine { class nixlDocaBckndReq : public nixlBackendReqH { private: public: - enum class CompletionState : uint8_t { IN_PROGRESS, COMPLETING, COMPLETE }; + enum class completion_state : uint8_t { IN_PROGRESS, COMPLETING, COMPLETE }; cudaStream_t stream; uint32_t devId; @@ -200,13 +200,16 @@ class nixlDocaEngine : public nixlBackendEngine { uintptr_t backendHandleGpu; size_t postedCount = 0; nixl_status_t postStatus = NIXL_SUCCESS; - std::atomic completionState{CompletionState::IN_PROGRESS}; + std::atomic completionState{completion_state::IN_PROGRESS}; nixlDocaBckndReq() : nixlBackendReqH() {} ~nixlDocaBckndReq() {} }; + void + retireRequest(nixlDocaBckndReq *request) const; + nixl_status_t progressThreadStart(); void From 56c4b8c5c02ad0f4f406071a4dacb66ad4cab1f2 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Mon, 7 Sep 2026 19:17:03 +0800 Subject: [PATCH 09/25] perf(gpunetio): progress independent QPs without global completion blocking Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 345 ++++++++++----- src/plugins/gpunetio/gpunetio_backend.h | 15 +- src/plugins/gpunetio/gpunetio_backend_aux.h | 73 +++- src/plugins/gpunetio/gpunetio_kernels.cu | 442 ++++++++++++++------ 4 files changed, 609 insertions(+), 266 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index de802be166..8d95fb3866 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -43,18 +43,25 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) } result = doca_log_backend_create_standard(); - if (result != DOCA_SUCCESS) throw std::invalid_argument("Can't initialize doca log"); + if (result != DOCA_SUCCESS) { + throw std::invalid_argument("Can't initialize doca log"); + } result = doca_log_backend_create_with_file_sdk(stderr, &sdk_log); - if (result != DOCA_SUCCESS) throw std::invalid_argument("Can't initialize doca log"); + if (result != DOCA_SUCCESS) { + throw std::invalid_argument("Can't initialize doca log"); + } result = doca_log_backend_set_sdk_level(sdk_log, DOCA_LOG_LEVEL_ERROR); - if (result != DOCA_SUCCESS) throw std::invalid_argument("Can't initialize doca log"); + if (result != DOCA_SUCCESS) { + throw std::invalid_argument("Can't initialize doca log"); + } NIXL_INFO << "DOCA network devices "; // Temporary: will extend to more GPUs in a dedicated PR - if (custom_params->count("network_devices") > 1) + if (custom_params->count("network_devices") > 1) { throw std::invalid_argument("Only 1 network device is allowed"); + } if (custom_params->count("network_devices") == 0 || (*custom_params)["network_devices"] == "" || (*custom_params)["network_devices"] == "all") { @@ -69,8 +76,9 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) if (custom_params->count("oob_interface") > 0) { NIXL_INFO << "DOCA network devices "; // Temporary: will extend to more GPUs in a dedicated PR - if (custom_params->count("oob_interface") > 1) + if (custom_params->count("oob_interface") > 1) { throw std::invalid_argument("Only 1 oob interface is allowed"); + } oobdev = absl::StrSplit((*custom_params)["oob_interface"], " "); NIXL_INFO << "Using oob interface" << oobdev[0]; @@ -79,8 +87,9 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) NIXL_INFO << "DOCA GPU devices: "; // Temporary: will extend to more GPUs in a dedicated PR - if (custom_params->count("gpu_devices") > 1) + if (custom_params->count("gpu_devices") > 1) { throw std::invalid_argument("Only 1 GPU device is allowed"); + } if (custom_params->count("gpu_devices") == 0 || (*custom_params)["gpu_devices"] == "" || (*custom_params)["gpu_devices"] == "all") { @@ -96,9 +105,12 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) NIXL_INFO << std::endl; nstreams = 0; - if (custom_params->count("cuda_streams") != 0 && (*custom_params)["cuda_streams"] != "") + if (custom_params->count("cuda_streams") != 0 && (*custom_params)["cuda_streams"] != "") { nstreams = std::stoi((*custom_params)["cuda_streams"]); - if (nstreams == 0) nstreams = DOCA_POST_STREAM_NUM; + } + if (nstreams == 0) { + nstreams = DOCA_POST_STREAM_NUM; + } NIXL_INFO << "CUDA streams used for pool mode: " << nstreams; @@ -116,7 +128,9 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) } pd = doca_verbs_bridge_verbs_pd_get_ibv_pd(verbs_pd); - if (pd == NULL) throw std::invalid_argument("Failed to get ibv_pd"); + if (pd == NULL) { + throw std::invalid_argument("Failed to get ibv_pd"); + } result = doca_rdma_bridge_open_dev_from_pd(pd, &ddev); if (result != DOCA_SUCCESS) { @@ -141,8 +155,9 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) if (port_attr.link_layer == IBV_LINK_LAYER_INFINIBAND) { result = create_verbs_ah_attr( verbs_context, gid_index, DOCA_VERBS_ADDR_TYPE_IB_NO_GRH, &verbs_ah_attr); - if (result != DOCA_SUCCESS) + if (result != DOCA_SUCCESS) { throw std::invalid_argument("Failed to create doca verbs ah attributes"); + } lid = port_attr.lid; } else { @@ -168,8 +183,9 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) cudaFree(0); result = doca_gpu_create(pciBusId, &item.second); - if (result != DOCA_SUCCESS) + if (result != DOCA_SUCCESS) { NIXL_ERROR << "Failed to create DOCA GPU device " << doca_error_get_descr(result); + } } if (oobdev.size() > 0 && oobdev[0] != "") { @@ -179,7 +195,8 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) NIXL_DEBUG << "Eth IP address " << static_cast(ipv4_addr[0]) << " " << static_cast(ipv4_addr[1]) << " " << static_cast(ipv4_addr[2]) << " " - << static_cast(ipv4_addr[3]) << " " << "ifface " << oobdev[0].c_str(); + << static_cast(ipv4_addr[3]) << " " + << "ifface " << oobdev[0].c_str(); } else { doca_devinfo_get_ipv4_addr( doca_dev_as_devinfo(ddev), (uint8_t *)ipv4_addr, DOCA_DEVINFO_IPV4_ADDR_SIZE); @@ -217,53 +234,26 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) cudaMemset(xferReqRingGpu, 0, sizeof(struct docaXferReqGpu) * DOCA_XFER_REQ_MAX), "Failed to memset GPU memory"); - result = doca_gpu_mem_alloc(gdevs[0].second, - sizeof(uint64_t), - 4096, - DOCA_GPU_MEM_TYPE_GPU, - (void **)&last_rsvd_flags, - nullptr); - if (result != DOCA_SUCCESS || last_rsvd_flags == nullptr) { - NIXL_ERROR << "Function doca_gpu_mem_alloc return " << doca_error_get_descr(result); - } - - nixlDocaEngineCheckCudaError(cudaMemset(last_rsvd_flags, 0, sizeof(uint64_t)), - "Failed to memset GPU memory"); - - result = doca_gpu_mem_alloc(gdevs[0].second, - sizeof(uint64_t), - 4096, - DOCA_GPU_MEM_TYPE_GPU, - (void **)&last_posted_flags, - nullptr); - if (result != DOCA_SUCCESS || last_posted_flags == nullptr) { - NIXL_ERROR << "Function doca_gpu_mem_alloc return " << doca_error_get_descr(result); - } - - nixlDocaEngineCheckCudaError(cudaMemset(last_posted_flags, 0, sizeof(uint64_t)), - "Failed to memset GPU memory"); - nixlDocaEngineCheckCudaError(cudaStreamCreateWithFlags(&wait_stream, cudaStreamNonBlocking), "Failed to create CUDA stream"); - for (int i = 0; i < nstreams; i++) + for (int i = 0; i < nstreams; i++) { nixlDocaEngineCheckCudaError( cudaStreamCreateWithFlags(&post_stream[i], cudaStreamNonBlocking), "Failed to create CUDA stream"); + } xferStream = 0; result = doca_gpu_mem_alloc(gdevs[0].second, - sizeof(struct docaXferCompletion) * DOCA_MAX_COMPLETION_INFLIGHT, + sizeof(struct docaProgressState), 4096, DOCA_GPU_MEM_TYPE_CPU_GPU, - (void **)&completion_list_gpu, - (void **)&completion_list_cpu); - if (result != DOCA_SUCCESS || completion_list_gpu == nullptr || - completion_list_cpu == nullptr) { + (void **)&progress_state_gpu, + (void **)&progress_state_cpu); + if (result != DOCA_SUCCESS || progress_state_gpu == nullptr || progress_state_cpu == nullptr) { NIXL_ERROR << "Function doca_gpu_mem_alloc return " << doca_error_get_descr(result); } - memset( - completion_list_cpu, 0, sizeof(struct docaXferCompletion) * DOCA_MAX_COMPLETION_INFLIGHT); + memset(progress_state_cpu, 0, sizeof(struct docaProgressState)); // DOCA_GPU_MEM_TYPE_GPU_CPU == GDRCopy result = doca_gpu_mem_alloc(gdevs[0].second, @@ -289,7 +279,7 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) } } - ((volatile uint8_t *)wait_exit_cpu)[0] = 0; + std::atomic_ref(*wait_exit_cpu).store(0, std::memory_order_release); result = doca_gpu_mem_alloc(gdevs[0].second, sizeof(struct docaNotif), @@ -313,35 +303,22 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) memset(notif_progress_cpu, 0, sizeof(struct docaNotif)); - result = doca_gpu_mem_alloc(gdevs[0].second, - sizeof(struct docaNotif), - 4096, - DOCA_GPU_MEM_TYPE_CPU_GPU, - (void **)¬if_send_gpu, - (void **)¬if_send_cpu); - if (result != DOCA_SUCCESS || notif_send_gpu == nullptr || notif_send_cpu == nullptr) { - NIXL_ERROR << "Function doca_gpu_mem_alloc return " << doca_error_get_descr(result); - } - - memset(notif_send_cpu, 0, sizeof(struct docaNotif)); - // We may need a GPU warmup with relevant DOCA engine kernels - doca_kernel_write(0, nullptr, nullptr, 0); - doca_kernel_read(0, nullptr, nullptr, 0); + doca_kernel_write(0, nullptr, nullptr, nullptr, nullptr, 0); + doca_kernel_read(0, nullptr, nullptr, nullptr, nullptr, 0); nixlDocaEngineCheckCudaError(cudaStreamSynchronize(0), "stream synchronize"); // Warmup doca_kernel_progress( - wait_stream, nullptr, notif_fill_gpu, notif_progress_gpu, notif_send_gpu, wait_exit_gpu); + wait_stream, nullptr, nullptr, notif_fill_gpu, notif_progress_gpu, wait_exit_gpu); nixlDocaEngineCheckCudaError(cudaStreamSynchronize(wait_stream), "stream synchronize"); doca_kernel_progress(wait_stream, - completion_list_gpu, + xferReqRingGpu, + progress_state_gpu, notif_fill_gpu, notif_progress_gpu, - notif_send_gpu, wait_exit_gpu); - lastPostedReq = 0; xferRingPos = 0; progressThreadStart(); @@ -358,14 +335,10 @@ nixlDocaEngine::~nixlDocaEngine() { NIXL_DEBUG << "Before progressThreadStop "; progressThreadStop(); - ((volatile uint8_t *)wait_exit_cpu)[0] = 1; + std::atomic_ref(*wait_exit_cpu).store(1, std::memory_order_release); NIXL_DEBUG << "Before cudaStreamSynchronize "; nixlDocaEngineCheckCudaError(cudaStreamSynchronize(wait_stream), "stream synchronize"); nixlDocaEngineCheckCudaError(cudaStreamDestroy(wait_stream), "stream destroy"); - doca_gpu_mem_free(gdevs[0].second, wait_exit_gpu); - doca_gpu_mem_free(gdevs[0].second, xferReqRingGpu); - doca_gpu_mem_free(gdevs[0].second, last_rsvd_flags); - doca_gpu_mem_free(gdevs[0].second, last_posted_flags); for (int i = 0; i < nstreams; i++) { NIXL_DEBUG << "Before cudaStreamSynchronize post_stream " << i; @@ -373,26 +346,35 @@ nixlDocaEngine::~nixlDocaEngine() { nixlDocaEngineCheckCudaError(cudaStreamDestroy(post_stream[i]), "stream destroy"); } + for (auto *progress : qp_progress_gpu_) { + doca_gpu_mem_free(gdevs[0].second, progress); + } + qp_progress_gpu_.clear(); + doca_gpu_mem_free(gdevs[0].second, xferReqRingGpu); + doca_gpu_mem_free(gdevs[0].second, progress_state_gpu); + doca_gpu_mem_free(gdevs[0].second, wait_exit_gpu); + NIXL_DEBUG << "Before nixlDocaDestroyNotif "; - for (auto notif : notifMap) + for (auto notif : notifMap) { nixlDocaDestroyNotif(gdevs[0].second, notif.second); + } doca_gpu_mem_free(gdevs[0].second, notif_fill_gpu); doca_gpu_mem_free(gdevs[0].second, notif_progress_gpu); - doca_gpu_mem_free(gdevs[0].second, notif_send_gpu); - doca_gpu_mem_free(gdevs[0].second, completion_list_gpu); NIXL_DEBUG << "Before qpMap.clear "; qpMap.clear(); result = doca_dev_close(ddev); - if (result != DOCA_SUCCESS) + if (result != DOCA_SUCCESS) { NIXL_ERROR << "Failed to close DOCA device " << doca_error_get_descr(result); + } result = doca_gpu_destroy(gdevs[0].second); - if (result != DOCA_SUCCESS) + if (result != DOCA_SUCCESS) { NIXL_ERROR << "Failed to close DOCA GPU device " << doca_error_get_descr(result); + } } /**************************************** @@ -457,8 +439,9 @@ nixlDocaEngine::nixlDocaInitNotif(const std::string &remote_agent, doca_dev *dev std::atomic_thread_fence(std::memory_order_seq_cst); ((volatile struct docaNotif *)notif_fill_cpu)->qp_gpu = qpMap[remote_agent]->qp_notif->get_qp_gpu_dev(); - while (((volatile struct docaNotif *)notif_fill_cpu)->qp_gpu != nullptr) + while (((volatile struct docaNotif *)notif_fill_cpu)->qp_gpu != nullptr) { ; + } NIXL_INFO << "nixlDocaInitNotif added new qp for " << remote_agent << std::endl; @@ -575,6 +558,7 @@ nixlDocaEngine::getGpuCudaId() { nixl_status_t nixlDocaEngine::addRdmaQp(const std::string &remote_agent) { + doca_error_t result; struct nixlDocaRdmaQp *rdma_qp; std::lock_guard lock(qpLock); @@ -589,7 +573,6 @@ nixlDocaEngine::addRdmaQp(const std::string &remote_agent) { NIXL_DEBUG << "DOCA addRdmaQp for remote " << remote_agent << std::endl; rdma_qp = new struct nixlDocaRdmaQp; - try { rdma_qp->qp_data = std::make_unique(gdevs[0].second, @@ -625,6 +608,55 @@ nixlDocaEngine::addRdmaQp(const std::string &remote_agent) { rdma_qp->qpn_notif = doca_verbs_qp_get_qpn(rdma_qp->qp_notif->get_qp()); + result = doca_gpu_mem_alloc(gdevs[0].second, + sizeof(struct docaQpProgress), + 4096, + DOCA_GPU_MEM_TYPE_GPU, + (void **)&rdma_qp->progress_gpu, + nullptr); + if (result != DOCA_SUCCESS || rdma_qp->progress_gpu == nullptr) { + NIXL_ERROR << "Failed to allocate QP progress state " << doca_error_get_descr(result); + delete rdma_qp; + return NIXL_ERR_BACKEND; + } + + cudaStream_t init_stream = nullptr; + int previous_cuda_device = 0; + bool cuda_device_switched = false; + cudaError_t cuda_result = cudaGetDevice(&previous_cuda_device); + if (cuda_result == cudaSuccess) { + cuda_result = cudaSetDevice(gdevs[0].first); + cuda_device_switched = cuda_result == cudaSuccess; + } + if (cuda_result == cudaSuccess) { + cuda_result = cudaStreamCreateWithFlags(&init_stream, cudaStreamNonBlocking); + } + if (cuda_result == cudaSuccess) { + cuda_result = + cudaMemsetAsync(rdma_qp->progress_gpu, 0, sizeof(struct docaQpProgress), init_stream); + } + if (cuda_result == cudaSuccess) { + cuda_result = cudaStreamSynchronize(init_stream); + } + if (cuda_result == cudaSuccess) { + cuda_result = cudaStreamDestroy(init_stream); + init_stream = nullptr; + } + if (init_stream != nullptr) { + cudaStreamDestroy(init_stream); + } + if (cuda_device_switched && previous_cuda_device != gdevs[0].first) { + cudaSetDevice(previous_cuda_device); + } + if (cuda_result != cudaSuccess) { + NIXL_ERROR << "Failed to initialize QP progress state " << cudaGetErrorString(cuda_result); + doca_gpu_mem_free(gdevs[0].second, rdma_qp->progress_gpu); + delete rdma_qp; + return NIXL_ERR_BACKEND; + } + + qp_progress_gpu_.push_back(rdma_qp->progress_gpu); + qpMap[remote_agent] = rdma_qp; NIXL_DEBUG << "DOCA addRdmaQp new QP added for " << remote_agent; @@ -1073,8 +1105,9 @@ nixlDocaEngine::loadRemoteMD(const nixlBlobDesc &input, md->conn = conn; std::stringstream ss(input.metaInfo.data()); - while (std::getline(ss, token, info_delimiter)) + while (std::getline(ss, token, info_delimiter)) { tokens.push_back(token); + } uint32_t rkey = static_cast(atoi(tokens[0].c_str())); uintptr_t addr = static_cast(atol(tokens[1].c_str())); @@ -1109,7 +1142,7 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, const std::string &remote_agent, nixlBackendReqH *&handle, const nixl_opt_b_args_t *opt_args) const { - uint32_t pos; + uint32_t pos = 0; nixlDocaBckndReq *treq; nixlDocaPrivateMetadata *lmd; nixlDocaPublicMetadata *rmd; @@ -1187,6 +1220,7 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, }; treq->positions.reserve((lcnt + DOCA_XFER_REQ_SIZE - 1) / DOCA_XFER_REQ_SIZE); + treq->generations.reserve((lcnt + DOCA_XFER_REQ_SIZE - 1) / DOCA_XFER_REQ_SIZE); if (!reserve_position()) { abandon_request(); return NIXL_ERR_BACKEND; @@ -1196,6 +1230,10 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, do { docaXferReqGpu staged_req{}; staged_req.has_notif_msg_idx = DOCA_NOTIF_NULL; + staged_req.generation = xferReqRingCpu[pos].generation + 1; + staged_req.state = DOCA_XFER_STATE_PREPARED; + staged_req.data_state = DOCA_XFER_DATA_NONE; + staged_req.notif_state = DOCA_XFER_NOTIF_NONE; while (desc_offset < lcnt && staged_req.num < DOCA_XFER_REQ_SIZE) { const uint32_t idx = staged_req.num; @@ -1212,11 +1250,13 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, staged_req.num++; } - staged_req.last_rsvd = last_rsvd_flags; - staged_req.last_posted = last_posted_flags; staged_req.qp_data = rdma_qp->qp_data->get_qp_gpu_dev(); staged_req.qp_notif = rdma_qp->qp_notif->get_qp_gpu_dev(); + staged_req.qp_progress = rdma_qp->progress_gpu; memcpy(&xferReqRingCpu[pos], &staged_req, sizeof(staged_req)); + std::atomic_ref(xferReqRingCpu[pos].state) + .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); + treq->generations.push_back(staged_req.generation); if (desc_offset < lcnt && !reserve_position()) { abandon_request(); @@ -1295,22 +1335,23 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, treq->completionState.store(nixlDocaBckndReq::completion_state::IN_PROGRESS, std::memory_order_release); for (uint32_t idx : treq->positions) { - std::lock_guard lock(postLock_); - const uint32_t completion_index = - lastPostedReq.load(std::memory_order_relaxed) & DOCA_MAX_COMPLETION_INFLIGHT_MASK; - xferReqRingCpu[idx].id = completion_index; - completion_list_cpu[completion_index].xferReqRingGpu = nullptr; - completion_list_cpu[completion_index].completed = 0; - const doca_error_t result = operation == NIXL_READ ? - doca_kernel_read(treq->stream, xferReqRingCpu[idx].qp_data, xferReqRingGpu, idx) : - doca_kernel_write(treq->stream, xferReqRingCpu[idx].qp_data, xferReqRingGpu, idx); + doca_kernel_read(treq->stream, + xferReqRingCpu[idx].qp_data, + xferReqRingGpu, + progress_state_gpu, + wait_exit_gpu, + idx) : + doca_kernel_write(treq->stream, + xferReqRingCpu[idx].qp_data, + xferReqRingGpu, + progress_state_gpu, + wait_exit_gpu, + idx); if (result != DOCA_SUCCESS) { treq->postStatus = NIXL_ERR_BACKEND; break; } - completion_list_cpu[completion_index].xferReqRingGpu = xferReqRingGpu + idx; - lastPostedReq.fetch_add(1, std::memory_order_relaxed); ++treq->postedCount; } @@ -1320,8 +1361,14 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, nixl_status_t nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; - uint32_t completion_index; - + if (treq->postStatus != NIXL_SUCCESS) { + return treq->postStatus; + } + if (std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != + 0) { + treq->postStatus = NIXL_ERR_BACKEND; + return NIXL_ERR_BACKEND; + } auto state = treq->completionState.load(std::memory_order_acquire); if (state == nixlDocaBckndReq::completion_state::COMPLETE) { return treq->postStatus; @@ -1333,9 +1380,17 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { for (size_t i = 0; i < treq->postedCount; ++i) { const uint32_t idx = treq->positions[i]; - completion_index = xferReqRingCpu[idx].id & (DOCA_MAX_COMPLETION_INFLIGHT_MASK); - - if (((volatile docaXferCompletion *)completion_list_cpu)[completion_index].completed != 1) { + const uint32_t req_state = + std::atomic_ref(xferReqRingCpu[idx].state).load(std::memory_order_acquire); + if (xferReqRingCpu[idx].generation != treq->generations[i]) { + treq->postStatus = NIXL_ERR_BACKEND; + return NIXL_ERR_BACKEND; + } + if (req_state == DOCA_XFER_STATE_ERROR) { + treq->postStatus = NIXL_ERR_BACKEND; + return NIXL_ERR_BACKEND; + } + if (req_state != DOCA_XFER_STATE_COMPLETE) { return NIXL_IN_PROG; } } @@ -1360,7 +1415,6 @@ nixlDocaEngine::retireRequest(nixlDocaBckndReq *request) const { expected, nixlDocaBckndReq::completion_state::COMPLETING, std::memory_order_acq_rel)) { for (size_t i = 0; i < request->postedCount; ++i) { const uint32_t idx = request->positions[i]; - *((volatile uint8_t *)&xferReqRingCpu[idx].in_use) = 0; NIXL_INFO << "DOCA retireRequest pos " << idx << " COMPLETED!"; } request->completionState.store(nixlDocaBckndReq::completion_state::COMPLETE, @@ -1375,18 +1429,22 @@ nixl_status_t nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { auto *treq = static_cast(handle); if (treq->postedCount == 0) { + const nixl_status_t status = treq->postStatus; for (uint32_t idx : treq->positions) { xferReqReserved_[idx].store(false, std::memory_order_release); } delete treq; - return NIXL_SUCCESS; + return status; } if (treq->completionState.load(std::memory_order_acquire) != nixlDocaBckndReq::completion_state::COMPLETE) { nixl_status_t status = checkXfer(handle); if (status == NIXL_IN_PROG) { - return NIXL_ERR_BACKEND; + return NIXL_IN_PROG; + } + if (status != NIXL_SUCCESS) { + return status; } } @@ -1409,11 +1467,22 @@ nixlDocaEngine::getNotifs(notif_list_t ¬if_list) { // while getNotifs is running std::lock_guard lock(notifLock); for (auto ¬if : notifMap) { + if (std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != + 0 || + std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + return NIXL_ERR_BACKEND; + } ((volatile struct docaNotif *)notif_progress_cpu)->qp_gpu = qpMap[notif.first]->qp_notif->get_qp_gpu_dev(); std::atomic_thread_fence(std::memory_order_seq_cst); - while (((volatile struct docaNotif *)notif_progress_cpu)->qp_gpu != nullptr) - ; + while (((volatile struct docaNotif *)notif_progress_cpu)->qp_gpu != nullptr) { + if (std::atomic_ref(progress_state_cpu->failed) + .load(std::memory_order_acquire) != 0 || + std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + return NIXL_ERR_BACKEND; + } + std::this_thread::yield(); + } num_msg = ((volatile struct docaNotif *)notif_progress_cpu)->msg_num; while (num_msg > 0) { recv_idx = notif.second->recv_pi.load() & (DOCA_MAX_NOTIF_INFLIGHT - 1); @@ -1450,7 +1519,7 @@ nixlDocaEngine::getNotifs(notif_list_t ¬if_list) { } else { NIXL_ERROR << "getNotifs error message at " << num_msg << " size " << msg_src.size() << " msg " << msg_src; - break; + return NIXL_ERR_BACKEND; } } } @@ -1462,6 +1531,7 @@ nixl_status_t nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg) const { struct nixlDocaNotif *notif; uint32_t buf_idx; + uint32_t pos = 0; uintptr_t msg_buf; auto searchNotif = notifMap.find(remote_agent); @@ -1489,20 +1559,67 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg std::string newMsg = msg_tag_start + std::to_string((int)msg.size()) + msg_tag_end + msg; buf_idx = (notif->send_pi.fetch_add(1) & (notif->elems_num - 1)); msg_buf = (uintptr_t)notif->send_addr + (buf_idx * notif->elems_size); - memcpy((void *)msg_buf, newMsg.c_str(), newMsg.size()); + memcpy((void *)msg_buf, newMsg.c_str(), newMsg.size() + 1); NIXL_DEBUG << "genNotif to " << remote_agent << " msg size " << std::to_string((int)msg.size()) << " msg " << newMsg << " at " << buf_idx << " msg_buf " << msg_buf << "\n"; - std::lock_guard lock(notifSendLock); - ((volatile struct docaNotif *)notif_send_cpu)->msg_buf = msg_buf; - ((volatile struct docaNotif *)notif_send_cpu)->msg_lkey = notif->send_mr->get_lkey(); - ((volatile struct docaNotif *)notif_send_cpu)->msg_size = newMsg.size(); - std::atomic_thread_fence(std::memory_order_seq_cst); - ((volatile struct docaNotif *)notif_send_cpu)->qp_gpu = - searchQp->second->qp_notif->get_qp_gpu_dev(); - while (((volatile struct docaNotif *)notif_send_cpu)->qp_gpu != nullptr) - ; + bool ring_reserved = false; + for (uint32_t attempt = 0; attempt < DOCA_XFER_REQ_MAX; ++attempt) { + const uint32_t candidate = xferRingPos.fetch_add(1) & DOCA_XFER_REQ_MASK; + bool expected = false; + if (xferReqReserved_[candidate].compare_exchange_strong( + expected, true, std::memory_order_acq_rel)) { + pos = candidate; + ring_reserved = true; + break; + } + if (attempt + 1 == DOCA_XFER_REQ_MAX) { + NIXL_ERROR << "GPUNETIO transfer ring exhausted while sending notification"; + return NIXL_ERR_BACKEND; + } + } + if (!ring_reserved) { + return NIXL_ERR_BACKEND; + } - return NIXL_SUCCESS; + docaXferReqGpu request{}; + request.generation = xferReqRingCpu[pos].generation + 1; + request.state = DOCA_XFER_STATE_NOTIF_PENDING; + request.notif_state = DOCA_XFER_NOTIF_PENDING; + request.has_notif_msg_idx = buf_idx; + request.msg_sz = newMsg.size() + 1; + request.lbuf_notif = msg_buf; + request.lkey_notif = notif->send_mr->get_lkey(); + request.qp_notif = searchQp->second->qp_notif->get_qp_gpu_dev(); + request.qp_progress = searchQp->second->progress_gpu; + memcpy(&xferReqRingCpu[pos], &request, sizeof(request)); + std::atomic_ref(xferReqRingCpu[pos].state) + .store(DOCA_XFER_STATE_NOTIF_PENDING, std::memory_order_release); + const doca_error_t result = + doca_kernel_publish_notif(post_stream[xferStream.fetch_add(1) & (nstreams - 1)], + xferReqRingGpu, + progress_state_gpu, + pos); + if (result != DOCA_SUCCESS) { + xferReqReserved_[pos].store(false, std::memory_order_release); + return NIXL_ERR_BACKEND; + } + + while (true) { + const uint32_t request_state = + std::atomic_ref(xferReqRingCpu[pos].state).load(std::memory_order_acquire); + if (request_state == DOCA_XFER_STATE_COMPLETE) { + xferReqReserved_[pos].store(false, std::memory_order_release); + return NIXL_SUCCESS; + } + if (request_state == DOCA_XFER_STATE_ERROR || + std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != + 0 || + std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + xferReqReserved_[pos].store(false, std::memory_order_release); + return NIXL_ERR_BACKEND; + } + std::this_thread::yield(); + } } diff --git a/src/plugins/gpunetio/gpunetio_backend.h b/src/plugins/gpunetio/gpunetio_backend.h index 0081ffc1f0..53a2bd5074 100644 --- a/src/plugins/gpunetio/gpunetio_backend.h +++ b/src/plugins/gpunetio/gpunetio_backend.h @@ -156,21 +156,18 @@ class nixlDocaEngine : public nixlBackendEngine { struct sockaddr oob_saddr; struct sockaddr oob_netmask; std::thread pthr; - uint64_t *last_rsvd_flags; - uint64_t *last_posted_flags; cudaStream_t post_stream[DOCA_POST_STREAM_NUM]; cudaStream_t wait_stream; mutable std::atomic xferStream; - mutable std::atomic lastPostedReq; - mutable std::mutex postLock_; struct docaXferReqGpu *xferReqRingGpu; struct docaXferReqGpu *xferReqRingCpu; mutable std::atomic xferRingPos; mutable std::array xferReqReserved_; - struct docaXferCompletion *completion_list_gpu; - struct docaXferCompletion *completion_list_cpu; + struct docaProgressState *progress_state_gpu; + struct docaProgressState *progress_state_cpu; + std::vector qp_progress_gpu_; uint32_t *wait_exit_gpu; uint32_t *wait_exit_cpu; struct docaNotif *notif_fill_gpu; @@ -178,9 +175,6 @@ class nixlDocaEngine : public nixlBackendEngine { struct docaNotif *notif_progress_gpu; struct docaNotif *notif_progress_cpu; - struct docaNotif *notif_send_gpu; - struct docaNotif *notif_send_cpu; - // Map of agent name to saved nixlDocaConnection info std::unordered_map remoteConnMap; std::unordered_map qpMap; @@ -197,6 +191,7 @@ class nixlDocaEngine : public nixlBackendEngine { cudaStream_t stream; uint32_t devId; std::vector positions; + std::vector generations; uintptr_t backendHandleGpu; size_t postedCount = 0; nixl_status_t postStatus = NIXL_SUCCESS; @@ -220,8 +215,6 @@ class nixlDocaEngine : public nixlBackendEngine { connectClientRdmaQp(int oob_sock_client, const std::string &remote_agent); nixl_status_t nixlDocaDestroyNotif(doca_gpu *gpu, struct nixlDocaNotif *notif); - - mutable std::mutex notifSendLock; }; #endif diff --git a/src/plugins/gpunetio/gpunetio_backend_aux.h b/src/plugins/gpunetio/gpunetio_backend_aux.h index 0efd9cbb59..c46240ad72 100644 --- a/src/plugins/gpunetio/gpunetio_backend_aux.h +++ b/src/plugins/gpunetio/gpunetio_backend_aux.h @@ -54,15 +54,12 @@ // Local includes #include "common/nixl_time.h" -constexpr uint32_t DOCA_MAX_COMPLETION_INFLIGHT = 128; -constexpr uint32_t DOCA_MAX_COMPLETION_INFLIGHT_MASK = (DOCA_MAX_COMPLETION_INFLIGHT - 1); constexpr uint32_t RDMA_SEND_QUEUE_SIZE = 2048; constexpr uint32_t RDMA_RECV_QUEUE_SIZE = (RDMA_SEND_QUEUE_SIZE * 2); constexpr uint32_t DOCA_POST_STREAM_NUM = 4; constexpr uint32_t DOCA_XFER_REQ_SIZE = 512; constexpr uint32_t DOCA_XFER_REQ_MAX = 32; constexpr uint32_t DOCA_XFER_REQ_MASK = (DOCA_XFER_REQ_MAX - 1); -constexpr uint32_t DOCA_ENG_MAX_CONN = 20; constexpr uint32_t DOCA_RDMA_CM_LOCAL_PORT_SERVER = 6544; constexpr uint32_t VERBS_TEST_HOP_LIMIT = 255; @@ -79,25 +76,68 @@ constexpr uint32_t DOCA_NOTIF_NULL = 0xFFFFFFFF; #endif struct docaXferReqGpu { - uint32_t id; uintptr_t lbuf[DOCA_XFER_REQ_SIZE]; uintptr_t rbuf[DOCA_XFER_REQ_SIZE]; size_t size[DOCA_XFER_REQ_SIZE]; uint32_t lkey[DOCA_XFER_REQ_SIZE]; uint32_t rkey[DOCA_XFER_REQ_SIZE]; uint16_t num; - uint8_t in_use; uint32_t conn_idx; uint32_t has_notif_msg_idx; uint32_t msg_sz; uint64_t last_wqe; + uint64_t data_ticket; + uint64_t notif_wqe; + uint64_t notif_ticket; + uint32_t generation; + uint32_t state; + uint32_t data_state; + uint32_t notif_state; uintptr_t lbuf_notif; uint32_t lkey_notif; - uint64_t *last_rsvd; - uint64_t *last_posted; nixl_xfer_op_t backendOp; /* Needed only in case of GPU device transfer */ doca_gpu_dev_verbs_qp *qp_data; doca_gpu_dev_verbs_qp *qp_notif; + struct docaQpProgress *qp_progress; +}; + +enum docaXferState : uint32_t { + DOCA_XFER_STATE_FREE, + DOCA_XFER_STATE_PREPARED, + DOCA_XFER_STATE_DATA_POSTED, + DOCA_XFER_STATE_NOTIF_PENDING, + DOCA_XFER_STATE_NOTIF_POSTED, + DOCA_XFER_STATE_COMPLETE, + DOCA_XFER_STATE_ERROR, +}; + +enum docaXferDataState : uint32_t { + DOCA_XFER_DATA_NONE, + DOCA_XFER_DATA_POSTED, + DOCA_XFER_DATA_COMPLETE, +}; + +enum docaXferNotifState : uint32_t { + DOCA_XFER_NOTIF_NONE, + DOCA_XFER_NOTIF_PENDING, + DOCA_XFER_NOTIF_POSTED, + DOCA_XFER_NOTIF_COMPLETE, +}; + +struct docaQpProgress { + uint32_t data_producer_lock; + uint32_t notif_producer_lock; + uint64_t next_data_ticket; + uint64_t head_data_ticket; + uint64_t next_notif_ticket; + uint64_t head_notif_ticket; +}; + +struct docaProgressState { + uint32_t active_bitmap; + uint32_t progress_cursor; + uint32_t failed; + uint32_t active_generation[DOCA_XFER_REQ_MAX]; }; struct nixlDocaNotif { @@ -111,11 +151,6 @@ struct nixlDocaNotif { std::unique_ptr recv_mr; }; -struct docaXferCompletion { - uint8_t completed; - struct docaXferReqGpu *xferReqRingGpu; -}; - struct docaNotif { doca_gpu_dev_verbs_qp *qp_gpu; uint32_t msg_lkey; @@ -176,6 +211,7 @@ struct nixlDocaRdmaQp { uint32_t qpn_notif; uint32_t rqpn_notif; uint32_t remote_gid_notif; + struct docaQpProgress *progress_gpu; }; struct nixlDocaEngine; @@ -210,18 +246,27 @@ doca_error_t doca_kernel_write(cudaStream_t stream, doca_gpu_dev_verbs_qp *qp_gpu, struct docaXferReqGpu *xferReqRing, + struct docaProgressState *progress_state, + uint32_t *exit_flag, uint32_t pos); doca_error_t doca_kernel_read(cudaStream_t stream, doca_gpu_dev_verbs_qp *qp_gpu, struct docaXferReqGpu *xferReqRing, + struct docaProgressState *progress_state, + uint32_t *exit_flag, uint32_t pos); doca_error_t +doca_kernel_publish_notif(cudaStream_t stream, + struct docaXferReqGpu *xfer_req_ring, + struct docaProgressState *progress_state, + uint32_t pos); +doca_error_t doca_kernel_progress(cudaStream_t stream, - struct docaXferCompletion *completion_list, + struct docaXferReqGpu *xfer_req_ring, + struct docaProgressState *progress_state, struct docaNotif *notif_fill, struct docaNotif *notif_progress, - struct docaNotif *notif_send_gpu, uint32_t *exit_flag); #endif /* GPUNETIO_BACKEND_AUX_H */ diff --git a/src/plugins/gpunetio/gpunetio_kernels.cu b/src/plugins/gpunetio/gpunetio_kernels.cu index 7e97cc7ae4..8fd2e751ad 100644 --- a/src/plugins/gpunetio/gpunetio_kernels.cu +++ b/src/plugins/gpunetio/gpunetio_kernels.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -24,6 +24,18 @@ #define ENABLE_DEBUG 0 +__device__ uint32_t +nixl_gpunetio_dev_load_host_state(uint32_t &state) { + return cuda::atomic_ref(state).load( + cuda::std::memory_order_acquire); +} + +__device__ void +nixl_gpunetio_dev_store_host_state(uint32_t &state, uint32_t value) { + cuda::atomic_ref(state).store( + value, cuda::std::memory_order_release); +} + __device__ inline void nixl_gpunetio_dev_cq_print_cqe_err(struct mlx5_cqe64 *cqe64) { struct mlx5_err_cqe_ex *err_cqe = (struct mlx5_err_cqe_ex *)cqe64; @@ -61,10 +73,13 @@ nixl_gpunetio_dev_priv_poll_one_cq_at(doca_gpu_dev_verbs_cq *cq, uint64_t cons_i bool observed_completion = !((opown & MLX5_CQE_OWNER_MASK) ^ !!(cons_index & cqe_num)); observed_completion = observed_completion && (opcode != MLX5_CQE_INVALID); - if (!observed_completion) return EBUSY; + if (!observed_completion) { + return EBUSY; + } - if ((opcode == MLX5_CQE_REQ_ERR || opcode == MLX5_CQE_RESP_ERR) * -EIO) + if ((opcode == MLX5_CQE_REQ_ERR || opcode == MLX5_CQE_RESP_ERR) * -EIO) { nixl_gpunetio_dev_cq_print_cqe_err(cqe64); + } return ((opcode == MLX5_CQE_REQ_ERR || opcode == MLX5_CQE_RESP_ERR) * -EIO); } @@ -93,26 +108,109 @@ nixl_gpunetio_dev_poll_one_cq_at(doca_gpu_dev_verbs_cq *cq, uint64_t cons_index) return status; } +__device__ bool +nixl_gpunetio_dev_has_sq_credit(doca_gpu_dev_verbs_qp *qp, uint32_t count) { + auto *cq = doca_gpu_dev_verbs_qp_get_cq_sq(qp); + const uint64_t reserved = atomicAdd((unsigned long long *)&qp->sq_rsvd_index, 0); + const uint64_t completed = atomicAdd((unsigned long long *)&cq->cqe_ci, 0); + return reserved + count <= completed + __ldg(&qp->sq_wqe_num); +} + +__device__ void +nixl_gpunetio_dev_fail_request(docaXferReqGpu *request, + docaProgressState *progress_state, + uint32_t pos, + uint32_t *exit_flag) { + cuda::atomic_ref(progress_state->active_bitmap) + .fetch_and(~(1U << pos), cuda::std::memory_order_release); + nixl_gpunetio_dev_store_host_state(request->state, DOCA_XFER_STATE_ERROR); + nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); + nixl_gpunetio_dev_store_host_state(*exit_flag, 1U); +} + +__device__ void +nixl_gpunetio_dev_complete_request(docaXferReqGpu *request, + docaProgressState *progress_state, + uint32_t pos) { + cuda::atomic_ref(progress_state->active_bitmap) + .fetch_and(~(1U << pos), cuda::std::memory_order_release); + nixl_gpunetio_dev_store_host_state(request->state, DOCA_XFER_STATE_COMPLETE); +} + +__device__ bool +nixl_gpunetio_dev_reserve_data(docaXferReqGpu *request, + uint32_t count, + uint32_t *exit_flag, + uint64_t *base_wqe_idx) { + docaQpProgress *progress = request->qp_progress; + while (nixl_gpunetio_dev_load_host_state(*exit_flag) == 0U) { + if (atomicCAS(&progress->data_producer_lock, 0U, 1U) != 0U) { + continue; + } + if (!nixl_gpunetio_dev_has_sq_credit(request->qp_data, count)) { + atomicExch(&progress->data_producer_lock, 0U); + continue; + } + + *base_wqe_idx = + doca_gpu_dev_verbs_reserve_wq_slots(request->qp_data, count); + request->data_ticket = atomicAdd((unsigned long long *)&progress->next_data_ticket, 1ULL); + return true; + } + return false; +} + +__device__ void +nixl_gpunetio_dev_publish_data(docaXferReqGpu *request, + docaProgressState *progress_state, + uint32_t pos) { + atomicExch(&request->data_state, DOCA_XFER_DATA_POSTED); + nixl_gpunetio_dev_store_host_state(request->state, DOCA_XFER_STATE_DATA_POSTED); + progress_state->active_generation[pos] = request->generation; + cuda::atomic_ref(progress_state->active_bitmap) + .fetch_or(1U << pos, cuda::std::memory_order_release); +} + __global__ void -kernel_read(doca_gpu_dev_verbs_qp *qp, struct docaXferReqGpu *xferReqRing, uint32_t pos) { +kernel_read(doca_gpu_dev_verbs_qp *qp, + struct docaXferReqGpu *xferReqRing, + docaProgressState *progress_state, + uint32_t *exit_flag, + uint32_t pos) { uint64_t wqe_idx = 0; doca_gpu_dev_verbs_wqe *wqe_ptr; enum doca_gpu_dev_verbs_wqe_ctrl_flags cflag = DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE; uint32_t tot_wqe, idx = 0; __shared__ uint64_t base_wqe_idx; + __shared__ uint64_t last_wqe_idx; + __shared__ uint32_t reserved; // Warmup - if (xferReqRing == nullptr) return; + if (xferReqRing == nullptr) { + return; + } tot_wqe = xferReqRing[pos].num; if (threadIdx.x == 0) { - if (qp->need_mcst == true) - base_wqe_idx = doca_gpu_dev_verbs_reserve_wq_slots(qp, tot_wqe + 1); - else - base_wqe_idx = doca_gpu_dev_verbs_reserve_wq_slots(qp, tot_wqe); + if (nixl_gpunetio_dev_load_host_state(xferReqRing[pos].state) != DOCA_XFER_STATE_PREPARED) { + reserved = 0U; + nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos, exit_flag); + } else { + const uint32_t count = tot_wqe + (qp->need_mcst ? 1 : 0); + reserved = + nixl_gpunetio_dev_reserve_data(&xferReqRing[pos], count, exit_flag, &base_wqe_idx); + if (reserved == 0U) { + nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos, exit_flag); + } + } } __syncthreads(); + if (reserved == 0U) { + return; + } for (idx = threadIdx.x; idx < tot_wqe; idx += blockDim.x) { wqe_idx = base_wqe_idx + idx; @@ -130,54 +228,73 @@ kernel_read(doca_gpu_dev_verbs_qp *qp, struct docaXferReqGpu *xferReqRing, uint3 } __syncthreads(); - if ((idx - blockDim.x) == (tot_wqe - 1)) { + if (threadIdx.x == 0) { + last_wqe_idx = base_wqe_idx + tot_wqe - 1; if (qp->need_mcst == true) { - wqe_idx++; - wqe_ptr = doca_gpu_dev_verbs_get_wqe_ptr(qp, wqe_idx); + ++last_wqe_idx; + wqe_ptr = doca_gpu_dev_verbs_get_wqe_ptr(qp, last_wqe_idx); doca_gpu_dev_verbs_wqe_prepare_dump(qp, wqe_ptr, - wqe_idx, + last_wqe_idx, DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE, (uint64_t)(xferReqRing[pos].lbuf[tot_wqe - 1]), xferReqRing[pos].lkey[tot_wqe - 1], 1); } - doca_gpu_dev_verbs_mark_wqes_ready(qp, base_wqe_idx, wqe_idx); - doca_gpu_dev_verbs_submit(qp, wqe_idx + 1); - // Wait for final CQE in block of iterations - if (doca_gpu_dev_verbs_poll_cq_at(doca_gpu_dev_verbs_qp_get_cq_sq(qp), wqe_idx) != 0) - printf("kernel_read: Error CQE!\n"); - - DOCA_GPUNETIO_VOLATILE(xferReqRing[pos].last_wqe) = wqe_idx; - doca_gpu_dev_verbs_fence_release(); - DOCA_GPUNETIO_VOLATILE(xferReqRing[pos].in_use) = 1; + doca_gpu_dev_verbs_mark_wqes_ready(qp, base_wqe_idx, last_wqe_idx); + doca_gpu_dev_verbs_submit(qp, last_wqe_idx + 1); + xferReqRing[pos].last_wqe = last_wqe_idx; + atomicExch(&xferReqRing[pos].qp_progress->data_producer_lock, 0U); + nixl_gpunetio_dev_publish_data(&xferReqRing[pos], progress_state, pos); } #if ENABLE_DEBUG == 1 - if (threadIdx.x == 0) + if (threadIdx.x == 0) { printf(">>>>>>> CUDA rdma read kernel pos %d posted %d buffers from base_wqe_idx %ld\n", pos, xferReqRing[pos].num, base_wqe_idx); + } #endif } __global__ void -kernel_write(doca_gpu_dev_verbs_qp *qp, struct docaXferReqGpu *xferReqRing, uint32_t pos) { +kernel_write(doca_gpu_dev_verbs_qp *qp, + struct docaXferReqGpu *xferReqRing, + docaProgressState *progress_state, + uint32_t *exit_flag, + uint32_t pos) { uint64_t wqe_idx = 0; doca_gpu_dev_verbs_wqe *wqe_ptr; enum doca_gpu_dev_verbs_wqe_ctrl_flags cflag = DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE; uint32_t tot_wqe, idx = 0; __shared__ uint64_t base_wqe_idx; + __shared__ uint32_t reserved; // Warmup - if (xferReqRing == nullptr) return; + if (xferReqRing == nullptr) { + return; + } tot_wqe = xferReqRing[pos].num; - if (threadIdx.x == 0) base_wqe_idx = doca_gpu_dev_verbs_reserve_wq_slots(qp, tot_wqe); + if (threadIdx.x == 0) { + if (nixl_gpunetio_dev_load_host_state(xferReqRing[pos].state) != DOCA_XFER_STATE_PREPARED) { + reserved = 0U; + nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos, exit_flag); + } else { + reserved = nixl_gpunetio_dev_reserve_data( + &xferReqRing[pos], tot_wqe, exit_flag, &base_wqe_idx); + if (reserved == 0U) { + nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos, exit_flag); + } + } + } __syncthreads(); + if (reserved == 0U) { + return; + } for (idx = threadIdx.x; idx < tot_wqe; idx += blockDim.x) { wqe_idx = base_wqe_idx + idx; @@ -205,101 +322,164 @@ kernel_write(doca_gpu_dev_verbs_qp *qp, struct docaXferReqGpu *xferReqRing, uint } __syncthreads(); - if ((idx - blockDim.x) == (tot_wqe - 1)) { - doca_gpu_dev_verbs_mark_wqes_ready(qp, base_wqe_idx, wqe_idx); - doca_gpu_dev_verbs_submit(qp, wqe_idx + 1); - - DOCA_GPUNETIO_VOLATILE(xferReqRing[pos].last_wqe) = wqe_idx; - doca_gpu_dev_verbs_fence_release(); - DOCA_GPUNETIO_VOLATILE(xferReqRing[pos].in_use) = 1; + if (threadIdx.x == 0) { + const uint64_t last_wqe_idx = base_wqe_idx + tot_wqe - 1; + doca_gpu_dev_verbs_mark_wqes_ready(qp, base_wqe_idx, last_wqe_idx); + doca_gpu_dev_verbs_submit(qp, last_wqe_idx + 1); + xferReqRing[pos].last_wqe = last_wqe_idx; + atomicExch(&xferReqRing[pos].qp_progress->data_producer_lock, 0U); + nixl_gpunetio_dev_publish_data(&xferReqRing[pos], progress_state, pos); } #if ENABLE_DEBUG == 1 - if (threadIdx.x == 0) + if (threadIdx.x == 0) { printf(">>>>>>> CUDA rdma write kernel pos %d posted %d buffers from base_wqe_idx %ld\n", pos, xferReqRing[pos].num, base_wqe_idx); + } #endif } __global__ void -kernel_progress(struct docaXferCompletion *completion_list, +kernel_publish_notif(docaXferReqGpu *xfer_req_ring, + docaProgressState *progress_state, + uint32_t pos) { + if (nixl_gpunetio_dev_load_host_state(xfer_req_ring[pos].state) != + DOCA_XFER_STATE_NOTIF_PENDING) { + return; + } + progress_state->active_generation[pos] = xfer_req_ring[pos].generation; + cuda::atomic_ref(progress_state->active_bitmap) + .fetch_or(1U << pos, cuda::std::memory_order_release); +} + +__device__ bool +nixl_gpunetio_dev_try_post_notif(docaXferReqGpu *request, uint32_t *exit_flag) { + docaQpProgress *progress = request->qp_progress; + doca_gpu_dev_verbs_qp *qp = request->qp_notif; + + // Data-coupled and standalone notifications share this per-peer SQ owner. + if (atomicCAS(&progress->notif_producer_lock, 0U, 1U) != 0U) { + return false; + } + if (nixl_gpunetio_dev_load_host_state(*exit_flag) != 0U || + !nixl_gpunetio_dev_has_sq_credit(qp, 1)) { + atomicExch(&progress->notif_producer_lock, 0U); + return false; + } + + const uint64_t wqe_idx = + doca_gpu_dev_verbs_reserve_wq_slots(qp, 1); + doca_gpu_dev_verbs_wqe *wqe_ptr = doca_gpu_dev_verbs_get_wqe_ptr(qp, wqe_idx); + doca_gpu_dev_verbs_wqe_prepare_send(qp, + wqe_ptr, + wqe_idx, + DOCA_GPUNETIO_MLX5_OPCODE_SEND, + DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE, + 0, + request->lbuf_notif, + request->lkey_notif, + request->msg_sz); + doca_gpu_dev_verbs_mark_wqes_ready(qp, wqe_idx, wqe_idx); + doca_gpu_dev_verbs_submit(qp, wqe_idx + 1); + request->notif_wqe = wqe_idx; + request->notif_ticket = atomicAdd((unsigned long long *)&progress->next_notif_ticket, 1ULL); + atomicExch(&request->notif_state, DOCA_XFER_NOTIF_POSTED); + nixl_gpunetio_dev_store_host_state(request->state, DOCA_XFER_STATE_NOTIF_POSTED); + atomicExch(&progress->notif_producer_lock, 0U); + return true; +} + +__global__ void +kernel_progress(struct docaXferReqGpu *xfer_req_ring, + struct docaProgressState *progress_state, struct docaNotif *notif_fill, struct docaNotif *notif_progress, - struct docaNotif *notif_send_gpu, uint32_t *exit_flag) { - uint32_t index = 0; - doca_gpu_dev_verbs_ticket_t out_ticket; - - // Warmup - if (completion_list == nullptr) return; + if (xfer_req_ring == nullptr) { + return; + } - // Wait Xfer & notify + // One scheduler owns all data-SQ and notification-SQ completion credits. if (blockIdx.x == 0) { - while (DOCA_GPUNETIO_VOLATILE(*exit_flag) == 0) { - // Check xfer completion and send notif - if (DOCA_GPUNETIO_VOLATILE(completion_list[index].xferReqRingGpu) != nullptr) { - if (DOCA_GPUNETIO_VOLATILE(completion_list[index].completed) == 0 && - DOCA_GPUNETIO_VOLATILE(completion_list[index].xferReqRingGpu->in_use) == 1) { - // Wait for final CQE in block of iterations - int poll_status = nixl_gpunetio_dev_poll_one_cq_at< + while (nixl_gpunetio_dev_load_host_state(*exit_flag) == 0U) { + const uint32_t active = + cuda::atomic_ref(progress_state->active_bitmap) + .load(cuda::std::memory_order_acquire); + const uint32_t start = atomicAdd(&progress_state->progress_cursor, 0U); + + for (uint32_t offset = 0; offset < DOCA_XFER_REQ_MAX; ++offset) { + const uint32_t pos = (start + offset) & DOCA_XFER_REQ_MASK; + if ((active & (1U << pos)) == 0U) { + continue; + } + + docaXferReqGpu *request = &xfer_req_ring[pos]; + if (progress_state->active_generation[pos] != request->generation) { + nixl_gpunetio_dev_fail_request(request, progress_state, pos, exit_flag); + break; + } + docaQpProgress *progress = request->qp_progress; + const uint32_t request_state = nixl_gpunetio_dev_load_host_state(request->state); + + if (request_state == DOCA_XFER_STATE_DATA_POSTED && + request->data_ticket == + atomicAdd((unsigned long long *)&progress->head_data_ticket, 0ULL)) { + const int poll_status = nixl_gpunetio_dev_poll_one_cq_at< DOCA_GPUNETIO_VERBS_RESOURCE_SHARING_MODE_GPU, DOCA_GPUNETIO_VERBS_QP_SQ>( - doca_gpu_dev_verbs_qp_get_cq_sq( - completion_list[index].xferReqRingGpu->qp_data), - DOCA_GPUNETIO_VOLATILE(completion_list[index].xferReqRingGpu->last_wqe)); - if (poll_status == EBUSY) { - /* Not completed yet, continue progress loop to check exit_flag */ - continue; - } else if (poll_status != 0) { - DOCA_GPUNETIO_VOLATILE(*exit_flag) = 1; - printf( - "kernel_progress: block %d error CQE! poll_status %d wqe %ld index " - "%d\n", - blockIdx.x, - poll_status, - DOCA_GPUNETIO_VOLATILE(completion_list[index].xferReqRingGpu->last_wqe), - index); - break; - } else { - if (DOCA_GPUNETIO_VOLATILE( - completion_list[index].xferReqRingGpu->has_notif_msg_idx) != - DOCA_NOTIF_NULL) { -#if ENABLE_DEBUG == 1 - printf("Notif after completion at %d id %d sz %d\n", - index, - DOCA_GPUNETIO_VOLATILE( - completion_list[index].xferReqRingGpu->has_notif_msg_idx), - (int)completion_list[index].xferReqRingGpu->msg_sz); -#endif - doca_gpu_dev_verbs_send( - completion_list[index].xferReqRingGpu->qp_notif, - doca_gpu_dev_verbs_addr{ - .addr = (uint64_t)(completion_list[index] - .xferReqRingGpu->lbuf_notif), - .key = completion_list[index].xferReqRingGpu->lkey_notif}, - completion_list[index].xferReqRingGpu->msg_sz, - &out_ticket); - - doca_gpu_dev_verbs_wait(completion_list[index].xferReqRingGpu->qp_notif, - &out_ticket); -#if ENABLE_DEBUG == 1 - printf("Notif correctly sent %ld\n", out_ticket); -#endif + doca_gpu_dev_verbs_qp_get_cq_sq(request->qp_data), request->last_wqe); + if (poll_status == 0) { + atomicAdd((unsigned long long *)&progress->head_data_ticket, 1ULL); + atomicExch(&request->data_state, DOCA_XFER_DATA_COMPLETE); + if (request->has_notif_msg_idx == DOCA_NOTIF_NULL) { + nixl_gpunetio_dev_complete_request(request, progress_state, pos); + continue; + } else { + atomicExch(&request->notif_state, DOCA_XFER_NOTIF_PENDING); + nixl_gpunetio_dev_store_host_state(request->state, + DOCA_XFER_STATE_NOTIF_PENDING); } + } else if (poll_status != EBUSY) { + nixl_gpunetio_dev_fail_request(request, progress_state, pos, exit_flag); + break; + } + } - DOCA_GPUNETIO_VOLATILE(completion_list[index].completed) = 1; - index = (index + 1) & DOCA_MAX_COMPLETION_INFLIGHT_MASK; + if (nixl_gpunetio_dev_load_host_state(request->state) == + DOCA_XFER_STATE_NOTIF_PENDING) { + nixl_gpunetio_dev_try_post_notif(request, exit_flag); + } + + if (nixl_gpunetio_dev_load_host_state(request->state) == + DOCA_XFER_STATE_NOTIF_POSTED && + request->notif_ticket == + atomicAdd((unsigned long long *)&progress->head_notif_ticket, 0ULL)) { + const int poll_status = nixl_gpunetio_dev_poll_one_cq_at< + DOCA_GPUNETIO_VERBS_RESOURCE_SHARING_MODE_GPU, + DOCA_GPUNETIO_VERBS_QP_SQ>( + doca_gpu_dev_verbs_qp_get_cq_sq(request->qp_notif), request->notif_wqe); + if (poll_status == 0) { + atomicAdd((unsigned long long *)&progress->head_notif_ticket, 1ULL); + atomicExch(&request->notif_state, DOCA_XFER_NOTIF_COMPLETE); + nixl_gpunetio_dev_complete_request(request, progress_state, pos); + continue; + } else if (poll_status != EBUSY) { + nixl_gpunetio_dev_fail_request(request, progress_state, pos, exit_flag); + break; } } } + atomicExch(&progress_state->progress_cursor, (start + 1) & DOCA_XFER_REQ_MASK); } - } // Block 0 + } // Receive notif: fill recv in new queue and progress queue if (blockIdx.x == 1) { - while (DOCA_GPUNETIO_VOLATILE(*exit_flag) == 0) { + while (nixl_gpunetio_dev_load_host_state(*exit_flag) == 0U) { // Check received notifications if (DOCA_GPUNETIO_VOLATILE(notif_progress->qp_gpu) != nullptr) { uint32_t msg_last = DOCA_GPUNETIO_VOLATILE(notif_progress->msg_last); @@ -307,7 +487,7 @@ kernel_progress(struct docaXferCompletion *completion_list, nixl_gpunetio_dev_poll_one_cq_at( doca_gpu_dev_verbs_qp_get_cq_rq(notif_progress->qp_gpu), msg_last); - if (ret != EBUSY) { + if (ret == 0) { #if ENABLE_DEBUG == 1 printf("kernel received notification at %d ret %d\n", msg_last, ret); #endif @@ -330,13 +510,17 @@ kernel_progress(struct docaXferCompletion *completion_list, doca_gpu_dev_verbs_fence_release(); DOCA_GPUNETIO_VOLATILE(notif_progress->qp_gpu) = nullptr; - } else { + } else if (ret == EBUSY) { #if ENABLE_DEBUG == 1 printf("kernel received notification EBUSY at %d ret %d\n", msg_last, ret); #endif DOCA_GPUNETIO_VOLATILE(notif_progress->msg_num) = 0; doca_gpu_dev_verbs_fence_release(); DOCA_GPUNETIO_VOLATILE(notif_progress->qp_gpu) = nullptr; + } else { + nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); + nixl_gpunetio_dev_store_host_state(*exit_flag, 1U); + DOCA_GPUNETIO_VOLATILE(notif_progress->qp_gpu) = nullptr; } } @@ -362,37 +546,14 @@ kernel_progress(struct docaXferCompletion *completion_list, } } } - - // Send standalone notifications - if (blockIdx.x == 2) { - while (DOCA_GPUNETIO_VOLATILE(*exit_flag) == 0) { - if (DOCA_GPUNETIO_VOLATILE(notif_send_gpu->qp_gpu) != nullptr) { - doca_gpu_dev_verbs_send( - notif_send_gpu->qp_gpu, - doca_gpu_dev_verbs_addr{.addr = (uint64_t)notif_send_gpu->msg_buf, - .key = notif_send_gpu->msg_lkey}, - notif_send_gpu->msg_size, - &out_ticket); - - doca_gpu_dev_verbs_wait(notif_send_gpu->qp_gpu, &out_ticket); -#if ENABLE_DEBUG == 1 - printf("Notif correctly sent %ld addr %lx msg_lkey %x qp %p size %d\n", - out_ticket, - notif_send_gpu->msg_buf, - notif_send_gpu->msg_lkey, - (void *)notif_send_gpu->qp_gpu, - (int)notif_send_gpu->msg_size); -#endif - DOCA_GPUNETIO_VOLATILE(notif_send_gpu->qp_gpu) = nullptr; - } - } - } } doca_error_t doca_kernel_write(cudaStream_t stream, doca_gpu_dev_verbs_qp *qp, struct docaXferReqGpu *xferReqRing, + struct docaProgressState *progress_state, + uint32_t *exit_flag, uint32_t pos) { cudaError_t result = cudaSuccess; @@ -404,7 +565,8 @@ doca_kernel_write(cudaStream_t stream, return DOCA_ERROR_BAD_STATE; } - kernel_write<<<1, DOCA_XFER_REQ_SIZE, 0, stream>>>(qp, xferReqRing, pos); + kernel_write<<<1, DOCA_XFER_REQ_SIZE, 0, stream>>>( + qp, xferReqRing, progress_state, exit_flag, pos); result = cudaGetLastError(); if (result != cudaSuccess) { fprintf( @@ -419,6 +581,8 @@ doca_error_t doca_kernel_read(cudaStream_t stream, doca_gpu_dev_verbs_qp *qp, struct docaXferReqGpu *xferReqRing, + struct docaProgressState *progress_state, + uint32_t *exit_flag, uint32_t pos) { cudaError_t result = cudaSuccess; @@ -430,7 +594,31 @@ doca_kernel_read(cudaStream_t stream, return DOCA_ERROR_BAD_STATE; } - kernel_read<<<1, DOCA_XFER_REQ_SIZE, 0, stream>>>(qp, xferReqRing, pos); + kernel_read<<<1, DOCA_XFER_REQ_SIZE, 0, stream>>>( + qp, xferReqRing, progress_state, exit_flag, pos); + result = cudaGetLastError(); + if (result != cudaSuccess) { + fprintf( + stderr, "[%s:%d] cuda failed with %s", __FILE__, __LINE__, cudaGetErrorString(result)); + return DOCA_ERROR_BAD_STATE; + } + + return DOCA_SUCCESS; +} + +doca_error_t +doca_kernel_publish_notif(cudaStream_t stream, + struct docaXferReqGpu *xfer_req_ring, + struct docaProgressState *progress_state, + uint32_t pos) { + cudaError_t result = cudaGetLastError(); + if (result != cudaSuccess) { + fprintf( + stderr, "[%s:%d] cuda failed with %s", __FILE__, __LINE__, cudaGetErrorString(result)); + return DOCA_ERROR_BAD_STATE; + } + + kernel_publish_notif<<<1, 1, 0, stream>>>(xfer_req_ring, progress_state, pos); result = cudaGetLastError(); if (result != cudaSuccess) { fprintf( @@ -443,10 +631,10 @@ doca_kernel_read(cudaStream_t stream, doca_error_t doca_kernel_progress(cudaStream_t stream, - struct docaXferCompletion *completion_list, + struct docaXferReqGpu *xfer_req_ring, + struct docaProgressState *progress_state, struct docaNotif *notif_fill, struct docaNotif *notif_progress, - struct docaNotif *notif_send_gpu, uint32_t *exit_flag) { cudaError_t result = cudaSuccess; @@ -458,8 +646,8 @@ doca_kernel_progress(cudaStream_t stream, return DOCA_ERROR_BAD_STATE; } - kernel_progress<<<3, 1, 0, stream>>>( - completion_list, notif_fill, notif_progress, notif_send_gpu, exit_flag); + kernel_progress<<<2, 1, 0, stream>>>( + xfer_req_ring, progress_state, notif_fill, notif_progress, exit_flag); result = cudaGetLastError(); if (result != cudaSuccess) { fprintf( From 467cada8c75a9529e6b85f56cf8394bb2ceb74f1 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Mon, 7 Sep 2026 19:39:35 +0800 Subject: [PATCH 10/25] fix(gpunetio): rearm requests and drain owned work after errors Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 94 ++++++++++++++++------- src/plugins/gpunetio/gpunetio_backend.h | 1 - src/plugins/gpunetio/gpunetio_kernels.cu | 66 ++++++++++------ 3 files changed, 111 insertions(+), 50 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index 8d95fb3866..4e4c8cec2c 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -645,7 +645,7 @@ nixlDocaEngine::addRdmaQp(const std::string &remote_agent) { if (init_stream != nullptr) { cudaStreamDestroy(init_stream); } - if (cuda_device_switched && previous_cuda_device != gdevs[0].first) { + if (cuda_device_switched && previous_cuda_device != static_cast(gdevs[0].first)) { cudaSetDevice(previous_cuda_device); } if (cuda_result != cudaSuccess) { @@ -1326,10 +1326,44 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, if (operation != NIXL_READ && operation != NIXL_WRITE) { return NIXL_ERR_INVALID_PARAM; } + if (std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != + 0 || + std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + return NIXL_ERR_BACKEND; + } if (treq->postStatus != NIXL_SUCCESS) { return treq->postStatus; } + const auto completion_state = treq->completionState.load(std::memory_order_acquire); + if (completion_state == nixlDocaBckndReq::completion_state::COMPLETING || + (completion_state == nixlDocaBckndReq::completion_state::IN_PROGRESS && + treq->postedCount != 0)) { + return NIXL_IN_PROG; + } + if (completion_state == nixlDocaBckndReq::completion_state::COMPLETE) { + for (size_t i = 0; i < treq->positions.size(); ++i) { + const uint32_t idx = treq->positions[i]; + if (std::atomic_ref(xferReqRingCpu[idx].state) + .load(std::memory_order_acquire) != DOCA_XFER_STATE_COMPLETE) { + treq->postStatus = NIXL_ERR_BACKEND; + return NIXL_ERR_BACKEND; + } + + auto &request = xferReqRingCpu[idx]; + request.generation++; + request.last_wqe = 0; + request.data_ticket = 0; + request.notif_wqe = 0; + request.notif_ticket = 0; + request.data_state = DOCA_XFER_DATA_NONE; + request.notif_state = DOCA_XFER_NOTIF_NONE; + treq->generations[i] = request.generation; + std::atomic_ref(request.state) + .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); + } + } + treq->postedCount = 0; treq->postStatus = NIXL_SUCCESS; treq->completionState.store(nixlDocaBckndReq::completion_state::IN_PROGRESS, @@ -1350,51 +1384,61 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, idx); if (result != DOCA_SUCCESS) { treq->postStatus = NIXL_ERR_BACKEND; + std::atomic_ref(progress_state_cpu->failed) + .store(1, std::memory_order_release); + for (size_t pending = treq->postedCount; pending < treq->positions.size(); ++pending) { + std::atomic_ref(xferReqRingCpu[treq->positions[pending]].state) + .store(DOCA_XFER_STATE_ERROR, std::memory_order_release); + } break; } ++treq->postedCount; } - return treq->postedCount == 0 ? treq->postStatus : NIXL_IN_PROG; + return treq->postStatus == NIXL_SUCCESS ? NIXL_IN_PROG : treq->postStatus; } nixl_status_t nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; - if (treq->postStatus != NIXL_SUCCESS) { - return treq->postStatus; - } - if (std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != - 0) { - treq->postStatus = NIXL_ERR_BACKEND; - return NIXL_ERR_BACKEND; - } auto state = treq->completionState.load(std::memory_order_acquire); if (state == nixlDocaBckndReq::completion_state::COMPLETE) { - return treq->postStatus; + return std::atomic_ref(progress_state_cpu->failed) + .load(std::memory_order_acquire) == 0 && + treq->postStatus == NIXL_SUCCESS ? + NIXL_SUCCESS : + NIXL_ERR_BACKEND; } if (state == nixlDocaBckndReq::completion_state::COMPLETING) { treq->completionState.wait(state, std::memory_order_acquire); - return treq->postStatus; + return checkXfer(handle); } - for (size_t i = 0; i < treq->postedCount; ++i) { + bool request_error = false; + for (size_t i = 0; i < treq->positions.size(); ++i) { const uint32_t idx = treq->positions[i]; const uint32_t req_state = std::atomic_ref(xferReqRingCpu[idx].state).load(std::memory_order_acquire); if (xferReqRingCpu[idx].generation != treq->generations[i]) { - treq->postStatus = NIXL_ERR_BACKEND; - return NIXL_ERR_BACKEND; + request_error = true; + continue; } if (req_state == DOCA_XFER_STATE_ERROR) { - treq->postStatus = NIXL_ERR_BACKEND; - return NIXL_ERR_BACKEND; + request_error = true; + continue; } if (req_state != DOCA_XFER_STATE_COMPLETE) { return NIXL_IN_PROG; } } + if (request_error || treq->postStatus != NIXL_SUCCESS || + std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != + 0) { + treq->postStatus = NIXL_ERR_BACKEND; + return NIXL_ERR_BACKEND; + } + retireRequest(treq); return treq->postStatus; } @@ -1429,12 +1473,11 @@ nixl_status_t nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { auto *treq = static_cast(handle); if (treq->postedCount == 0) { - const nixl_status_t status = treq->postStatus; for (uint32_t idx : treq->positions) { xferReqReserved_[idx].store(false, std::memory_order_release); } delete treq; - return status; + return NIXL_SUCCESS; } if (treq->completionState.load(std::memory_order_acquire) != @@ -1443,9 +1486,6 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { if (status == NIXL_IN_PROG) { return NIXL_IN_PROG; } - if (status != NIXL_SUCCESS) { - return status; - } } for (uint32_t idx : treq->positions) { @@ -1602,6 +1642,8 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg progress_state_gpu, pos); if (result != DOCA_SUCCESS) { + std::atomic_ref(xferReqRingCpu[pos].state) + .store(DOCA_XFER_STATE_ERROR, std::memory_order_release); xferReqReserved_[pos].store(false, std::memory_order_release); return NIXL_ERR_BACKEND; } @@ -1613,13 +1655,13 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg xferReqReserved_[pos].store(false, std::memory_order_release); return NIXL_SUCCESS; } - if (request_state == DOCA_XFER_STATE_ERROR || - std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != - 0 || - std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + if (request_state == DOCA_XFER_STATE_ERROR) { xferReqReserved_[pos].store(false, std::memory_order_release); return NIXL_ERR_BACKEND; } + if (std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + return NIXL_ERR_BACKEND; + } std::this_thread::yield(); } } diff --git a/src/plugins/gpunetio/gpunetio_backend.h b/src/plugins/gpunetio/gpunetio_backend.h index 53a2bd5074..f95a58a3fe 100644 --- a/src/plugins/gpunetio/gpunetio_backend.h +++ b/src/plugins/gpunetio/gpunetio_backend.h @@ -204,7 +204,6 @@ class nixlDocaEngine : public nixlBackendEngine { void retireRequest(nixlDocaBckndReq *request) const; - nixl_status_t progressThreadStart(); void diff --git a/src/plugins/gpunetio/gpunetio_kernels.cu b/src/plugins/gpunetio/gpunetio_kernels.cu index 8fd2e751ad..f485e8c4ae 100644 --- a/src/plugins/gpunetio/gpunetio_kernels.cu +++ b/src/plugins/gpunetio/gpunetio_kernels.cu @@ -101,7 +101,7 @@ __device__ int nixl_gpunetio_dev_poll_one_cq_at(doca_gpu_dev_verbs_cq *cq, uint64_t cons_index) { int status = nixl_gpunetio_dev_priv_poll_one_cq_at(cq, cons_index); - if (status == 0) { + if (status != EBUSY) { doca_gpu_dev_verbs_fence_acquire(); doca_gpu_dev_verbs_atomic_max(&cq->cqe_ci, cons_index + 1); } @@ -117,15 +117,20 @@ nixl_gpunetio_dev_has_sq_credit(doca_gpu_dev_verbs_qp *qp, uint32_t count) { } __device__ void -nixl_gpunetio_dev_fail_request(docaXferReqGpu *request, - docaProgressState *progress_state, - uint32_t pos, - uint32_t *exit_flag) { +nixl_gpunetio_dev_terminal_error(docaXferReqGpu *request, + docaProgressState *progress_state, + uint32_t pos) { cuda::atomic_ref(progress_state->active_bitmap) .fetch_and(~(1U << pos), cuda::std::memory_order_release); nixl_gpunetio_dev_store_host_state(request->state, DOCA_XFER_STATE_ERROR); +} + +__device__ void +nixl_gpunetio_dev_fail_request(docaXferReqGpu *request, + docaProgressState *progress_state, + uint32_t pos) { + nixl_gpunetio_dev_terminal_error(request, progress_state, pos); nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); - nixl_gpunetio_dev_store_host_state(*exit_flag, 1U); } __device__ void @@ -140,13 +145,19 @@ nixl_gpunetio_dev_complete_request(docaXferReqGpu *request, __device__ bool nixl_gpunetio_dev_reserve_data(docaXferReqGpu *request, uint32_t count, + docaProgressState *progress_state, uint32_t *exit_flag, uint64_t *base_wqe_idx) { docaQpProgress *progress = request->qp_progress; - while (nixl_gpunetio_dev_load_host_state(*exit_flag) == 0U) { + while (nixl_gpunetio_dev_load_host_state(*exit_flag) == 0U && + nixl_gpunetio_dev_load_host_state(progress_state->failed) == 0U) { if (atomicCAS(&progress->data_producer_lock, 0U, 1U) != 0U) { continue; } + if (nixl_gpunetio_dev_load_host_state(progress_state->failed) != 0U) { + atomicExch(&progress->data_producer_lock, 0U); + return false; + } if (!nixl_gpunetio_dev_has_sq_credit(request->qp_data, count)) { atomicExch(&progress->data_producer_lock, 0U); continue; @@ -197,13 +208,13 @@ kernel_read(doca_gpu_dev_verbs_qp *qp, if (threadIdx.x == 0) { if (nixl_gpunetio_dev_load_host_state(xferReqRing[pos].state) != DOCA_XFER_STATE_PREPARED) { reserved = 0U; - nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos, exit_flag); + nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos); } else { const uint32_t count = tot_wqe + (qp->need_mcst ? 1 : 0); - reserved = - nixl_gpunetio_dev_reserve_data(&xferReqRing[pos], count, exit_flag, &base_wqe_idx); + reserved = nixl_gpunetio_dev_reserve_data( + &xferReqRing[pos], count, progress_state, exit_flag, &base_wqe_idx); if (reserved == 0U) { - nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos, exit_flag); + nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos); } } } @@ -282,12 +293,12 @@ kernel_write(doca_gpu_dev_verbs_qp *qp, if (threadIdx.x == 0) { if (nixl_gpunetio_dev_load_host_state(xferReqRing[pos].state) != DOCA_XFER_STATE_PREPARED) { reserved = 0U; - nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos, exit_flag); + nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos); } else { reserved = nixl_gpunetio_dev_reserve_data( - &xferReqRing[pos], tot_wqe, exit_flag, &base_wqe_idx); + &xferReqRing[pos], tot_wqe, progress_state, exit_flag, &base_wqe_idx); if (reserved == 0U) { - nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos, exit_flag); + nixl_gpunetio_dev_fail_request(&xferReqRing[pos], progress_state, pos); } } } @@ -355,7 +366,9 @@ kernel_publish_notif(docaXferReqGpu *xfer_req_ring, } __device__ bool -nixl_gpunetio_dev_try_post_notif(docaXferReqGpu *request, uint32_t *exit_flag) { +nixl_gpunetio_dev_try_post_notif(docaXferReqGpu *request, + docaProgressState *progress_state, + uint32_t *exit_flag) { docaQpProgress *progress = request->qp_progress; doca_gpu_dev_verbs_qp *qp = request->qp_notif; @@ -364,6 +377,7 @@ nixl_gpunetio_dev_try_post_notif(docaXferReqGpu *request, uint32_t *exit_flag) { return false; } if (nixl_gpunetio_dev_load_host_state(*exit_flag) != 0U || + nixl_gpunetio_dev_load_host_state(progress_state->failed) != 0U || !nixl_gpunetio_dev_has_sq_credit(qp, 1)) { atomicExch(&progress->notif_producer_lock, 0U); return false; @@ -419,8 +433,9 @@ kernel_progress(struct docaXferReqGpu *xfer_req_ring, docaXferReqGpu *request = &xfer_req_ring[pos]; if (progress_state->active_generation[pos] != request->generation) { - nixl_gpunetio_dev_fail_request(request, progress_state, pos, exit_flag); - break; + // Do not release a possibly newer owner on a stale generation. + nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); + continue; } docaQpProgress *progress = request->qp_progress; const uint32_t request_state = nixl_gpunetio_dev_load_host_state(request->state); @@ -444,14 +459,19 @@ kernel_progress(struct docaXferReqGpu *xfer_req_ring, DOCA_XFER_STATE_NOTIF_PENDING); } } else if (poll_status != EBUSY) { - nixl_gpunetio_dev_fail_request(request, progress_state, pos, exit_flag); - break; + atomicAdd((unsigned long long *)&progress->head_data_ticket, 1ULL); + nixl_gpunetio_dev_fail_request(request, progress_state, pos); + continue; } } if (nixl_gpunetio_dev_load_host_state(request->state) == DOCA_XFER_STATE_NOTIF_PENDING) { - nixl_gpunetio_dev_try_post_notif(request, exit_flag); + if (nixl_gpunetio_dev_load_host_state(progress_state->failed) != 0U) { + nixl_gpunetio_dev_terminal_error(request, progress_state, pos); + continue; + } + nixl_gpunetio_dev_try_post_notif(request, progress_state, exit_flag); } if (nixl_gpunetio_dev_load_host_state(request->state) == @@ -468,8 +488,9 @@ kernel_progress(struct docaXferReqGpu *xfer_req_ring, nixl_gpunetio_dev_complete_request(request, progress_state, pos); continue; } else if (poll_status != EBUSY) { - nixl_gpunetio_dev_fail_request(request, progress_state, pos, exit_flag); - break; + atomicAdd((unsigned long long *)&progress->head_notif_ticket, 1ULL); + nixl_gpunetio_dev_fail_request(request, progress_state, pos); + continue; } } } @@ -519,7 +540,6 @@ kernel_progress(struct docaXferReqGpu *xfer_req_ring, DOCA_GPUNETIO_VOLATILE(notif_progress->qp_gpu) = nullptr; } else { nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); - nixl_gpunetio_dev_store_host_state(*exit_flag, 1U); DOCA_GPUNETIO_VOLATILE(notif_progress->qp_gpu) = nullptr; } } From d8a53f5021704bf90d2740348f4078783ed5cef7 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Mon, 7 Sep 2026 19:39:35 +0800 Subject: [PATCH 11/25] test(gpunetio): exercise progress ordering and publish paired benchmarks Signed-off-by: Yunxiao Ning --- test/gtest/plugins/gpunetio/README.md | 69 + test/gtest/plugins/gpunetio/meson.build | 20 + .../plugins/gpunetio/qp_progress_gtest.cu | 1240 +++++++++++++++++ test/gtest/plugins/meson.build | 4 + 4 files changed, 1333 insertions(+) create mode 100644 test/gtest/plugins/gpunetio/README.md create mode 100644 test/gtest/plugins/gpunetio/meson.build create mode 100644 test/gtest/plugins/gpunetio/qp_progress_gtest.cu diff --git a/test/gtest/plugins/gpunetio/README.md b/test/gtest/plugins/gpunetio/README.md new file mode 100644 index 0000000000..a07cd7361e --- /dev/null +++ b/test/gtest/plugins/gpunetio/README.md @@ -0,0 +1,69 @@ +# GPUNETIO QP-progress GoogleTest harness + +`gpunetio_qp_progress_gtest` is a paired, environment-configured integration +executable. It is intentionally not part of the default test suite: it needs +one source GPU and two target GPUs, a numeric target IPv4 address, unique OOB +ports, and a shared directory used only to exchange serialized NIXL metadata. +Control/verification acknowledgements use TCP, outside the measured interval. + +The GPUNETIO GTest parent wires this directory. The local `meson.build` builds +the paired executable and registers it with CTest; absent role configuration is +an explicit GoogleTest skip before any CUDA query. + +## Required environment + +Both processes need `NIXL_PLUGIN_DIR` and `LD_LIBRARY_PATH` pointing to the +same feature-tree build. Set the following values without placing site-specific +values in this repository: + +| Variable | Source | Target | Meaning | +| --- | --- | --- | --- | +| `NIXL_QP_PROGRESS_ROLE` | `source` | `target` or `target-fault` | Process role | +| `NIXL_QP_PROGRESS_COORD_DIR` | same path | same path | Fresh shared metadata directory | +| `NIXL_QP_PROGRESS_TARGET_IPV4` | required | unset | Numeric IPv4 of the target process | +| `NIXL_QP_PROGRESS_CONTROL_PORT` | same | same | TCP control port | +| `NIXL_QP_PROGRESS_SOURCE_OOB_PORT` | same | same | Source agent listener port | +| `NIXL_QP_PROGRESS_TARGET_A_OOB_PORT` | same | same | Target GPU 0 listener port | +| `NIXL_QP_PROGRESS_TARGET_B_OOB_PORT` | same | same | Target GPU 1 listener port | +| `NIXL_QP_PROGRESS_NETWORK_DEVICE` | optional | optional | GPUNETIO network device | +| `NIXL_QP_PROGRESS_OOB_INTERFACE` | optional | optional | OOB interface | +| `NIXL_QP_PROGRESS_GID_INDEX` | optional | optional | GID index | +| `NIXL_QP_PROGRESS_CONTROL_BYTES` | optional | optional | Single-active-QP control size; defaults to 4096 | + +Run target first, then source with the same `--gtest_filter`. Use a fresh +process for each fault trial (`target-fault` plus `source` filtered to the +fault case). The harness skips if its role, coordinate directory, GPU count, +or source target IPv4 is unavailable. + +```bash +# Terminal/host with two visible target GPUs. +NIXL_QP_PROGRESS_ROLE=target \ + gpunetio_qp_progress_gtest --gtest_filter=QpProgress.PerformanceMixed2MiBAnd4KiB + +# Terminal/host with one visible source GPU. Use the target's numeric IPv4. +NIXL_QP_PROGRESS_ROLE=source NIXL_QP_PROGRESS_TARGET_IPV4= \ + gpunetio_qp_progress_gtest --gtest_filter=QpProgress.PerformanceMixed2MiBAnd4KiB +``` + +The performance case runs one 20-warmup/100-measured pair. Run three fresh +source/target process pairs externally for independent repeats. +It writes `qp_progress_performance.json` into the coordinate directory. The +JSON reports p50/p99 transfer-window latency, actual payload and marker bytes, +and separate wall time. It does not claim CQ timing or backend-internal timing. + +`OutstandingDescriptorsAndAttachedStreamOrder` uses 513 descriptors per +request (512 data plus an epoch marker) with descriptor merging disabled. It +holds one delayed source-peer request while completing and releasing 129 +fast-peer requests, forcing live 32-slot ring reuse without exceeding capacity. +The delay is a test-only CUDA kernel enqueued before the attached transfer; it +is not a backend option or a production hook. The target verifies every fast +payload and the delayed payload after a system-acquire of each exact epoch +marker, then checks data-coupled notifications plus one standalone notification +per peer. `SingleQpReadWriteControl` provides the paired one-QP WRITE and READ +control performance path. It re-posts one prepared WRITE and one prepared READ +handle across changing payload epochs, with one active QP at a time, and emits +source API p50/p99 windows separately from the mixed WRITE performance JSON. +Set `NIXL_QP_PROGRESS_CONTROL_BYTES=2097152` to repeat the same control at 2 MiB. +`RemoteDeregisterReturnsBackendError` requires its fresh fault role and checks +bounded `NIXL_ERR_BACKEND` and release after target B memory is deregistered +after source metadata consumption and before its post. diff --git a/test/gtest/plugins/gpunetio/meson.build b/test/gtest/plugins/gpunetio/meson.build new file mode 100644 index 0000000000..fbb7f5d5e4 --- /dev/null +++ b/test/gtest/plugins/gpunetio/meson.build @@ -0,0 +1,20 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +if not cuda_dep.found() + message('CUDA not found, skipping gpunetio_qp_progress_gtest build') + subdir_done() +endif + +gpunetio_qp_progress_gtest = executable( + 'gpunetio_qp_progress_gtest', + 'qp_progress_gtest.cu', + include_directories: [nixl_inc_dirs, utils_inc_dirs, gtest_inc_dirs], + dependencies: [nixl_dep, nixl_infra, nixl_common_dep, cuda_dep, gtest_dep, gmock_dep], + link_with: [nixl_build_lib], + install: true, +) + +# The executable is always registered. Without the paired-role environment its +# GoogleTests explicitly SKIP, so CTest exposes rather than silently omits it. +test('gpunetio_qp_progress_gtest', gpunetio_qp_progress_gtest, is_parallel: false, timeout: 180) diff --git a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu new file mode 100644 index 0000000000..4aac639b8a --- /dev/null +++ b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu @@ -0,0 +1,1240 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +// Paired GPUNETIO integration harness. It observes only public NIXL/CUDA results. +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { +namespace fs = std::filesystem; +using Clock = std::chrono::steady_clock; + +constexpr size_t kLargeBytes = 2ULL << 20; +constexpr size_t kSmallBytes = 4ULL << 10; +constexpr size_t kStressDescriptors = 512; +constexpr size_t kStressRequests = 7; +constexpr size_t kFastChurnRequests = 129; +constexpr int kWarmup = 20; +constexpr int kMeasured = 100; +constexpr int kRepeats = 1; +constexpr auto kTimeout = std::chrono::seconds(30); + +struct Config { + std::string role; + fs::path coord; + std::string target_ipv4; + std::string network_device; + std::string oob_interface; + std::string gid_index; + int control_port = 0; + int source_port = 0; + int target_a_port = 0; + int target_b_port = 0; + size_t control_bytes = kSmallBytes; +}; + +[[noreturn]] void +Fail(const std::string &message) { + throw std::runtime_error(message); +} + +void +CheckCuda(cudaError_t status, const char *what) { + if (status != cudaSuccess) { + Fail(std::string(what) + ": " + cudaGetErrorString(status)); + } +} + +void +CheckNixl(nixl_status_t status, const char *what) { + if (status != NIXL_SUCCESS) { + Fail(std::string(what) + ": status=" + std::to_string(static_cast(status))); + } +} + +uint64_t +NowNs() { + return std::chrono::duration_cast(Clock::now().time_since_epoch()) + .count(); +} + +std::optional +Env(const char *key) { + const char *value = std::getenv(key); + return value == nullptr || *value == '\0' ? std::nullopt : std::optional(value); +} + +int +EnvPort(const char *key) { + const auto value = Env(key); + if (!value) { + return 0; + } + try { + const int port = std::stoi(*value); + return port > 0 && port <= 65535 ? port : 0; + } + catch (...) { + return 0; + } +} + +size_t +EnvBytes(const char *key, size_t fallback) { + const auto value = Env(key); + if (!value) { + return fallback; + } + try { + const size_t bytes = std::stoull(*value); + return bytes == 0 ? fallback : bytes; + } + catch (...) { + return fallback; + } +} + +Config +GetConfig() { + Config config; + config.role = Env("NIXL_QP_PROGRESS_ROLE").value_or(""); + config.coord = Env("NIXL_QP_PROGRESS_COORD_DIR").value_or(""); + config.target_ipv4 = Env("NIXL_QP_PROGRESS_TARGET_IPV4").value_or(""); + config.network_device = Env("NIXL_QP_PROGRESS_NETWORK_DEVICE").value_or(""); + config.oob_interface = Env("NIXL_QP_PROGRESS_OOB_INTERFACE").value_or(""); + config.gid_index = Env("NIXL_QP_PROGRESS_GID_INDEX").value_or(""); + config.control_port = EnvPort("NIXL_QP_PROGRESS_CONTROL_PORT"); + config.source_port = EnvPort("NIXL_QP_PROGRESS_SOURCE_OOB_PORT"); + config.target_a_port = EnvPort("NIXL_QP_PROGRESS_TARGET_A_OOB_PORT"); + config.target_b_port = EnvPort("NIXL_QP_PROGRESS_TARGET_B_OOB_PORT"); + config.control_bytes = EnvBytes("NIXL_QP_PROGRESS_CONTROL_BYTES", kSmallBytes); + return config; +} + +bool +IsTarget(const Config &config) { + return config.role == "target" || config.role == "target-fault"; +} + +std::optional +ConfigProblem(const Config &config, bool fault) { + if (config.role.empty() || config.coord.empty() || config.control_port == 0 || + config.source_port == 0 || config.target_a_port == 0 || config.target_b_port == 0) { + return "set NIXL_QP_PROGRESS_ROLE, COORD_DIR, and all paired ports"; + } + if (config.role != "source" && !IsTarget(config)) { + return "NIXL_QP_PROGRESS_ROLE must be source, target, or target-fault"; + } + if (fault != (config.role == "target-fault") && IsTarget(config)) { + return "role does not match this fault/non-fault case"; + } + if (config.role == "source" && config.target_ipv4.empty()) { + return "source requires numeric NIXL_QP_PROGRESS_TARGET_IPV4"; + } + int devices = 0; + const cudaError_t cuda_status = cudaGetDeviceCount(&devices); + if (cuda_status != cudaSuccess) { + return "CUDA runtime/device unavailable"; + } + if ((config.role == "source" && devices < 1) || (IsTarget(config) && devices < 2)) { + return "paired harness needs one source GPU or two target GPUs"; + } + return std::nullopt; +} + +class Control { +public: + explicit Control(const Config &config) { + fd_ = socket(AF_INET, SOCK_STREAM, 0); + if (fd_ < 0) { + Fail("control socket"); + } + sockaddr_in address{}; + address.sin_family = AF_INET; + address.sin_port = htons(config.control_port); + if (IsTarget(config)) { + int one = 1; + setsockopt(fd_, SOL_SOCKET, SO_REUSEADDR, &one, sizeof(one)); + address.sin_addr.s_addr = INADDR_ANY; + if (bind(fd_, reinterpret_cast(&address), sizeof(address)) != 0 || + listen(fd_, 1) != 0) { + Fail("control bind/listen"); + } + const int accepted = accept(fd_, nullptr, nullptr); + close(fd_); + fd_ = accepted; + if (fd_ < 0) { + Fail("control accept"); + } + } else { + if (inet_pton(AF_INET, config.target_ipv4.c_str(), &address.sin_addr) != 1) { + Fail("NIXL_QP_PROGRESS_TARGET_IPV4 must be numeric IPv4"); + } + const auto deadline = Clock::now() + kTimeout; + while (connect(fd_, reinterpret_cast(&address), sizeof(address)) != 0) { + if (Clock::now() >= deadline) { + Fail("control connect timeout"); + } + close(fd_); + fd_ = socket(AF_INET, SOCK_STREAM, 0); + if (fd_ < 0) { + Fail("control retry socket"); + } + std::this_thread::sleep_for(std::chrono::milliseconds(50)); + } + } + timeval timeout{static_cast(kTimeout.count()), 0}; + setsockopt(fd_, SOL_SOCKET, SO_RCVTIMEO, &timeout, sizeof(timeout)); + setsockopt(fd_, SOL_SOCKET, SO_SNDTIMEO, &timeout, sizeof(timeout)); + } + + ~Control() { + if (fd_ >= 0) { + close(fd_); + } + } + + void + Send(char message) const { + if (send(fd_, &message, 1, MSG_NOSIGNAL) != 1) { + Fail("control send"); + } + } + + void + Expect(char expected) const { + char actual = 0; + if (recv(fd_, &actual, 1, MSG_WAITALL) != 1 || actual != expected) { + Fail("control receive/order"); + } + } + +private: + int fd_ = -1; +}; + +void +AtomicWrite(const fs::path &path, const std::string &data) { + const fs::path temporary = path.string() + ".tmp." + std::to_string(getpid()); + std::ofstream output(temporary, std::ios::binary); + if (!output) { + Fail("open metadata for write: " + temporary.string()); + } + output.write(data.data(), static_cast(data.size())); + output.close(); + fs::rename(temporary, path); +} + +std::string +ReadFile(const fs::path &path) { + std::ifstream input(path, std::ios::binary); + if (!input) { + Fail("open metadata for read: " + path.string()); + } + return {std::istreambuf_iterator(input), std::istreambuf_iterator()}; +} + +__device__ uint8_t +Pattern(uint64_t seed, size_t offset) { + return static_cast((seed + offset * 1315423911ULL + (offset >> 7U) * 17ULL) & 0xffU); +} + +__global__ void +FillKernel(uint8_t *data, size_t bytes, uint64_t seed) { + for (size_t offset = blockIdx.x * blockDim.x + threadIdx.x; offset < bytes; + offset += blockDim.x * gridDim.x) { + data[offset] = Pattern(seed, offset); + } +} + +__global__ void +EpochKernel(uint64_t *epoch, uint64_t value) { + if (blockIdx.x == 0 && threadIdx.x == 0) { + *epoch = value; + } +} + +__global__ void +DelayKernel(uint64_t cycles) { + const uint64_t begin = clock64(); + while (clock64() - begin < cycles) {} +} + +__global__ void +VerifyKernel(const uint8_t *data, + size_t bytes, + uint64_t seed, + const uint64_t *epoch, + uint64_t expected_epoch, + unsigned long long *mismatches) { + __shared__ int epoch_matches; + if (threadIdx.x == 0) { + cuda::atomic_ref marker( + *const_cast(reinterpret_cast(epoch))); + epoch_matches = marker.load(cuda::memory_order_acquire) == expected_epoch; + if (!epoch_matches) { + atomicAdd(mismatches, 1ULL); + } + } + __syncthreads(); + if (!epoch_matches) { + return; + } + unsigned long long local = 0; + for (size_t offset = blockIdx.x * blockDim.x + threadIdx.x; offset < bytes; + offset += blockDim.x * gridDim.x) { + local += data[offset] != Pattern(seed, offset); + } + if (local != 0) { + atomicAdd(mismatches, local); + } +} + +class DeviceMemory { +public: + DeviceMemory(size_t bytes, int device) : bytes_(bytes), device_(device) { + CheckCuda(cudaSetDevice(device_), "cudaSetDevice allocation"); + CheckCuda(cudaMalloc(&data_, bytes_), "cudaMalloc data"); + CheckCuda(cudaMalloc(&epoch_, sizeof(uint64_t)), "cudaMalloc epoch"); + } + + ~DeviceMemory() { + cudaSetDevice(device_); + if (epoch_ != nullptr) { + cudaFree(epoch_); + } + if (data_ != nullptr) { + cudaFree(data_); + } + } + + DeviceMemory(const DeviceMemory &) = delete; + DeviceMemory & + operator=(const DeviceMemory &) = delete; + + uint8_t * + data() const { + return data_; + } + + uint64_t * + epoch() const { + return epoch_; + } + + size_t + bytes() const { + return bytes_; + } + + int + device() const { + return device_; + } + +private: + size_t bytes_; + int device_; + uint8_t *data_ = nullptr; + uint64_t *epoch_ = nullptr; +}; + +class VerifyCounter { +public: + explicit VerifyCounter(int device) : device_(device) { + CheckCuda(cudaSetDevice(device_), "cudaSetDevice counter"); + CheckCuda(cudaMalloc(&device_counter_, sizeof(*device_counter_)), + "cudaMalloc verify counter"); + CheckCuda(cudaMallocHost(&host_counter_, sizeof(*host_counter_)), + "cudaMallocHost verify counter"); + } + + ~VerifyCounter() { + cudaSetDevice(device_); + if (host_counter_ != nullptr) { + cudaFreeHost(host_counter_); + } + if (device_counter_ != nullptr) { + cudaFree(device_counter_); + } + } + + uint64_t + Verify(const DeviceMemory &memory, uint64_t seed, uint64_t epoch, cudaStream_t stream = 0) { + CheckCuda(cudaSetDevice(device_), "cudaSetDevice verify"); + CheckCuda(cudaMemsetAsync(device_counter_, 0, sizeof(*device_counter_), stream), + "clear verify counter"); + VerifyKernel<<(65535, (memory.bytes() + 255) / 256), 256, 0, stream>>>( + memory.data(), memory.bytes(), seed, memory.epoch(), epoch, device_counter_); + CheckCuda(cudaGetLastError(), "launch verify kernel"); + CheckCuda(cudaMemcpyAsync(host_counter_, + device_counter_, + sizeof(*host_counter_), + cudaMemcpyDeviceToHost, + stream), + "copy verify counter"); + CheckCuda(cudaStreamSynchronize(stream), "verify stream synchronization"); + return *host_counter_; + } + +private: + int device_; + unsigned long long *device_counter_ = nullptr; + unsigned long long *host_counter_ = nullptr; +}; + +void +Fill(const DeviceMemory &memory, uint64_t seed, uint64_t epoch, cudaStream_t stream = 0) { + CheckCuda(cudaSetDevice(memory.device()), "cudaSetDevice fill"); + FillKernel<<(65535, (memory.bytes() + 255) / 256), 256, 0, stream>>>( + memory.data(), memory.bytes(), seed); + CheckCuda(cudaGetLastError(), "launch fill kernel"); + EpochKernel<<<1, 1, 0, stream>>>(memory.epoch(), epoch); + CheckCuda(cudaGetLastError(), "launch epoch kernel"); +} + +nixl_b_params_t +BackendParams(const Config &config, int port, int device) { + nixl_b_params_t params{{"gpu_devices", std::to_string(device)}, + {"cuda_streams", "4"}, + {"data_qp_count", "1"}, + {"oob_port", std::to_string(port)}}; + if (!config.network_device.empty()) { + params["network_devices"] = config.network_device; + } + if (!config.oob_interface.empty()) { + params["oob_interface"] = config.oob_interface; + } + if (!config.gid_index.empty()) { + params["gid_index"] = config.gid_index; + } + return params; +} + +nixl_opt_args_t +BackendOnly(nixlBackendH *backend) { + nixl_opt_args_t options; + options.backends.push_back(backend); + return options; +} + +nixl_reg_dlist_t +Registration(const DeviceMemory &memory) { + nixl_reg_dlist_t registration(VRAM_SEG); + registration.addDesc( + nixlBlobDesc(reinterpret_cast(memory.data()), memory.bytes(), memory.device())); + registration.addDesc(nixlBlobDesc( + reinterpret_cast(memory.epoch()), sizeof(uint64_t), memory.device())); + return registration; +} + +nixl_xfer_dlist_t +TransferList(const DeviceMemory &memory) { + nixl_xfer_dlist_t list(VRAM_SEG); + list.addDesc( + nixlBasicDesc(reinterpret_cast(memory.data()), memory.bytes(), memory.device())); + list.addDesc(nixlBasicDesc( + reinterpret_cast(memory.epoch()), sizeof(uint64_t), memory.device())); + return list; +} + +nixl_xfer_dlist_t +RemoteList(uintptr_t data, size_t bytes, uintptr_t epoch, int device) { + nixl_xfer_dlist_t list(VRAM_SEG); + list.addDesc(nixlBasicDesc(data, bytes, device)); + list.addDesc(nixlBasicDesc(epoch, sizeof(uint64_t), device)); + return list; +} + +struct Endpoint { + std::string name; + nixlAgent agent; + nixlBackendH *backend = nullptr; + nixl_opt_args_t options; + nixl_reg_dlist_t registration; + bool registration_active = false; + + Endpoint(const Config &config, std::string endpoint_name, int port, const DeviceMemory &memory) + : name(std::move(endpoint_name)), + agent(name, nixlAgentConfig(true)), + registration(Registration(memory)) { + CheckCuda(cudaSetDevice(memory.device()), "cudaSetDevice endpoint"); + CheckNixl( + agent.createBackend("GPUNETIO", BackendParams(config, port, memory.device()), backend), + "create GPUNETIO backend"); + options = BackendOnly(backend); + CheckNixl(agent.registerMem(registration, &options), "register GPUNETIO memory"); + registration_active = true; + } + + void + Deregister() { + if (registration_active) { + CheckNixl(agent.deregisterMem(registration, &options), "deregister GPUNETIO memory"); + registration_active = false; + } + } + + ~Endpoint() { + if (registration_active) { + agent.deregisterMem(registration, &options); + } + } +}; + +struct RemoteAddresses { + uintptr_t a_data = 0; + uintptr_t a_epoch = 0; + uintptr_t b_data = 0; + uintptr_t b_epoch = 0; +}; + +void +PublishMetadata(const Config &config, + const Endpoint &a, + const DeviceMemory &a_memory, + const Endpoint &b, + const DeviceMemory &b_memory) { + fs::create_directories(config.coord); + std::string metadata; + CheckNixl(a.agent.getLocalMD(metadata), "get target A metadata"); + AtomicWrite(config.coord / "target-a.md", metadata); + CheckNixl(b.agent.getLocalMD(metadata), "get target B metadata"); + AtomicWrite(config.coord / "target-b.md", metadata); + AtomicWrite(config.coord / "target-addresses.txt", + std::to_string(reinterpret_cast(a_memory.data())) + " " + + std::to_string(reinterpret_cast(a_memory.epoch())) + " " + + std::to_string(reinterpret_cast(b_memory.data())) + " " + + std::to_string(reinterpret_cast(b_memory.epoch())) + "\n"); +} + +RemoteAddresses +LoadMetadata(const Config &config, nixlAgent &source) { + std::string name; + CheckNixl(source.loadRemoteMD(ReadFile(config.coord / "target-a.md"), name), + "load target A metadata"); + if (name != "qp-progress-a") { + Fail("target A metadata name mismatch"); + } + CheckNixl(source.loadRemoteMD(ReadFile(config.coord / "target-b.md"), name), + "load target B metadata"); + if (name != "qp-progress-b") { + Fail("target B metadata name mismatch"); + } + RemoteAddresses addresses; + std::istringstream input(ReadFile(config.coord / "target-addresses.txt")); + input >> addresses.a_data >> addresses.a_epoch >> addresses.b_data >> addresses.b_epoch; + if (!input) { + Fail("invalid target address metadata"); + } + return addresses; +} + +nixl_opt_args_t +Attached(nixlBackendH *backend, cudaStream_t stream, std::optional notification) { + nixl_opt_args_t options = BackendOnly(backend); + options.customParam.resize(sizeof(stream)); + std::memcpy(options.customParam.data(), &stream, sizeof(stream)); + options.skipDescMerge = true; + if (notification) { + options.notif = *notification; + } + return options; +} + +nixl_status_t +Wait(nixlAgent &agent, nixlXferReqH *request, std::chrono::milliseconds timeout) { + const auto deadline = Clock::now() + timeout; + nixl_status_t status = NIXL_IN_PROG; + while (status == NIXL_IN_PROG && Clock::now() < deadline) { + status = agent.getXferStatus(request); + } + return status; +} + +void +Release(nixlAgent &agent, nixlXferReqH *request) { + CheckNixl(agent.releaseXferReq(request), "release transfer request"); +} + +uint64_t +SeedA(uint64_t round) { + return 0xa500000000000031ULL + (round << 8U); +} + +uint64_t +SeedB(uint64_t round) { + return 0xb400000000000073ULL + (round << 8U); +} + +uint64_t +EpochA(uint64_t round) { + return 0xe100000000000000ULL + round; +} + +uint64_t +EpochB(uint64_t round) { + return 0xe200000000000000ULL + round; +} + +struct TransferTimes { + uint64_t start_ns = 0; + uint64_t done_ns = 0; +}; + +TransferTimes +PostAndWait(nixlAgent &source, nixlXferReqH *request) { + TransferTimes times{NowNs(), 0}; + const nixl_status_t post = source.postXferReq(request); + if (post != NIXL_SUCCESS && post != NIXL_IN_PROG) { + Fail("post transfer: status=" + std::to_string(static_cast(post))); + } + const nixl_status_t status = + post == NIXL_SUCCESS ? post : Wait(source, request, std::chrono::seconds(20)); + if (status != NIXL_SUCCESS) { + Fail("transfer completion: status=" + std::to_string(static_cast(status))); + } + times.done_ns = NowNs(); + return times; +} + +uint64_t +Percentile(std::vector values, double percentile) { + std::sort(values.begin(), values.end()); + const size_t index = static_cast((values.size() - 1) * percentile); + return values[index]; +} + +void +WritePerformanceJson(const Config &config, + const std::vector &b_windows, + const std::vector &a_windows, + const std::vector &max_windows, + uint64_t payload_bytes, + uint64_t marker_bytes, + uint64_t wall_ns) { + AtomicWrite( + config.coord / "qp_progress_performance.json", + "{\"result\":\"PASS\",\"warmup\":" + std::to_string(kWarmup) + + ",\"measured\":" + std::to_string(kMeasured) + + ",\"external_pair_repeats_required\":3" + ",\"b_post_to_done_p50_ns\":" + + std::to_string(Percentile(b_windows, 0.50)) + + ",\"b_post_to_done_p99_ns\":" + std::to_string(Percentile(b_windows, 0.99)) + + ",\"a_post_to_done_p50_ns\":" + std::to_string(Percentile(a_windows, 0.50)) + + ",\"max_window_p50_ns\":" + std::to_string(Percentile(max_windows, 0.50)) + + ",\"actual_payload_bytes\":" + std::to_string(payload_bytes) + + ",\"actual_marker_bytes\":" + std::to_string(marker_bytes) + + ",\"actual_transfer_window_bytes\":" + std::to_string(payload_bytes + marker_bytes) + + ",\"walltime_ns\":" + std::to_string(wall_ns) + + ",\"walltime_scope\":\"caller only; not serving\"" + + ",\"window_definition\":\"source post to source API completion; no CQ timestamps\"}\n"); +} + +void +TargetPerformance(const Config &config) { + // All target resources exist before the persistent backend is constructed. + DeviceMemory a_memory(kLargeBytes, 0); + DeviceMemory b_memory(kSmallBytes, 1); + VerifyCounter a_counter(0); + VerifyCounter b_counter(1); + Endpoint a(config, "qp-progress-a", config.target_a_port, a_memory); + Endpoint b(config, "qp-progress-b", config.target_b_port, b_memory); + Control control(config); + PublishMetadata(config, a, a_memory, b, b_memory); + control.Send('R'); + for (uint64_t repeat = 0; repeat < kRepeats; ++repeat) { + for (uint64_t round = 0; round < kWarmup + kMeasured; ++round) { + control.Expect('D'); + ASSERT_EQ(a_counter.Verify( + a_memory, SeedA(repeat * 1000 + round), EpochA(repeat * 1000 + round)), + 0U); + ASSERT_EQ(b_counter.Verify( + b_memory, SeedB(repeat * 1000 + round), EpochB(repeat * 1000 + round)), + 0U); + control.Send('V'); + } + } + control.Expect('C'); + control.Send('K'); +} + +void +SourcePerformance(const Config &config) { + DeviceMemory a_memory(kLargeBytes, 0); + DeviceMemory b_memory(kSmallBytes, 0); + Endpoint source(config, "qp-progress-source", config.source_port, a_memory); + // Both source buffers are allocated before backend construction; B is registered before posts. + nixl_reg_dlist_t b_registration = Registration(b_memory); + CheckNixl(source.agent.registerMem(b_registration, &source.options), + "register source B memory"); + Control control(config); + control.Expect('R'); + const RemoteAddresses remote = LoadMetadata(config, source.agent); + std::vector b_windows, a_windows, max_windows; + const uint64_t wall_start = NowNs(); + for (uint64_t repeat = 0; repeat < kRepeats; ++repeat) { + for (uint64_t round = 0; round < kWarmup + kMeasured; ++round) { + const uint64_t token = repeat * 1000 + round; + Fill(a_memory, SeedA(token), EpochA(token)); + Fill(b_memory, SeedB(token), EpochB(token)); + CheckCuda(cudaStreamSynchronize(0), "prepare performance payloads"); + nixlXferReqH *a_request = nullptr; + nixlXferReqH *b_request = nullptr; + const auto a_local = TransferList(a_memory); + const auto b_local = TransferList(b_memory); + const auto a_remote = RemoteList(remote.a_data, kLargeBytes, remote.a_epoch, 0); + const auto b_remote = RemoteList(remote.b_data, kSmallBytes, remote.b_epoch, 1); + CheckNixl( + source.agent.createXferReq( + NIXL_WRITE, a_local, a_remote, "qp-progress-a", a_request, &source.options), + "create performance A"); + CheckNixl( + source.agent.createXferReq( + NIXL_WRITE, b_local, b_remote, "qp-progress-b", b_request, &source.options), + "create performance B"); + const uint64_t start = NowNs(); + const nixl_status_t a_post = source.agent.postXferReq(a_request); + const uint64_t b_post_ns = NowNs(); + const nixl_status_t b_post = source.agent.postXferReq(b_request); + ASSERT_TRUE(a_post == NIXL_SUCCESS || a_post == NIXL_IN_PROG); + ASSERT_TRUE(b_post == NIXL_SUCCESS || b_post == NIXL_IN_PROG); + nixl_status_t a_status = a_post; + nixl_status_t b_status = b_post; + uint64_t a_done = a_status == NIXL_SUCCESS ? NowNs() : 0; + uint64_t b_done = b_status == NIXL_SUCCESS ? NowNs() : 0; + const auto deadline = Clock::now() + std::chrono::seconds(20); + while ((a_status == NIXL_IN_PROG || b_status == NIXL_IN_PROG) && + Clock::now() < deadline) { + if (b_status == NIXL_IN_PROG) { + b_status = source.agent.getXferStatus(b_request); + if (b_status == NIXL_SUCCESS) { + b_done = NowNs(); + } + } + if (a_status == NIXL_IN_PROG) { + a_status = source.agent.getXferStatus(a_request); + if (a_status == NIXL_SUCCESS) { + a_done = NowNs(); + } + } + } + ASSERT_EQ(a_status, NIXL_SUCCESS); + ASSERT_EQ(b_status, NIXL_SUCCESS); + if (round >= kWarmup) { + b_windows.push_back(b_done - b_post_ns); + a_windows.push_back(a_done - start); + max_windows.push_back(std::max(a_done, b_done) - start); + } + Release(source.agent, b_request); + Release(source.agent, a_request); + control.Send('D'); + control.Expect('V'); + } + } + const uint64_t wall_ns = NowNs() - wall_start; + WritePerformanceJson(config, + b_windows, + a_windows, + max_windows, + kRepeats * kMeasured * (kLargeBytes + kSmallBytes), + kRepeats * kMeasured * 2 * sizeof(uint64_t), + wall_ns); + CheckNixl(source.agent.invalidateRemoteMD("qp-progress-a"), "invalidate target A metadata"); + CheckNixl(source.agent.invalidateRemoteMD("qp-progress-b"), "invalidate target B metadata"); + CheckNixl(source.agent.deregisterMem(b_registration, &source.options), + "deregister source B memory"); + control.Send('C'); + control.Expect('K'); +} + +// The stress layout creates 512 payload descriptors plus one exact epoch descriptor per request. +struct StressMemory { + std::vector> slots; + + explicit StressMemory(int device) { + slots.reserve(kStressRequests); + for (size_t slot = 0; slot < kStressRequests; ++slot) { + slots.emplace_back( + std::make_unique(kStressDescriptors * kSmallBytes, device)); + } + } +}; + +nixl_reg_dlist_t +StressRegistration(const StressMemory &memory) { + nixl_reg_dlist_t list(VRAM_SEG); + for (const auto &slot : memory.slots) { + list.addDesc( + nixlBlobDesc(reinterpret_cast(slot->data()), slot->bytes(), slot->device())); + list.addDesc(nixlBlobDesc( + reinterpret_cast(slot->epoch()), sizeof(uint64_t), slot->device())); + } + return list; +} + +nixl_xfer_dlist_t +StressList(const DeviceMemory &memory) { + nixl_xfer_dlist_t list(VRAM_SEG); + for (size_t descriptor = 0; descriptor < kStressDescriptors; ++descriptor) { + list.addDesc( + nixlBasicDesc(reinterpret_cast(memory.data() + descriptor * kSmallBytes), + kSmallBytes, + memory.device())); + } + list.addDesc(nixlBasicDesc( + reinterpret_cast(memory.epoch()), sizeof(uint64_t), memory.device())); + return list; +} + +nixl_xfer_dlist_t +StressRemoteList(uintptr_t data, size_t bytes, uintptr_t epoch, int device) { + nixl_xfer_dlist_t list(VRAM_SEG); + for (size_t descriptor = 0; descriptor < kStressDescriptors; ++descriptor) { + list.addDesc(nixlBasicDesc(data + descriptor * kSmallBytes, kSmallBytes, device)); + } + list.addDesc(nixlBasicDesc(epoch, sizeof(uint64_t), device)); + return list; +} + +void +TargetStress(const Config &config) { + StressMemory a_memory(0); + StressMemory b_memory(1); + std::vector> a_counters; + std::vector> b_counters; + for (size_t index = 0; index < kStressRequests; ++index) { + a_counters.emplace_back(std::make_unique(0)); + b_counters.emplace_back(std::make_unique(1)); + } + // Register all slots before persistent engines; source uses matching full registrations. + DeviceMemory a_anchor(1, 0); + DeviceMemory b_anchor(1, 1); + Endpoint a(config, "qp-progress-a", config.target_a_port, a_anchor); + Endpoint b(config, "qp-progress-b", config.target_b_port, b_anchor); + const nixl_reg_dlist_t a_registration = StressRegistration(a_memory); + const nixl_reg_dlist_t b_registration = StressRegistration(b_memory); + CheckNixl(a.agent.registerMem(a_registration, &a.options), "register target A stress slots"); + CheckNixl(b.agent.registerMem(b_registration, &b.options), "register target B stress slots"); + Control control(config); + PublishMetadata(config, a, a_anchor, b, b_anchor); + // Slot addresses are separate metadata because serialized NIXL metadata carries registration + // information. + std::ofstream addresses(config.coord / "stress-addresses.txt"); + for (size_t index = 0; index < kStressRequests; ++index) { + addresses << reinterpret_cast(a_memory.slots[index]->data()) << ' ' + << reinterpret_cast(a_memory.slots[index]->epoch()) << ' ' + << reinterpret_cast(b_memory.slots[index]->data()) << ' ' + << reinterpret_cast(b_memory.slots[index]->epoch()) << '\n'; + } + addresses.close(); + control.Send('R'); + nixl_notifs_t a_notifications; + nixl_notifs_t b_notifications; + auto poll_notifications = [&] { + ASSERT_EQ(a.agent.getNotifs(a_notifications), NIXL_SUCCESS); + ASSERT_EQ(b.agent.getNotifs(b_notifications), NIXL_SUCCESS); + }; + for (size_t index = 0; index < kStressRequests; ++index) { + control.Expect('B'); + ASSERT_EQ(b_counters[index]->Verify(*b_memory.slots[index], SeedB(index), EpochB(index)), + 0U); + poll_notifications(); + control.Send('V'); + } + for (size_t index = 1; index < kStressRequests; index += 2) { + control.Expect('A'); + ASSERT_EQ(a_counters[index]->Verify(*a_memory.slots[index], SeedA(index), EpochA(index)), + 0U); + poll_notifications(); + control.Send('V'); + } + for (size_t index = 0; index < kStressRequests; index += 2) { + control.Expect('A'); + ASSERT_EQ(a_counters[index]->Verify(*a_memory.slots[index], SeedA(index), EpochA(index)), + 0U); + poll_notifications(); + control.Send('V'); + } + for (size_t index = 0; index < kFastChurnRequests; ++index) { + control.Expect('B'); + ASSERT_EQ(b_counters[0]->Verify(*b_memory.slots[0], + SeedB(kStressRequests + index), + EpochB(kStressRequests + index)), + 0U); + poll_notifications(); + control.Send('V'); + } + control.Expect('D'); + const auto deadline = Clock::now() + std::chrono::seconds(10); + while ( + (a_notifications["qp-progress-source"].size() < kStressRequests + 1 || + b_notifications["qp-progress-source"].size() < kStressRequests + kFastChurnRequests + 1) && + Clock::now() < deadline) { + poll_notifications(); + } + const auto &a_messages = a_notifications["qp-progress-source"]; + const auto &b_messages = b_notifications["qp-progress-source"]; + ASSERT_EQ(a_messages.size(), kStressRequests + 1); + ASSERT_EQ(b_messages.size(), kStressRequests + kFastChurnRequests + 1); + for (size_t index = 0; index < kStressRequests; ++index) { + EXPECT_NE( + std::find(a_messages.begin(), a_messages.end(), "data-a-" + std::to_string(index)), + a_messages.end()); + EXPECT_NE( + std::find(b_messages.begin(), b_messages.end(), "data-b-" + std::to_string(index)), + b_messages.end()); + } + for (size_t index = 0; index < kFastChurnRequests; ++index) { + EXPECT_NE(std::find(b_messages.begin(), + b_messages.end(), + "data-b-" + std::to_string(kStressRequests + index)), + b_messages.end()); + } + EXPECT_NE(std::find(a_messages.begin(), a_messages.end(), "standalone-a"), a_messages.end()); + EXPECT_NE(std::find(b_messages.begin(), b_messages.end(), "standalone-b"), b_messages.end()); + control.Send('V'); + control.Expect('C'); + CheckNixl(a.agent.deregisterMem(a_registration, &a.options), + "deregister target A stress slots"); + CheckNixl(b.agent.deregisterMem(b_registration, &b.options), + "deregister target B stress slots"); + control.Send('K'); +} + +std::array +ReadStressAddresses(const Config &config) { + std::array addresses{}; + std::ifstream input(config.coord / "stress-addresses.txt"); + for (auto &address : addresses) { + input >> address.a_data >> address.a_epoch >> address.b_data >> address.b_epoch; + if (!input) { + Fail("invalid stress address metadata"); + } + } + return addresses; +} + +void +SourceStress(const Config &config) { + StressMemory a_memory(0); + StressMemory b_memory(0); + DeviceMemory anchor(1, 0); + Endpoint source(config, "qp-progress-source", config.source_port, anchor); + const nixl_reg_dlist_t a_registration = StressRegistration(a_memory); + const nixl_reg_dlist_t b_registration = StressRegistration(b_memory); + CheckNixl(source.agent.registerMem(a_registration, &source.options), + "register source A stress slots"); + CheckNixl(source.agent.registerMem(b_registration, &source.options), + "register source B stress slots"); + cudaStream_t delayed_stream = nullptr; + cudaStream_t fast_stream = nullptr; + CheckCuda(cudaStreamCreateWithFlags(&delayed_stream, cudaStreamNonBlocking), + "create delayed stream"); + CheckCuda(cudaStreamCreateWithFlags(&fast_stream, cudaStreamNonBlocking), "create fast stream"); + Control control(config); + control.Expect('R'); + LoadMetadata(config, source.agent); + const auto remote = ReadStressAddresses(config); + std::array a_requests{}; + std::array b_requests{}; + // Alternating attached streams on the same A peer reverse CPU post order versus SQ readiness. + DelayKernel<<<1, 1, 0, delayed_stream>>>(50000000ULL); + CheckCuda(cudaGetLastError(), "launch one test-only delayed stream kernel"); + for (size_t index = 0; index < kStressRequests; ++index) { + cudaStream_t stream = index % 2 == 0 ? delayed_stream : fast_stream; + Fill(*a_memory.slots[index], SeedA(index), EpochA(index), stream); + auto options = Attached(source.backend, stream, "data-a-" + std::to_string(index)); + CheckNixl(source.agent.createXferReq(NIXL_WRITE, + StressList(*a_memory.slots[index]), + StressRemoteList(remote[index].a_data, + a_memory.slots[index]->bytes(), + remote[index].a_epoch, + 0), + "qp-progress-a", + a_requests[index], + &options), + "create alternating A stress request"); + ASSERT_TRUE(source.agent.postXferReq(a_requests[index]) >= NIXL_SUCCESS); + } + for (size_t index = 0; index < kStressRequests; ++index) { + Fill(*b_memory.slots[index], SeedB(index), EpochB(index), fast_stream); + auto options = Attached(source.backend, fast_stream, "data-b-" + std::to_string(index)); + CheckNixl(source.agent.createXferReq(NIXL_WRITE, + StressList(*b_memory.slots[index]), + StressRemoteList(remote[index].b_data, + b_memory.slots[index]->bytes(), + remote[index].b_epoch, + 1), + "qp-progress-b", + b_requests[index], + &options), + "create initial B stress request"); + ASSERT_TRUE(source.agent.postXferReq(b_requests[index]) >= NIXL_SUCCESS); + } + for (size_t index = 0; index < kStressRequests; ++index) { + ASSERT_EQ(Wait(source.agent, b_requests[index], std::chrono::seconds(10)), NIXL_SUCCESS); + Release(source.agent, b_requests[index]); + control.Send('B'); + control.Expect('V'); + } + for (size_t index = 1; index < kStressRequests; index += 2) { + ASSERT_EQ(Wait(source.agent, a_requests[index], std::chrono::seconds(10)), NIXL_SUCCESS); + Release(source.agent, a_requests[index]); + control.Send('A'); + control.Expect('V'); + } + for (size_t index = 0; index < kStressRequests; index += 2) { + ASSERT_EQ(Wait(source.agent, a_requests[index], std::chrono::seconds(10)), NIXL_SUCCESS); + Release(source.agent, a_requests[index]); + control.Send('A'); + control.Expect('V'); + } + for (size_t index = 0; index < kFastChurnRequests; ++index) { + const size_t token = kStressRequests + index; + Fill(*b_memory.slots[0], SeedB(token), EpochB(token), fast_stream); + nixlXferReqH *b_request = nullptr; + auto b_options = Attached(source.backend, fast_stream, "data-b-" + std::to_string(token)); + CheckNixl(source.agent.createXferReq( + NIXL_WRITE, + StressList(*b_memory.slots[0]), + StressRemoteList( + remote[0].b_data, b_memory.slots[0]->bytes(), remote[0].b_epoch, 1), + "qp-progress-b", + b_request, + &b_options), + "create fast stress request"); + ASSERT_TRUE(source.agent.postXferReq(b_request) >= NIXL_SUCCESS); + ASSERT_EQ(Wait(source.agent, b_request, std::chrono::seconds(10)), NIXL_SUCCESS); + Release(source.agent, b_request); + control.Send('B'); + control.Expect('V'); + } + CheckNixl(source.agent.genNotif("qp-progress-a", "standalone-a", &source.options), + "generate standalone notification for target A"); + CheckNixl(source.agent.genNotif("qp-progress-b", "standalone-b", &source.options), + "generate standalone notification for target B"); + // Attached transfer completion is the progress/reuse boundary. + control.Send('D'); + control.Expect('V'); + CheckNixl(source.agent.deregisterMem(a_registration, &source.options), + "deregister source A stress slots"); + CheckNixl(source.agent.deregisterMem(b_registration, &source.options), + "deregister source B stress slots"); + CheckCuda(cudaStreamDestroy(fast_stream), "destroy fast stream"); + CheckCuda(cudaStreamDestroy(delayed_stream), "destroy delayed stream"); + control.Send('C'); + control.Expect('K'); +} + +void +TargetReadWriteControl(const Config &config) { + DeviceMemory write_memory(config.control_bytes, 0); + DeviceMemory read_memory(config.control_bytes, 1); + VerifyCounter write_counter(0); + Endpoint write_target(config, "qp-progress-a", config.target_a_port, write_memory); + Endpoint read_target(config, "qp-progress-b", config.target_b_port, read_memory); + Control control(config); + PublishMetadata(config, write_target, write_memory, read_target, read_memory); + control.Send('R'); + for (uint64_t round = 0; round < kWarmup + kMeasured; ++round) { + control.Expect('W'); + ASSERT_EQ(write_counter.Verify(write_memory, SeedA(round), EpochA(round)), 0U); + control.Send('V'); + Fill(read_memory, SeedB(round), EpochB(round)); + CheckCuda(cudaStreamSynchronize(0), "prepare control READ payload"); + control.Send('P'); + control.Expect('D'); + } + control.Expect('C'); + control.Send('K'); +} + +void +SourceReadWriteControl(const Config &config) { + DeviceMemory write_memory(config.control_bytes, 0); + DeviceMemory read_memory(config.control_bytes, 0); + VerifyCounter read_counter(0); + Endpoint source(config, "qp-progress-source", config.source_port, write_memory); + const nixl_reg_dlist_t read_registration = Registration(read_memory); + CheckNixl(source.agent.registerMem(read_registration, &source.options), + "register control READ memory"); + Control control(config); + control.Expect('R'); + const RemoteAddresses remote = LoadMetadata(config, source.agent); + nixlXferReqH *write_request = nullptr; + nixlXferReqH *read_request = nullptr; + CheckNixl(source.agent.createXferReq( + NIXL_WRITE, + TransferList(write_memory), + RemoteList(remote.a_data, config.control_bytes, remote.a_epoch, 0), + "qp-progress-a", + write_request, + &source.options), + "create repostable control WRITE"); + CheckNixl(source.agent.createXferReq( + NIXL_READ, + TransferList(read_memory), + RemoteList(remote.b_data, config.control_bytes, remote.b_epoch, 1), + "qp-progress-b", + read_request, + &source.options), + "create repostable control READ"); + std::vector write_windows; + std::vector read_windows; + for (uint64_t round = 0; round < kWarmup + kMeasured; ++round) { + Fill(write_memory, SeedA(round), EpochA(round)); + CheckCuda(cudaStreamSynchronize(0), "prepare control WRITE payload"); + const TransferTimes write_times = PostAndWait(source.agent, write_request); + control.Send('W'); + control.Expect('V'); + control.Expect('P'); + const TransferTimes read_times = PostAndWait(source.agent, read_request); + ASSERT_EQ(read_counter.Verify(read_memory, SeedB(round), EpochB(round)), 0U); + if (round >= kWarmup) { + write_windows.push_back(write_times.done_ns - write_times.start_ns); + read_windows.push_back(read_times.done_ns - read_times.start_ns); + } + control.Send('D'); + } + AtomicWrite(config.coord / "qp_progress_single_qp_control.json", + "{\"result\":\"PASS\",\"warmup\":" + std::to_string(kWarmup) + + ",\"measured\":" + std::to_string(kMeasured) + + ",\"bytes_per_op\":" + std::to_string(config.control_bytes + sizeof(uint64_t)) + + ",\"write_p50_ns\":" + std::to_string(Percentile(write_windows, 0.50)) + + ",\"write_p99_ns\":" + std::to_string(Percentile(write_windows, 0.99)) + + ",\"read_p50_ns\":" + std::to_string(Percentile(read_windows, 0.50)) + + ",\"read_p99_ns\":" + std::to_string(Percentile(read_windows, 0.99)) + + ",\"window_definition\":\"source post to source API completion\"}\n"); + Release(source.agent, read_request); + Release(source.agent, write_request); + CheckNixl(source.agent.invalidateRemoteMD("qp-progress-a"), "invalidate control A metadata"); + CheckNixl(source.agent.invalidateRemoteMD("qp-progress-b"), "invalidate control B metadata"); + CheckNixl(source.agent.deregisterMem(read_registration, &source.options), + "deregister control READ memory"); + control.Send('C'); + control.Expect('K'); +} + +void +TargetFault(const Config &config) { + DeviceMemory a_memory(kLargeBytes, 0); + DeviceMemory b_memory(kSmallBytes, 1); + Endpoint a(config, "qp-progress-a", config.target_a_port, a_memory); + Endpoint b(config, "qp-progress-b", config.target_b_port, b_memory); + Control control(config); + PublishMetadata(config, a, a_memory, b, b_memory); + control.Send('R'); + control.Expect('M'); + b.Deregister(); + control.Send('F'); + control.Expect('E'); +} + +void +SourceFault(const Config &config) { + DeviceMemory b_memory(kSmallBytes, 0); + Endpoint source(config, "qp-progress-source", config.source_port, b_memory); + Control control(config); + control.Expect('R'); + const RemoteAddresses remote = LoadMetadata(config, source.agent); + control.Send('M'); + control.Expect('F'); + Fill(b_memory, SeedB(0), EpochB(0)); + CheckCuda(cudaStreamSynchronize(0), "prepare fault payload"); + nixlXferReqH *request = nullptr; + CheckNixl(source.agent.createXferReq(NIXL_WRITE, + TransferList(b_memory), + RemoteList(remote.b_data, kSmallBytes, remote.b_epoch, 1), + "qp-progress-b", + request, + &source.options), + "create fault request"); + const nixl_status_t post = source.agent.postXferReq(request); + const nixl_status_t status = + post == NIXL_IN_PROG ? Wait(source.agent, request, std::chrono::seconds(10)) : post; + ASSERT_EQ(status, NIXL_ERR_BACKEND); + ASSERT_EQ(source.agent.releaseXferReq(request), NIXL_SUCCESS); + control.Send('E'); +} + +TEST(QpProgress, PerformanceMixed2MiBAnd4KiB) { + const Config config = GetConfig(); + if (const auto problem = ConfigProblem(config, false)) { + GTEST_SKIP() << *problem; + } + if (config.role == "source") { + SourcePerformance(config); + } else { + TargetPerformance(config); + } +} + +TEST(QpProgress, OutstandingDescriptorsAndAttachedStreamOrder) { + const Config config = GetConfig(); + if (const auto problem = ConfigProblem(config, false)) { + GTEST_SKIP() << *problem; + } + if (config.role == "source") { + SourceStress(config); + } else { + TargetStress(config); + } +} + +TEST(QpProgress, SingleQpReadWriteControl) { + const Config config = GetConfig(); + if (const auto problem = ConfigProblem(config, false)) { + GTEST_SKIP() << *problem; + } + if (config.role == "source") { + SourceReadWriteControl(config); + } else { + TargetReadWriteControl(config); + } +} + +TEST(QpProgress, RemoteDeregisterReturnsBackendError) { + const Config config = GetConfig(); + if (const auto problem = ConfigProblem(config, true)) { + GTEST_SKIP() << *problem; + } + if (config.role == "source") { + SourceFault(config); + } else { + TargetFault(config); + } +} + +} // namespace + +int +main(int argc, char **argv) { + ::testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); +} diff --git a/test/gtest/plugins/meson.build b/test/gtest/plugins/meson.build index cb4573e283..9dbd6c6695 100644 --- a/test/gtest/plugins/meson.build +++ b/test/gtest/plugins/meson.build @@ -21,6 +21,10 @@ if enabled_plugins.get('AZURE_BLOB') subdir('azure_blob') endif +if enabled_plugins.get('GPUNETIO') and cuda_dep.found() + subdir('gpunetio') +endif + if not enabled_plugins.get('OBJ') message('OBJ plugin not enabled, skipping plugins_gtest build') subdir_done() From b89405692b24cd588ee2c953454829c8bbc3e9b3 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Mon, 7 Sep 2026 19:51:59 +0800 Subject: [PATCH 12/25] fix(gpunetio): reject failed progress-state allocation Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index 4e4c8cec2c..91ffd9fcbf 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -250,7 +250,7 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) (void **)&progress_state_gpu, (void **)&progress_state_cpu); if (result != DOCA_SUCCESS || progress_state_gpu == nullptr || progress_state_cpu == nullptr) { - NIXL_ERROR << "Function doca_gpu_mem_alloc return " << doca_error_get_descr(result); + throw std::runtime_error("Failed to allocate GPUNETIO progress state"); } memset(progress_state_cpu, 0, sizeof(struct docaProgressState)); From 2d2e11a4bb3bc97e4c2260ad368d0e42924404ab Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Mon, 7 Sep 2026 20:00:22 +0800 Subject: [PATCH 13/25] test(gpunetio): report measured throughput and operation counts Signed-off-by: Yunxiao Ning --- .../gtest/plugins/gpunetio/qp_progress_gtest.cu | 17 +++++++++++++++-- 1 file changed, 15 insertions(+), 2 deletions(-) diff --git a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu index 4aac639b8a..949aab6118 100644 --- a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu +++ b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu @@ -23,6 +23,7 @@ #include #include #include +#include #include #include #include @@ -638,6 +639,10 @@ WritePerformanceJson(const Config &config, uint64_t payload_bytes, uint64_t marker_bytes, uint64_t wall_ns) { + const uint64_t transfer_ns = + std::accumulate(max_windows.begin(), max_windows.end(), uint64_t{0}); + const double transfer_gbps = double(payload_bytes + marker_bytes) / transfer_ns; + const double wall_gbps = double(payload_bytes + marker_bytes) / wall_ns; AtomicWrite( config.coord / "qp_progress_performance.json", "{\"result\":\"PASS\",\"warmup\":" + std::to_string(kWarmup) + @@ -651,7 +656,12 @@ WritePerformanceJson(const Config &config, ",\"actual_payload_bytes\":" + std::to_string(payload_bytes) + ",\"actual_marker_bytes\":" + std::to_string(marker_bytes) + ",\"actual_transfer_window_bytes\":" + std::to_string(payload_bytes + marker_bytes) + - ",\"walltime_ns\":" + std::to_string(wall_ns) + + ",\"transfer_window_ns\":" + std::to_string(transfer_ns) + + ",\"transfer_window_GBps\":" + std::to_string(transfer_gbps) + + ",\"wall_GBps\":" + std::to_string(wall_gbps) + + ",\"completed_requests\":" + std::to_string(2 * (kWarmup + kMeasured)) + + ",\"measured_requests\":" + std::to_string(2 * kMeasured) + + ",\"payload_and_marker_mismatches\":0" + ",\"walltime_ns\":" + std::to_string(wall_ns) + ",\"walltime_scope\":\"caller only; not serving\"" + ",\"window_definition\":\"source post to source API completion; no CQ timestamps\"}\n"); } @@ -697,7 +707,7 @@ SourcePerformance(const Config &config) { control.Expect('R'); const RemoteAddresses remote = LoadMetadata(config, source.agent); std::vector b_windows, a_windows, max_windows; - const uint64_t wall_start = NowNs(); + uint64_t wall_start = 0; for (uint64_t repeat = 0; repeat < kRepeats; ++repeat) { for (uint64_t round = 0; round < kWarmup + kMeasured; ++round) { const uint64_t token = repeat * 1000 + round; @@ -719,6 +729,9 @@ SourcePerformance(const Config &config) { NIXL_WRITE, b_local, b_remote, "qp-progress-b", b_request, &source.options), "create performance B"); const uint64_t start = NowNs(); + if (round == kWarmup) { + wall_start = start; + } const nixl_status_t a_post = source.agent.postXferReq(a_request); const uint64_t b_post_ns = NowNs(); const nixl_status_t b_post = source.agent.postXferReq(b_request); From dd3857a2c9e1d59192e014d1d97002a3c5c77770 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Mon, 7 Sep 2026 20:27:28 +0800 Subject: [PATCH 14/25] test(gpunetio): verify multi-chunk READ in paired stress case Signed-off-by: Yunxiao Ning --- test/gtest/plugins/gpunetio/README.md | 22 +++++++------ test/gtest/plugins/gpunetio/meson.build | 2 +- .../plugins/gpunetio/qp_progress_gtest.cu | 32 +++++++++++++++++++ 3 files changed, 46 insertions(+), 10 deletions(-) diff --git a/test/gtest/plugins/gpunetio/README.md b/test/gtest/plugins/gpunetio/README.md index a07cd7361e..557f715554 100644 --- a/test/gtest/plugins/gpunetio/README.md +++ b/test/gtest/plugins/gpunetio/README.md @@ -1,13 +1,13 @@ # GPUNETIO QP-progress GoogleTest harness `gpunetio_qp_progress_gtest` is a paired, environment-configured integration -executable. It is intentionally not part of the default test suite: it needs +executable. Running its paired cases needs one source GPU and two target GPUs, a numeric target IPv4 address, unique OOB ports, and a shared directory used only to exchange serialized NIXL metadata. Control/verification acknowledgements use TCP, outside the measured interval. The GPUNETIO GTest parent wires this directory. The local `meson.build` builds -the paired executable and registers it with CTest; absent role configuration is +the paired executable and registers it with Meson; absent role configuration is an explicit GoogleTest skip before any CUDA query. ## Required environment @@ -53,13 +53,17 @@ and separate wall time. It does not claim CQ timing or backend-internal timing. `OutstandingDescriptorsAndAttachedStreamOrder` uses 513 descriptors per request (512 data plus an epoch marker) with descriptor merging disabled. It -holds one delayed source-peer request while completing and releasing 129 -fast-peer requests, forcing live 32-slot ring reuse without exceeding capacity. -The delay is a test-only CUDA kernel enqueued before the attached transfer; it -is not a backend option or a production hook. The target verifies every fast -payload and the delayed payload after a system-acquire of each exact epoch -marker, then checks data-coupled notifications plus one standalone notification -per peer. `SingleQpReadWriteControl` provides the paired one-QP WRITE and READ +posts seven requests per peer, then releases the alternating delayed/fast A +requests before performing 129 sequential B-peer reuse posts. The delay is a +test-only CUDA kernel enqueued before the attached transfer; it is not a backend +option or a production hook. The target verifies every payload after a +system-acquire of its exact epoch marker, then checks data-coupled notifications +plus one standalone notification per peer. With descriptor merging disabled, +the initial 14 requests occupy 28 of the 32 ring entries (513 descriptors use +two entries each). After those requests are released, the harness performs one +matched 513-descriptor READ into the reused B source buffer and GPU-verifies its +peer-distinct payload and epoch. `SingleQpReadWriteControl` provides the paired +one-active-QP WRITE and READ control performance path. It re-posts one prepared WRITE and one prepared READ handle across changing payload epochs, with one active QP at a time, and emits source API p50/p99 windows separately from the mixed WRITE performance JSON. diff --git a/test/gtest/plugins/gpunetio/meson.build b/test/gtest/plugins/gpunetio/meson.build index fbb7f5d5e4..9d9a418085 100644 --- a/test/gtest/plugins/gpunetio/meson.build +++ b/test/gtest/plugins/gpunetio/meson.build @@ -16,5 +16,5 @@ gpunetio_qp_progress_gtest = executable( ) # The executable is always registered. Without the paired-role environment its -# GoogleTests explicitly SKIP, so CTest exposes rather than silently omits it. +# GoogleTests explicitly SKIP, so Meson exposes rather than silently omits it. test('gpunetio_qp_progress_gtest', gpunetio_qp_progress_gtest, is_parallel: false, timeout: 180) diff --git a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu index 949aab6118..92a4c0bb0d 100644 --- a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu +++ b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu @@ -903,6 +903,12 @@ TargetStress(const Config &config) { poll_notifications(); control.Send('V'); } + control.Expect('R'); + const size_t read_token = kStressRequests + kFastChurnRequests; + Fill(*b_memory.slots[0], SeedB(read_token), EpochB(read_token)); + CheckCuda(cudaStreamSynchronize(0), "prepare multi-chunk READ payload"); + control.Send('P'); + control.Expect('Q'); control.Expect('D'); const auto deadline = Clock::now() + std::chrono::seconds(10); while ( @@ -957,6 +963,7 @@ void SourceStress(const Config &config) { StressMemory a_memory(0); StressMemory b_memory(0); + VerifyCounter read_counter(0); DeviceMemory anchor(1, 0); Endpoint source(config, "qp-progress-source", config.source_port, anchor); const nixl_reg_dlist_t a_registration = StressRegistration(a_memory); @@ -1048,6 +1055,31 @@ SourceStress(const Config &config) { control.Send('B'); control.Expect('V'); } + // Every B request has completed, released, and been target-verified, so slot 0 is safe to + // reuse as the local destination of one matched 513-descriptor READ. + control.Send('R'); + control.Expect('P'); + nixlXferReqH *read_request = nullptr; + const auto read_options = Attached(source.backend, fast_stream, std::nullopt); + CheckNixl( + source.agent.createXferReq( + NIXL_READ, + StressList(*b_memory.slots[0]), + StressRemoteList(remote[0].b_data, b_memory.slots[0]->bytes(), remote[0].b_epoch, 1), + "qp-progress-b", + read_request, + &read_options), + "create multi-chunk stress READ"); + const nixl_status_t read_post = source.agent.postXferReq(read_request); + ASSERT_TRUE(read_post == NIXL_SUCCESS || read_post == NIXL_IN_PROG); + ASSERT_EQ(read_post == NIXL_SUCCESS ? + read_post : + Wait(source.agent, read_request, std::chrono::seconds(10)), + NIXL_SUCCESS); + Release(source.agent, read_request); + const size_t read_token = kStressRequests + kFastChurnRequests; + ASSERT_EQ(read_counter.Verify(*b_memory.slots[0], SeedB(read_token), EpochB(read_token)), 0U); + control.Send('Q'); CheckNixl(source.agent.genNotif("qp-progress-a", "standalone-a", &source.options), "generate standalone notification for target A"); CheckNixl(source.agent.genNotif("qp-progress-b", "standalone-b", &source.options), From 8481aec66306d53a00f906a6127ff56103220c63 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 09:47:13 +0800 Subject: [PATCH 15/25] fix(gpunetio): select engine GPU before QP allocation Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 51 +++++++++++++++++------ 1 file changed, 38 insertions(+), 13 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index 91ffd9fcbf..bb84e2870a 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -26,6 +26,36 @@ const char info_delimiter = '-'; +namespace { +class cudaDeviceGuard { +public: + explicit cudaDeviceGuard(uint32_t device) : device_(static_cast(device)) { + status_ = cudaGetDevice(&previous_device_); + if (status_ == cudaSuccess) { + status_ = cudaSetDevice(device_); + restore_ = status_ == cudaSuccess && previous_device_ != device_; + } + } + + ~cudaDeviceGuard() { + if (restore_) { + cudaSetDevice(previous_device_); + } + } + + cudaError_t + status() const { + return status_; + } + +private: + int device_; + int previous_device_ = 0; + bool restore_ = false; + cudaError_t status_ = cudaSuccess; +}; +} // namespace + /**************************************** * Constructor/Destructor *****************************************/ @@ -572,6 +602,13 @@ nixlDocaEngine::addRdmaQp(const std::string &remote_agent) { NIXL_DEBUG << "DOCA addRdmaQp for remote " << remote_agent << std::endl; + cudaDeviceGuard cuda_device(gdevs[0].first); + if (cuda_device.status() != cudaSuccess) { + NIXL_ERROR << "Failed to select CUDA device " << gdevs[0].first + << " for QP setup: " << cudaGetErrorString(cuda_device.status()); + return NIXL_ERR_BACKEND; + } + rdma_qp = new struct nixlDocaRdmaQp; try { rdma_qp->qp_data = @@ -621,16 +658,7 @@ nixlDocaEngine::addRdmaQp(const std::string &remote_agent) { } cudaStream_t init_stream = nullptr; - int previous_cuda_device = 0; - bool cuda_device_switched = false; - cudaError_t cuda_result = cudaGetDevice(&previous_cuda_device); - if (cuda_result == cudaSuccess) { - cuda_result = cudaSetDevice(gdevs[0].first); - cuda_device_switched = cuda_result == cudaSuccess; - } - if (cuda_result == cudaSuccess) { - cuda_result = cudaStreamCreateWithFlags(&init_stream, cudaStreamNonBlocking); - } + cudaError_t cuda_result = cudaStreamCreateWithFlags(&init_stream, cudaStreamNonBlocking); if (cuda_result == cudaSuccess) { cuda_result = cudaMemsetAsync(rdma_qp->progress_gpu, 0, sizeof(struct docaQpProgress), init_stream); @@ -645,9 +673,6 @@ nixlDocaEngine::addRdmaQp(const std::string &remote_agent) { if (init_stream != nullptr) { cudaStreamDestroy(init_stream); } - if (cuda_device_switched && previous_cuda_device != static_cast(gdevs[0].first)) { - cudaSetDevice(previous_cuda_device); - } if (cuda_result != cudaSuccess) { NIXL_ERROR << "Failed to initialize QP progress state " << cudaGetErrorString(cuda_result); doca_gpu_mem_free(gdevs[0].second, rdma_qp->progress_gpu); From 5ec487f9e2bdabec58e35c83d2143526d537c582 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 09:52:08 +0800 Subject: [PATCH 16/25] test(gpunetio): report bulk-transfer tail alongside isolation latency Signed-off-by: Yunxiao Ning --- test/gtest/plugins/gpunetio/qp_progress_gtest.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu index 92a4c0bb0d..640b4a7d78 100644 --- a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu +++ b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu @@ -652,7 +652,9 @@ WritePerformanceJson(const Config &config, std::to_string(Percentile(b_windows, 0.50)) + ",\"b_post_to_done_p99_ns\":" + std::to_string(Percentile(b_windows, 0.99)) + ",\"a_post_to_done_p50_ns\":" + std::to_string(Percentile(a_windows, 0.50)) + + ",\"a_post_to_done_p99_ns\":" + std::to_string(Percentile(a_windows, 0.99)) + ",\"max_window_p50_ns\":" + std::to_string(Percentile(max_windows, 0.50)) + + ",\"max_window_p99_ns\":" + std::to_string(Percentile(max_windows, 0.99)) + ",\"actual_payload_bytes\":" + std::to_string(payload_bytes) + ",\"actual_marker_bytes\":" + std::to_string(marker_bytes) + ",\"actual_transfer_window_bytes\":" + std::to_string(payload_bytes + marker_bytes) + From bf280c1e02c9b4673543ff6f74c2edf851c94cb0 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 10:31:58 +0800 Subject: [PATCH 17/25] perf(gpunetio): localize progress state and trim request rearm Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 20 +++++++++++--------- 1 file changed, 11 insertions(+), 9 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index bb84e2870a..f9725ab3d1 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -276,9 +276,17 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) result = doca_gpu_mem_alloc(gdevs[0].second, sizeof(struct docaProgressState), 4096, - DOCA_GPU_MEM_TYPE_CPU_GPU, + DOCA_GPU_MEM_TYPE_GPU_CPU, (void **)&progress_state_gpu, (void **)&progress_state_cpu); + if (result != DOCA_SUCCESS || progress_state_gpu == nullptr || progress_state_cpu == nullptr) { + result = doca_gpu_mem_alloc(gdevs[0].second, + sizeof(struct docaProgressState), + 4096, + DOCA_GPU_MEM_TYPE_CPU_GPU, + (void **)&progress_state_gpu, + (void **)&progress_state_cpu); + } if (result != DOCA_SUCCESS || progress_state_gpu == nullptr || progress_state_cpu == nullptr) { throw std::runtime_error("Failed to allocate GPUNETIO progress state"); } @@ -1376,14 +1384,8 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, } auto &request = xferReqRingCpu[idx]; - request.generation++; - request.last_wqe = 0; - request.data_ticket = 0; - request.notif_wqe = 0; - request.notif_ticket = 0; - request.data_state = DOCA_XFER_DATA_NONE; - request.notif_state = DOCA_XFER_NOTIF_NONE; - treq->generations[i] = request.generation; + // Submission rewrites WQE/ticket fields before publishing their state. + request.generation = ++treq->generations[i]; std::atomic_ref(request.state) .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); } From f2b43128147d2bab0c4a9d17897e14c25d6c4cd3 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 10:31:58 +0800 Subject: [PATCH 18/25] docs(gpunetio): provide paired progress reproduction commands Signed-off-by: Yunxiao Ning --- test/gtest/plugins/gpunetio/README.md | 34 ++++++++++++++++++++++++--- 1 file changed, 31 insertions(+), 3 deletions(-) diff --git a/test/gtest/plugins/gpunetio/README.md b/test/gtest/plugins/gpunetio/README.md index 557f715554..11d7cc0534 100644 --- a/test/gtest/plugins/gpunetio/README.md +++ b/test/gtest/plugins/gpunetio/README.md @@ -12,6 +12,30 @@ an explicit GoogleTest skip before any CUDA query. ## Required environment +Build with CUDA, DOCA, GoogleTest/GoogleMock and the normal NIXL development +dependencies available. Use the same compiler options for both comparison arms. +The root build disables tests for `buildtype=release`, so use an optimized +`debugoptimized` build with assertions enabled: + +```bash +meson setup build --buildtype=debugoptimized -Doptimization=3 -Ddebug=false \ + -Db_ndebug=false -Denable_plugins=GPUNETIO -Dbuild_tests=true \ + -Dbuild_examples=false -Dbuild_docs=false -Dnixl_cuda_arch_list=90 \ + -Dcpp_args=-Wno-error=maybe-uninitialized --wrap-mode=forcefallback +ninja -C build src/core/libnixl.so src/plugins/gpunetio/libplugin_GPUNETIO.so \ + test/gtest/plugins/gpunetio/gpunetio_qp_progress_gtest +export QP_TEST="$PWD/build/test/gtest/plugins/gpunetio/gpunetio_qp_progress_gtest" +export NIXL_PLUGIN_DIR="$PWD/build/src/plugins/gpunetio" +export LD_LIBRARY_PATH="$(find "$PWD/build/src" -name 'lib*.so' -printf '%h\n' \ + | sort -u | paste -sd:):${LD_LIBRARY_PATH:-}" +git rev-parse HEAD +sha256sum "$QP_TEST" "$NIXL_PLUGIN_DIR/libplugin_GPUNETIO.so" +``` + +Architecture 90 and the GCC warning override describe the tested H20 build; +adapt these to the GPU/compiler in use. Include the DOCA library directory in +`LD_LIBRARY_PATH` if it is not already in the system loader configuration. + Both processes need `NIXL_PLUGIN_DIR` and `LD_LIBRARY_PATH` pointing to the same feature-tree build. Set the following values without placing site-specific values in this repository: @@ -38,11 +62,11 @@ or source target IPv4 is unavailable. ```bash # Terminal/host with two visible target GPUs. NIXL_QP_PROGRESS_ROLE=target \ - gpunetio_qp_progress_gtest --gtest_filter=QpProgress.PerformanceMixed2MiBAnd4KiB + "$QP_TEST" --gtest_filter=QpProgress.PerformanceMixed2MiBAnd4KiB # Terminal/host with one visible source GPU. Use the target's numeric IPv4. -NIXL_QP_PROGRESS_ROLE=source NIXL_QP_PROGRESS_TARGET_IPV4= \ - gpunetio_qp_progress_gtest --gtest_filter=QpProgress.PerformanceMixed2MiBAnd4KiB +NIXL_QP_PROGRESS_ROLE=source \ + "$QP_TEST" --gtest_filter=QpProgress.PerformanceMixed2MiBAnd4KiB ``` The performance case runs one 20-warmup/100-measured pair. Run three fresh @@ -50,6 +74,10 @@ source/target process pairs externally for independent repeats. It writes `qp_progress_performance.json` into the coordinate directory. The JSON reports p50/p99 transfer-window latency, actual payload and marker bytes, and separate wall time. It does not claim CQ timing or backend-internal timing. +Both endpoints must print `PASSED`, not `SKIPPED`. Keep all three JSON files per +arm. Compare medians of the three per-run metrics, reporting small-transfer and +bulk-transfer tails plus aggregate rate; an isolation win is not automatically +a throughput win. Profiled runs must not enter the performance comparison. `OutstandingDescriptorsAndAttachedStreamOrder` uses 513 descriptors per request (512 data plus an epoch marker) with descriptor merging disabled. It From 78724d7ba4cd262a7abe11a0550922d48e11b551 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 10:41:33 +0800 Subject: [PATCH 19/25] style(gpunetio): align control-memory fallback arguments Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index f9725ab3d1..de8f7480fb 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -281,11 +281,11 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) (void **)&progress_state_cpu); if (result != DOCA_SUCCESS || progress_state_gpu == nullptr || progress_state_cpu == nullptr) { result = doca_gpu_mem_alloc(gdevs[0].second, - sizeof(struct docaProgressState), - 4096, - DOCA_GPU_MEM_TYPE_CPU_GPU, - (void **)&progress_state_gpu, - (void **)&progress_state_cpu); + sizeof(struct docaProgressState), + 4096, + DOCA_GPU_MEM_TYPE_CPU_GPU, + (void **)&progress_state_gpu, + (void **)&progress_state_cpu); } if (result != DOCA_SUCCESS || progress_state_gpu == nullptr || progress_state_cpu == nullptr) { throw std::runtime_error("Failed to allocate GPUNETIO progress state"); From 3b7d578ccbf591e4fdbdb00ab470dfc2f133dfb5 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 13:31:16 +0800 Subject: [PATCH 20/25] perf(gpunetio): publish CPU-local generation-tagged completion records Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 107 +++++++++++++++----- src/plugins/gpunetio/gpunetio_backend.h | 7 ++ src/plugins/gpunetio/gpunetio_backend_aux.h | 11 ++ src/plugins/gpunetio/gpunetio_kernels.cu | 37 +++++-- 4 files changed, 126 insertions(+), 36 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index de8f7480fb..e097a04ff2 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -292,6 +292,18 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) } memset(progress_state_cpu, 0, sizeof(struct docaProgressState)); + result = doca_gpu_mem_alloc(gdevs[0].second, + sizeof(docaHostState), + 4096, + DOCA_GPU_MEM_TYPE_CPU_GPU, + (void **)&host_state_gpu_, + (void **)&host_state_cpu_); + if (result != DOCA_SUCCESS || host_state_gpu_ == nullptr || host_state_cpu_ == nullptr) { + throw std::runtime_error("Failed to allocate GPUNETIO host completion state"); + } + memset(host_state_cpu_, 0, sizeof(docaHostState)); + progress_state_cpu->host = host_state_gpu_; + std::atomic_thread_fence(std::memory_order_seq_cst); // DOCA_GPU_MEM_TYPE_GPU_CPU == GDRCopy result = doca_gpu_mem_alloc(gdevs[0].second, @@ -373,6 +385,7 @@ nixlDocaEngine::~nixlDocaEngine() { NIXL_DEBUG << "Before progressThreadStop "; progressThreadStop(); + stopped_.store(1, std::memory_order_release); std::atomic_ref(*wait_exit_cpu).store(1, std::memory_order_release); NIXL_DEBUG << "Before cudaStreamSynchronize "; nixlDocaEngineCheckCudaError(cudaStreamSynchronize(wait_stream), "stream synchronize"); @@ -390,6 +403,7 @@ nixlDocaEngine::~nixlDocaEngine() { qp_progress_gpu_.clear(); doca_gpu_mem_free(gdevs[0].second, xferReqRingGpu); doca_gpu_mem_free(gdevs[0].second, progress_state_gpu); + doca_gpu_mem_free(gdevs[0].second, host_state_gpu_); doca_gpu_mem_free(gdevs[0].second, wait_exit_gpu); NIXL_DEBUG << "Before nixlDocaDestroyNotif "; @@ -1219,6 +1233,7 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, rdma_qp = search->second; treq = new nixlDocaBckndReq; + treq->qp_data = rdma_qp->qp_data->get_qp_gpu_dev(); auto abandon_request = [&]() { for (uint32_t reserved_pos : treq->positions) { xferReqReserved_[reserved_pos].store(false, std::memory_order_release); @@ -1263,7 +1278,7 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, do { docaXferReqGpu staged_req{}; staged_req.has_notif_msg_idx = DOCA_NOTIF_NULL; - staged_req.generation = xferReqRingCpu[pos].generation + 1; + staged_req.generation = ++xferReqGenerations_[pos]; staged_req.state = DOCA_XFER_STATE_PREPARED; staged_req.data_state = DOCA_XFER_DATA_NONE; staged_req.notif_state = DOCA_XFER_NOTIF_NONE; @@ -1290,6 +1305,9 @@ nixlDocaEngine::prepXfer(const nixl_xfer_op_t &operation, std::atomic_ref(xferReqRingCpu[pos].state) .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); treq->generations.push_back(staged_req.generation); + host_state_cpu_->completions[pos].generation = staged_req.generation; + std::atomic_ref(host_state_cpu_->completions[pos].state) + .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); if (desc_offset < lcnt && !reserve_position()) { abandon_request(); @@ -1359,9 +1377,9 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, if (operation != NIXL_READ && operation != NIXL_WRITE) { return NIXL_ERR_INVALID_PARAM; } - if (std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != + if (std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) != 0 || - std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + stopped_.load(std::memory_order_acquire) != 0) { return NIXL_ERR_BACKEND; } if (treq->postStatus != NIXL_SUCCESS) { @@ -1377,15 +1395,20 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, if (completion_state == nixlDocaBckndReq::completion_state::COMPLETE) { for (size_t i = 0; i < treq->positions.size(); ++i) { const uint32_t idx = treq->positions[i]; - if (std::atomic_ref(xferReqRingCpu[idx].state) - .load(std::memory_order_acquire) != DOCA_XFER_STATE_COMPLETE) { + auto &completion = host_state_cpu_->completions[idx]; + if (std::atomic_ref(completion.state) + .load(std::memory_order_acquire) != DOCA_XFER_STATE_COMPLETE || + completion.generation != treq->generations[i]) { treq->postStatus = NIXL_ERR_BACKEND; return NIXL_ERR_BACKEND; } auto &request = xferReqRingCpu[idx]; // Submission rewrites WQE/ticket fields before publishing their state. - request.generation = ++treq->generations[i]; + request.generation = ++xferReqGenerations_[idx]; + treq->generations[i] = xferReqGenerations_[idx]; + std::atomic_ref(host_state_cpu_->completions[idx].state) + .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); std::atomic_ref(request.state) .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); } @@ -1398,24 +1421,27 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, for (uint32_t idx : treq->positions) { const doca_error_t result = operation == NIXL_READ ? doca_kernel_read(treq->stream, - xferReqRingCpu[idx].qp_data, + treq->qp_data, xferReqRingGpu, progress_state_gpu, wait_exit_gpu, idx) : doca_kernel_write(treq->stream, - xferReqRingCpu[idx].qp_data, + treq->qp_data, xferReqRingGpu, progress_state_gpu, wait_exit_gpu, idx); if (result != DOCA_SUCCESS) { treq->postStatus = NIXL_ERR_BACKEND; - std::atomic_ref(progress_state_cpu->failed) - .store(1, std::memory_order_release); + markFailed(); for (size_t pending = treq->postedCount; pending < treq->positions.size(); ++pending) { std::atomic_ref(xferReqRingCpu[treq->positions[pending]].state) .store(DOCA_XFER_STATE_ERROR, std::memory_order_release); + auto &completion = host_state_cpu_->completions[treq->positions[pending]]; + completion.generation = treq->generations[pending]; + std::atomic_ref(completion.state) + .store(DOCA_XFER_STATE_ERROR, std::memory_order_release); } break; } @@ -1430,7 +1456,7 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; auto state = treq->completionState.load(std::memory_order_acquire); if (state == nixlDocaBckndReq::completion_state::COMPLETE) { - return std::atomic_ref(progress_state_cpu->failed) + return std::atomic_ref(host_state_cpu_->failed) .load(std::memory_order_acquire) == 0 && treq->postStatus == NIXL_SUCCESS ? NIXL_SUCCESS : @@ -1444,23 +1470,25 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { bool request_error = false; for (size_t i = 0; i < treq->positions.size(); ++i) { const uint32_t idx = treq->positions[i]; + auto &completion = host_state_cpu_->completions[idx]; const uint32_t req_state = - std::atomic_ref(xferReqRingCpu[idx].state).load(std::memory_order_acquire); - if (xferReqRingCpu[idx].generation != treq->generations[i]) { - request_error = true; - continue; + std::atomic_ref(completion.state).load(std::memory_order_acquire); + if (req_state != DOCA_XFER_STATE_COMPLETE && req_state != DOCA_XFER_STATE_ERROR) { + return NIXL_IN_PROG; + } + if (completion.generation != treq->generations[i]) { + markFailed(); + treq->postStatus = NIXL_ERR_BACKEND; + return NIXL_ERR_BACKEND; } if (req_state == DOCA_XFER_STATE_ERROR) { request_error = true; continue; } - if (req_state != DOCA_XFER_STATE_COMPLETE) { - return NIXL_IN_PROG; - } } if (request_error || treq->postStatus != NIXL_SUCCESS || - std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != + std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) != 0) { treq->postStatus = NIXL_ERR_BACKEND; return NIXL_ERR_BACKEND; @@ -1470,6 +1498,13 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { return treq->postStatus; } +void +nixlDocaEngine::markFailed() const { + std::atomic_ref(progress_state_cpu->failed).store(1, std::memory_order_release); + std::atomic_thread_fence(std::memory_order_seq_cst); + std::atomic_ref(host_state_cpu_->failed).store(1, std::memory_order_release); +} + void nixlDocaEngine::retireRequest(nixlDocaBckndReq *request) const { auto state = request->completionState.load(std::memory_order_acquire); @@ -1513,6 +1548,17 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { if (status == NIXL_IN_PROG) { return NIXL_IN_PROG; } + if (status != NIXL_SUCCESS) { + for (size_t i = 0; i < treq->positions.size(); ++i) { + auto &completion = host_state_cpu_->completions[treq->positions[i]]; + const auto terminal = std::atomic_ref(completion.state) + .load(std::memory_order_acquire); + if ((terminal != DOCA_XFER_STATE_COMPLETE && terminal != DOCA_XFER_STATE_ERROR) || + completion.generation != treq->generations[i]) { + return NIXL_IN_PROG; + } + } + } } for (uint32_t idx : treq->positions) { @@ -1534,18 +1580,18 @@ nixlDocaEngine::getNotifs(notif_list_t ¬if_list) { // while getNotifs is running std::lock_guard lock(notifLock); for (auto ¬if : notifMap) { - if (std::atomic_ref(progress_state_cpu->failed).load(std::memory_order_acquire) != + if (std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) != 0 || - std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + stopped_.load(std::memory_order_acquire) != 0) { return NIXL_ERR_BACKEND; } ((volatile struct docaNotif *)notif_progress_cpu)->qp_gpu = qpMap[notif.first]->qp_notif->get_qp_gpu_dev(); std::atomic_thread_fence(std::memory_order_seq_cst); while (((volatile struct docaNotif *)notif_progress_cpu)->qp_gpu != nullptr) { - if (std::atomic_ref(progress_state_cpu->failed) + if (std::atomic_ref(host_state_cpu_->failed) .load(std::memory_order_acquire) != 0 || - std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + stopped_.load(std::memory_order_acquire) != 0) { return NIXL_ERR_BACKEND; } std::this_thread::yield(); @@ -1651,7 +1697,7 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg } docaXferReqGpu request{}; - request.generation = xferReqRingCpu[pos].generation + 1; + request.generation = ++xferReqGenerations_[pos]; request.state = DOCA_XFER_STATE_NOTIF_PENDING; request.notif_state = DOCA_XFER_NOTIF_PENDING; request.has_notif_msg_idx = buf_idx; @@ -1663,6 +1709,9 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg memcpy(&xferReqRingCpu[pos], &request, sizeof(request)); std::atomic_ref(xferReqRingCpu[pos].state) .store(DOCA_XFER_STATE_NOTIF_PENDING, std::memory_order_release); + host_state_cpu_->completions[pos].generation = request.generation; + std::atomic_ref(host_state_cpu_->completions[pos].state) + .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); const doca_error_t result = doca_kernel_publish_notif(post_stream[xferStream.fetch_add(1) & (nstreams - 1)], xferReqRingGpu, @@ -1676,8 +1725,14 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg } while (true) { + auto &completion = host_state_cpu_->completions[pos]; const uint32_t request_state = - std::atomic_ref(xferReqRingCpu[pos].state).load(std::memory_order_acquire); + std::atomic_ref(completion.state).load(std::memory_order_acquire); + if ((request_state == DOCA_XFER_STATE_COMPLETE || request_state == DOCA_XFER_STATE_ERROR) && + completion.generation != request.generation) { + markFailed(); + return NIXL_ERR_BACKEND; + } if (request_state == DOCA_XFER_STATE_COMPLETE) { xferReqReserved_[pos].store(false, std::memory_order_release); return NIXL_SUCCESS; @@ -1686,7 +1741,7 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg xferReqReserved_[pos].store(false, std::memory_order_release); return NIXL_ERR_BACKEND; } - if (std::atomic_ref(*wait_exit_cpu).load(std::memory_order_acquire) != 0) { + if (stopped_.load(std::memory_order_acquire) != 0) { return NIXL_ERR_BACKEND; } std::this_thread::yield(); diff --git a/src/plugins/gpunetio/gpunetio_backend.h b/src/plugins/gpunetio/gpunetio_backend.h index f95a58a3fe..fbaaa72d9d 100644 --- a/src/plugins/gpunetio/gpunetio_backend.h +++ b/src/plugins/gpunetio/gpunetio_backend.h @@ -164,6 +164,10 @@ class nixlDocaEngine : public nixlBackendEngine { struct docaXferReqGpu *xferReqRingCpu; mutable std::atomic xferRingPos; mutable std::array xferReqReserved_; + mutable std::array xferReqGenerations_{}; + std::atomic stopped_{0}; + docaHostState *host_state_cpu_ = nullptr; + docaHostState *host_state_gpu_ = nullptr; struct docaProgressState *progress_state_gpu; struct docaProgressState *progress_state_cpu; @@ -192,6 +196,7 @@ class nixlDocaEngine : public nixlBackendEngine { uint32_t devId; std::vector positions; std::vector generations; + doca_gpu_dev_verbs_qp *qp_data = nullptr; uintptr_t backendHandleGpu; size_t postedCount = 0; nixl_status_t postStatus = NIXL_SUCCESS; @@ -204,6 +209,8 @@ class nixlDocaEngine : public nixlBackendEngine { void retireRequest(nixlDocaBckndReq *request) const; + void + markFailed() const; nixl_status_t progressThreadStart(); void diff --git a/src/plugins/gpunetio/gpunetio_backend_aux.h b/src/plugins/gpunetio/gpunetio_backend_aux.h index c46240ad72..cc1ae62c9a 100644 --- a/src/plugins/gpunetio/gpunetio_backend_aux.h +++ b/src/plugins/gpunetio/gpunetio_backend_aux.h @@ -133,7 +133,18 @@ struct docaQpProgress { uint64_t head_notif_ticket; }; +struct docaHostCompletion { + uint32_t generation; + uint32_t state; +}; + +struct docaHostState { + uint32_t failed; + docaHostCompletion completions[DOCA_XFER_REQ_MAX]; +}; + struct docaProgressState { + docaHostState *host; uint32_t active_bitmap; uint32_t progress_cursor; uint32_t failed; diff --git a/src/plugins/gpunetio/gpunetio_kernels.cu b/src/plugins/gpunetio/gpunetio_kernels.cu index f485e8c4ae..5846adb3c0 100644 --- a/src/plugins/gpunetio/gpunetio_kernels.cu +++ b/src/plugins/gpunetio/gpunetio_kernels.cu @@ -117,29 +117,46 @@ nixl_gpunetio_dev_has_sq_credit(doca_gpu_dev_verbs_qp *qp, uint32_t count) { } __device__ void -nixl_gpunetio_dev_terminal_error(docaXferReqGpu *request, - docaProgressState *progress_state, - uint32_t pos) { +nixl_gpunetio_dev_mark_failed(docaProgressState *progress_state) { + nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); + nixl_gpunetio_dev_store_host_state(progress_state->host->failed, 1U); +} + +__device__ void +nixl_gpunetio_dev_finish_request(docaXferReqGpu *request, + docaProgressState *progress_state, + uint32_t pos, + uint32_t terminal) { + const uint32_t generation = request->generation; + auto *completion = &progress_state->host->completions[pos]; cuda::atomic_ref(progress_state->active_bitmap) .fetch_and(~(1U << pos), cuda::std::memory_order_release); - nixl_gpunetio_dev_store_host_state(request->state, DOCA_XFER_STATE_ERROR); + nixl_gpunetio_dev_store_host_state(request->state, terminal); + completion->generation = generation; + // Last publication: CPU may rearm/release the slot immediately afterwards. + nixl_gpunetio_dev_store_host_state(completion->state, terminal); +} + +__device__ void +nixl_gpunetio_dev_terminal_error(docaXferReqGpu *request, + docaProgressState *progress_state, + uint32_t pos) { + nixl_gpunetio_dev_finish_request(request, progress_state, pos, DOCA_XFER_STATE_ERROR); } __device__ void nixl_gpunetio_dev_fail_request(docaXferReqGpu *request, docaProgressState *progress_state, uint32_t pos) { + nixl_gpunetio_dev_mark_failed(progress_state); nixl_gpunetio_dev_terminal_error(request, progress_state, pos); - nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); } __device__ void nixl_gpunetio_dev_complete_request(docaXferReqGpu *request, docaProgressState *progress_state, uint32_t pos) { - cuda::atomic_ref(progress_state->active_bitmap) - .fetch_and(~(1U << pos), cuda::std::memory_order_release); - nixl_gpunetio_dev_store_host_state(request->state, DOCA_XFER_STATE_COMPLETE); + nixl_gpunetio_dev_finish_request(request, progress_state, pos, DOCA_XFER_STATE_COMPLETE); } __device__ bool @@ -434,7 +451,7 @@ kernel_progress(struct docaXferReqGpu *xfer_req_ring, docaXferReqGpu *request = &xfer_req_ring[pos]; if (progress_state->active_generation[pos] != request->generation) { // Do not release a possibly newer owner on a stale generation. - nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); + nixl_gpunetio_dev_mark_failed(progress_state); continue; } docaQpProgress *progress = request->qp_progress; @@ -539,7 +556,7 @@ kernel_progress(struct docaXferReqGpu *xfer_req_ring, doca_gpu_dev_verbs_fence_release(); DOCA_GPUNETIO_VOLATILE(notif_progress->qp_gpu) = nullptr; } else { - nixl_gpunetio_dev_store_host_state(progress_state->failed, 1U); + nixl_gpunetio_dev_mark_failed(progress_state); DOCA_GPUNETIO_VOLATILE(notif_progress->qp_gpu) = nullptr; } } From d82a22f767f074bbf76217941739d13b6c046925 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 13:39:18 +0800 Subject: [PATCH 21/25] fix(gpunetio): preserve terminal ownership on host-side failure paths Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 73 ++++++++++++++--------- src/plugins/gpunetio/gpunetio_kernels.cu | 10 ++-- 2 files changed, 50 insertions(+), 33 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index e097a04ff2..d6848fa656 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -293,11 +293,11 @@ nixlDocaEngine::nixlDocaEngine(const nixlBackendInitParams *init_params) memset(progress_state_cpu, 0, sizeof(struct docaProgressState)); result = doca_gpu_mem_alloc(gdevs[0].second, - sizeof(docaHostState), - 4096, - DOCA_GPU_MEM_TYPE_CPU_GPU, - (void **)&host_state_gpu_, - (void **)&host_state_cpu_); + sizeof(docaHostState), + 4096, + DOCA_GPU_MEM_TYPE_CPU_GPU, + (void **)&host_state_gpu_, + (void **)&host_state_cpu_); if (result != DOCA_SUCCESS || host_state_gpu_ == nullptr || host_state_cpu_ == nullptr) { throw std::runtime_error("Failed to allocate GPUNETIO host completion state"); } @@ -1377,8 +1377,7 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, if (operation != NIXL_READ && operation != NIXL_WRITE) { return NIXL_ERR_INVALID_PARAM; } - if (std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) != - 0 || + if (std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) != 0 || stopped_.load(std::memory_order_acquire) != 0) { return NIXL_ERR_BACKEND; } @@ -1396,8 +1395,8 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, for (size_t i = 0; i < treq->positions.size(); ++i) { const uint32_t idx = treq->positions[i]; auto &completion = host_state_cpu_->completions[idx]; - if (std::atomic_ref(completion.state) - .load(std::memory_order_acquire) != DOCA_XFER_STATE_COMPLETE || + if (std::atomic_ref(completion.state).load(std::memory_order_acquire) != + DOCA_XFER_STATE_COMPLETE || completion.generation != treq->generations[i]) { treq->postStatus = NIXL_ERR_BACKEND; return NIXL_ERR_BACKEND; @@ -1409,6 +1408,7 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, treq->generations[i] = xferReqGenerations_[idx]; std::atomic_ref(host_state_cpu_->completions[idx].state) .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); + completion.generation = treq->generations[i]; std::atomic_ref(request.state) .store(DOCA_XFER_STATE_PREPARED, std::memory_order_release); } @@ -1419,19 +1419,18 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, treq->completionState.store(nixlDocaBckndReq::completion_state::IN_PROGRESS, std::memory_order_release); for (uint32_t idx : treq->positions) { - const doca_error_t result = operation == NIXL_READ ? - doca_kernel_read(treq->stream, - treq->qp_data, - xferReqRingGpu, - progress_state_gpu, - wait_exit_gpu, - idx) : - doca_kernel_write(treq->stream, - treq->qp_data, - xferReqRingGpu, - progress_state_gpu, - wait_exit_gpu, - idx); + const doca_error_t result = operation == NIXL_READ ? doca_kernel_read(treq->stream, + treq->qp_data, + xferReqRingGpu, + progress_state_gpu, + wait_exit_gpu, + idx) : + doca_kernel_write(treq->stream, + treq->qp_data, + xferReqRingGpu, + progress_state_gpu, + wait_exit_gpu, + idx); if (result != DOCA_SUCCESS) { treq->postStatus = NIXL_ERR_BACKEND; markFailed(); @@ -1456,8 +1455,8 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { nixlDocaBckndReq *treq = (nixlDocaBckndReq *)handle; auto state = treq->completionState.load(std::memory_order_acquire); if (state == nixlDocaBckndReq::completion_state::COMPLETE) { - return std::atomic_ref(host_state_cpu_->failed) - .load(std::memory_order_acquire) == 0 && + return std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) == + 0 && treq->postStatus == NIXL_SUCCESS ? NIXL_SUCCESS : NIXL_ERR_BACKEND; @@ -1488,8 +1487,7 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { } if (request_error || treq->postStatus != NIXL_SUCCESS || - std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) != - 0) { + std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) != 0) { treq->postStatus = NIXL_ERR_BACKEND; return NIXL_ERR_BACKEND; } @@ -1535,6 +1533,16 @@ nixl_status_t nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { auto *treq = static_cast(handle); if (treq->postedCount == 0) { + // Prepared-only handles, including failures before the first enqueue. + for (size_t i = 0; i < treq->positions.size(); ++i) { + auto &completion = host_state_cpu_->completions[treq->positions[i]]; + const auto state = + std::atomic_ref(completion.state).load(std::memory_order_acquire); + if ((state != DOCA_XFER_STATE_PREPARED && state != DOCA_XFER_STATE_ERROR) || + completion.generation != treq->generations[i]) { + return NIXL_IN_PROG; + } + } for (uint32_t idx : treq->positions) { xferReqReserved_[idx].store(false, std::memory_order_release); } @@ -1551,8 +1559,8 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { if (status != NIXL_SUCCESS) { for (size_t i = 0; i < treq->positions.size(); ++i) { auto &completion = host_state_cpu_->completions[treq->positions[i]]; - const auto terminal = std::atomic_ref(completion.state) - .load(std::memory_order_acquire); + const auto terminal = + std::atomic_ref(completion.state).load(std::memory_order_acquire); if ((terminal != DOCA_XFER_STATE_COMPLETE && terminal != DOCA_XFER_STATE_ERROR) || completion.generation != treq->generations[i]) { return NIXL_IN_PROG; @@ -1642,6 +1650,10 @@ nixlDocaEngine::getNotifs(notif_list_t ¬if_list) { nixl_status_t nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg) const { + if (std::atomic_ref(host_state_cpu_->failed).load(std::memory_order_acquire) != 0 || + stopped_.load(std::memory_order_acquire) != 0) { + return NIXL_ERR_BACKEND; + } struct nixlDocaNotif *notif; uint32_t buf_idx; uint32_t pos = 0; @@ -1718,8 +1730,12 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg progress_state_gpu, pos); if (result != DOCA_SUCCESS) { + markFailed(); std::atomic_ref(xferReqRingCpu[pos].state) .store(DOCA_XFER_STATE_ERROR, std::memory_order_release); + host_state_cpu_->completions[pos].generation = request.generation; + std::atomic_ref(host_state_cpu_->completions[pos].state) + .store(DOCA_XFER_STATE_ERROR, std::memory_order_release); xferReqReserved_[pos].store(false, std::memory_order_release); return NIXL_ERR_BACKEND; } @@ -1742,6 +1758,7 @@ nixlDocaEngine::genNotif(const std::string &remote_agent, const std::string &msg return NIXL_ERR_BACKEND; } if (stopped_.load(std::memory_order_acquire) != 0) { + // Only teardown sets stopped_; retain ownership until its stream drain. return NIXL_ERR_BACKEND; } std::this_thread::yield(); diff --git a/src/plugins/gpunetio/gpunetio_kernels.cu b/src/plugins/gpunetio/gpunetio_kernels.cu index 5846adb3c0..3f0fabe3c4 100644 --- a/src/plugins/gpunetio/gpunetio_kernels.cu +++ b/src/plugins/gpunetio/gpunetio_kernels.cu @@ -124,9 +124,9 @@ nixl_gpunetio_dev_mark_failed(docaProgressState *progress_state) { __device__ void nixl_gpunetio_dev_finish_request(docaXferReqGpu *request, - docaProgressState *progress_state, - uint32_t pos, - uint32_t terminal) { + docaProgressState *progress_state, + uint32_t pos, + uint32_t terminal) { const uint32_t generation = request->generation; auto *completion = &progress_state->host->completions[pos]; cuda::atomic_ref(progress_state->active_bitmap) @@ -139,8 +139,8 @@ nixl_gpunetio_dev_finish_request(docaXferReqGpu *request, __device__ void nixl_gpunetio_dev_terminal_error(docaXferReqGpu *request, - docaProgressState *progress_state, - uint32_t pos) { + docaProgressState *progress_state, + uint32_t pos) { nixl_gpunetio_dev_finish_request(request, progress_state, pos, DOCA_XFER_STATE_ERROR); } From 4c25f8d7b3174243c187bf5116d30bc6dcf08ea3 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 13:54:12 +0800 Subject: [PATCH 22/25] test(gpunetio): verify CPU fatal errors retain pending request ownership Signed-off-by: Yunxiao Ning --- .../plugins/gpunetio/qp_progress_gtest.cu | 162 ++++++++++++++++++ 1 file changed, 162 insertions(+) diff --git a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu index 640b4a7d78..a105f246f0 100644 --- a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu +++ b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu @@ -1230,6 +1230,144 @@ SourceFault(const Config &config) { control.Send('E'); } +enum class PreErrorAction { PostTransfer, GenerateNotification }; + +void +TargetPreError(const Config &config) { + constexpr uint64_t kSentinelRound = 0x5100; + DeviceMemory a_memory(kSmallBytes, 0); + DeviceMemory b_memory(kSmallBytes, 1); + cudaStream_t a_verify_stream = nullptr; + cudaStream_t b_verify_stream = nullptr; + CheckCuda(cudaSetDevice(0), "select pre-error target A device"); + CheckCuda(cudaStreamCreateWithFlags(&a_verify_stream, cudaStreamNonBlocking), + "create pre-error target A verify stream"); + CheckCuda(cudaSetDevice(1), "select pre-error target B device"); + CheckCuda(cudaStreamCreateWithFlags(&b_verify_stream, cudaStreamNonBlocking), + "create pre-error target B verify stream"); + VerifyCounter a_counter(0); + VerifyCounter b_counter(1); + Fill(a_memory, SeedA(kSentinelRound), EpochA(kSentinelRound), a_verify_stream); + Fill(b_memory, SeedB(kSentinelRound), EpochB(kSentinelRound), b_verify_stream); + CheckCuda(cudaSetDevice(0), "select pre-error target A prepare device"); + CheckCuda(cudaStreamSynchronize(a_verify_stream), "prepare pre-error target A sentinel"); + CheckCuda(cudaSetDevice(1), "select pre-error target B prepare device"); + CheckCuda(cudaStreamSynchronize(b_verify_stream), "prepare pre-error target B sentinel"); + + // All target GPU allocations, verification state, and streams precede the persistent engines. + Endpoint a(config, "qp-progress-a", config.target_a_port, a_memory); + Endpoint b(config, "qp-progress-b", config.target_b_port, b_memory); + Control control(config); + PublishMetadata(config, a, a_memory, b, b_memory); + control.Send('R'); + control.Expect('D'); + ASSERT_EQ( + a_counter.Verify(a_memory, SeedA(kSentinelRound), EpochA(kSentinelRound), a_verify_stream), + 0U); + ASSERT_EQ( + b_counter.Verify(b_memory, SeedB(kSentinelRound), EpochB(kSentinelRound), b_verify_stream), + 0U); + control.Send('V'); + control.Expect('C'); + control.Send('K'); + CheckCuda(cudaSetDevice(0), "select pre-error target A destroy device"); + CheckCuda(cudaStreamSynchronize(a_verify_stream), "drain pre-error target A verify stream"); + CheckCuda(cudaStreamDestroy(a_verify_stream), "destroy pre-error target A verify stream"); + CheckCuda(cudaSetDevice(1), "select pre-error target B destroy device"); + CheckCuda(cudaStreamSynchronize(b_verify_stream), "drain pre-error target B verify stream"); + CheckCuda(cudaStreamDestroy(b_verify_stream), "destroy pre-error target B verify stream"); +} + +void +SourcePreError(const Config &config, PreErrorAction action) { + constexpr uint64_t kSourceRound = 0x5200; + DeviceMemory a_memory(kSmallBytes, 0); + DeviceMemory b_memory(kSmallBytes, 0); + cudaStream_t delayed_stream = nullptr; + cudaStream_t fast_stream = nullptr; + CheckCuda(cudaStreamCreateWithFlags(&delayed_stream, cudaStreamNonBlocking), + "create pre-error delayed stream"); + CheckCuda(cudaStreamCreateWithFlags(&fast_stream, cudaStreamNonBlocking), + "create pre-error fast stream"); + nixl_reg_dlist_t b_registration = Registration(b_memory); + + // Both attached streams and all source GPU memory precede the persistent engine. + Endpoint source(config, "qp-progress-source", config.source_port, a_memory); + CheckNixl(source.agent.registerMem(b_registration, &source.options), + "register pre-error source B memory"); + Control control(config); + control.Expect('R'); + const RemoteAddresses remote = LoadMetadata(config, source.agent); + + CheckCuda(cudaSetDevice(0), "select pre-error source device"); + Fill(b_memory, SeedB(kSourceRound), EpochB(kSourceRound), fast_stream); + CheckCuda(cudaStreamSynchronize(fast_stream), "prepare pre-error B payload"); + + // A's transfer kernel is queued behind a test-only delay. The transfer variant prepares B + // before the deliberate CUDA launch-configuration error is left uncleared for + // doca_kernel_write. + DelayKernel<<<1, 1, 0, delayed_stream>>>(500000000ULL); + CheckCuda(cudaGetLastError(), "launch pre-error delay kernel"); + Fill(a_memory, SeedA(kSourceRound), EpochA(kSourceRound), delayed_stream); + + nixlXferReqH *a_request = nullptr; + nixlXferReqH *b_request = nullptr; + const auto a_options = Attached(source.backend, delayed_stream, std::nullopt); + const auto b_options = Attached(source.backend, fast_stream, std::nullopt); + CheckNixl(source.agent.createXferReq(NIXL_WRITE, + TransferList(a_memory), + RemoteList(remote.a_data, kSmallBytes, remote.a_epoch, 0), + "qp-progress-a", + a_request, + &a_options), + "create pre-error A request"); + if (action == PreErrorAction::PostTransfer) { + CheckNixl( + source.agent.createXferReq(NIXL_WRITE, + TransferList(b_memory), + RemoteList(remote.b_data, kSmallBytes, remote.b_epoch, 1), + "qp-progress-b", + b_request, + &b_options), + "prepare pre-error B request"); + } + const nixl_status_t a_post = source.agent.postXferReq(a_request); + ASSERT_TRUE(a_post == NIXL_SUCCESS || a_post == NIXL_IN_PROG); + + CheckCuda(cudaGetLastError(), "clear CUDA error before injection"); + DelayKernel<<<0, 1>>>(1); + ASSERT_EQ(cudaPeekAtLastError(), cudaErrorInvalidConfiguration); + + if (action == PreErrorAction::PostTransfer) { + ASSERT_EQ(source.agent.postXferReq(b_request), NIXL_ERR_BACKEND); + ASSERT_EQ(source.agent.releaseXferReq(b_request), NIXL_SUCCESS); + } else { + ASSERT_EQ(source.agent.genNotif("qp-progress-b", "pre-error", &source.options), + NIXL_ERR_BACKEND); + } + + // CPU ownership is retained while A is still behind DelayKernel; only the GPU terminal + // publication makes its slot releasable. + ASSERT_EQ(source.agent.releaseXferReq(a_request), NIXL_IN_PROG); + CheckCuda(cudaStreamSynchronize(delayed_stream), "drain pre-error delayed stream"); + ASSERT_EQ(Wait(source.agent, a_request, std::chrono::seconds(10)), NIXL_ERR_BACKEND); + ASSERT_EQ(source.agent.releaseXferReq(a_request), NIXL_SUCCESS); + CheckCuda(cudaStreamSynchronize(fast_stream), "drain pre-error attached stream"); + + control.Send('D'); + control.Expect('V'); + CheckNixl(source.agent.invalidateRemoteMD("qp-progress-a"), + "invalidate pre-error target A metadata"); + CheckNixl(source.agent.invalidateRemoteMD("qp-progress-b"), + "invalidate pre-error target B metadata"); + CheckNixl(source.agent.deregisterMem(b_registration, &source.options), + "deregister pre-error source B memory"); + CheckCuda(cudaStreamDestroy(fast_stream), "destroy pre-error fast stream"); + CheckCuda(cudaStreamDestroy(delayed_stream), "destroy pre-error delayed stream"); + control.Send('C'); + control.Expect('K'); +} + TEST(QpProgress, PerformanceMixed2MiBAnd4KiB) { const Config config = GetConfig(); if (const auto problem = ConfigProblem(config, false)) { @@ -1278,6 +1416,30 @@ TEST(QpProgress, RemoteDeregisterReturnsBackendError) { } } +TEST(QpProgress, CpuFatalLatchRejectsQueuedTransfer) { + const Config config = GetConfig(); + if (const auto problem = ConfigProblem(config, false)) { + GTEST_SKIP() << *problem; + } + if (config.role == "source") { + SourcePreError(config, PreErrorAction::PostTransfer); + } else { + TargetPreError(config); + } +} + +TEST(QpProgress, CpuFatalLatchRejectsQueuedNotification) { + const Config config = GetConfig(); + if (const auto problem = ConfigProblem(config, false)) { + GTEST_SKIP() << *problem; + } + if (config.role == "source") { + SourcePreError(config, PreErrorAction::GenerateNotification); + } else { + TargetPreError(config); + } +} + } // namespace int From ebdb09ad760947e109074ab895284306d929844e Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 14:06:30 +0800 Subject: [PATCH 23/25] fix(gpunetio): reject cancellation until owned work is terminal Signed-off-by: Yunxiao Ning --- src/plugins/gpunetio/gpunetio_backend.cpp | 12 +++++++----- src/plugins/gpunetio/gpunetio_kernels.cu | 1 - test/gtest/plugins/gpunetio/qp_progress_gtest.cu | 2 +- 3 files changed, 8 insertions(+), 7 deletions(-) diff --git a/src/plugins/gpunetio/gpunetio_backend.cpp b/src/plugins/gpunetio/gpunetio_backend.cpp index d6848fa656..25f832e26e 100644 --- a/src/plugins/gpunetio/gpunetio_backend.cpp +++ b/src/plugins/gpunetio/gpunetio_backend.cpp @@ -1447,7 +1447,9 @@ nixlDocaEngine::postXfer(const nixl_xfer_op_t &operation, ++treq->postedCount; } - return treq->postStatus == NIXL_SUCCESS ? NIXL_IN_PROG : treq->postStatus; + // A later launch can fail while earlier chunks still own GPU/NIC work. + // Keep the frontend handle in progress until checkXfer has drained them. + return treq->postedCount != 0 ? NIXL_IN_PROG : treq->postStatus; } nixl_status_t @@ -1478,7 +1480,7 @@ nixlDocaEngine::checkXfer(nixlBackendReqH *handle) const { if (completion.generation != treq->generations[i]) { markFailed(); treq->postStatus = NIXL_ERR_BACKEND; - return NIXL_ERR_BACKEND; + return NIXL_IN_PROG; } if (req_state == DOCA_XFER_STATE_ERROR) { request_error = true; @@ -1540,7 +1542,7 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { std::atomic_ref(completion.state).load(std::memory_order_acquire); if ((state != DOCA_XFER_STATE_PREPARED && state != DOCA_XFER_STATE_ERROR) || completion.generation != treq->generations[i]) { - return NIXL_IN_PROG; + return NIXL_ERR_REPOST_ACTIVE; } } for (uint32_t idx : treq->positions) { @@ -1554,7 +1556,7 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { nixlDocaBckndReq::completion_state::COMPLETE) { nixl_status_t status = checkXfer(handle); if (status == NIXL_IN_PROG) { - return NIXL_IN_PROG; + return NIXL_ERR_REPOST_ACTIVE; } if (status != NIXL_SUCCESS) { for (size_t i = 0; i < treq->positions.size(); ++i) { @@ -1563,7 +1565,7 @@ nixlDocaEngine::releaseReqH(nixlBackendReqH *handle) const { std::atomic_ref(completion.state).load(std::memory_order_acquire); if ((terminal != DOCA_XFER_STATE_COMPLETE && terminal != DOCA_XFER_STATE_ERROR) || completion.generation != treq->generations[i]) { - return NIXL_IN_PROG; + return NIXL_ERR_REPOST_ACTIVE; } } } diff --git a/src/plugins/gpunetio/gpunetio_kernels.cu b/src/plugins/gpunetio/gpunetio_kernels.cu index 3f0fabe3c4..a85bfad29d 100644 --- a/src/plugins/gpunetio/gpunetio_kernels.cu +++ b/src/plugins/gpunetio/gpunetio_kernels.cu @@ -131,7 +131,6 @@ nixl_gpunetio_dev_finish_request(docaXferReqGpu *request, auto *completion = &progress_state->host->completions[pos]; cuda::atomic_ref(progress_state->active_bitmap) .fetch_and(~(1U << pos), cuda::std::memory_order_release); - nixl_gpunetio_dev_store_host_state(request->state, terminal); completion->generation = generation; // Last publication: CPU may rearm/release the slot immediately afterwards. nixl_gpunetio_dev_store_host_state(completion->state, terminal); diff --git a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu index a105f246f0..52c1a375b4 100644 --- a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu +++ b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu @@ -1348,7 +1348,7 @@ SourcePreError(const Config &config, PreErrorAction action) { // CPU ownership is retained while A is still behind DelayKernel; only the GPU terminal // publication makes its slot releasable. - ASSERT_EQ(source.agent.releaseXferReq(a_request), NIXL_IN_PROG); + ASSERT_EQ(source.agent.releaseXferReq(a_request), NIXL_ERR_REPOST_ACTIVE); CheckCuda(cudaStreamSynchronize(delayed_stream), "drain pre-error delayed stream"); ASSERT_EQ(Wait(source.agent, a_request, std::chrono::seconds(10)), NIXL_ERR_BACKEND); ASSERT_EQ(source.agent.releaseXferReq(a_request), NIXL_SUCCESS); From b260dd160455031a2fe61d419a228a98c674de69 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 14:11:50 +0800 Subject: [PATCH 24/25] test(gpunetio): vary payload bytes across epochs and descriptor offsets Signed-off-by: Yunxiao Ning --- test/gtest/plugins/gpunetio/qp_progress_gtest.cu | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu index 52c1a375b4..bbff31b13a 100644 --- a/test/gtest/plugins/gpunetio/qp_progress_gtest.cu +++ b/test/gtest/plugins/gpunetio/qp_progress_gtest.cu @@ -262,7 +262,12 @@ ReadFile(const fs::path &path) { __device__ uint8_t Pattern(uint64_t seed, size_t offset) { - return static_cast((seed + offset * 1315423911ULL + (offset >> 7U) * 17ULL) & 0xffU); + // Mix the full epoch/source seed and offset: stale payloads and misplaced + // chunks must not alias merely because their low eight bits are equal. + uint64_t value = seed + (offset + 1) * 0x9e3779b97f4a7c15ULL; + value = (value ^ (value >> 30)) * 0xbf58476d1ce4e5b9ULL; + value = (value ^ (value >> 27)) * 0x94d049bb133111ebULL; + return static_cast(value ^ (value >> 31)); } __global__ void From 292249b8fd099cad4224139f04dd82de1ed2d092 Mon Sep 17 00:00:00 2001 From: Yunxiao Ning Date: Tue, 8 Sep 2026 14:31:45 +0800 Subject: [PATCH 25/25] docs(gpunetio): describe CPU-error ownership and fresh-payload checks Signed-off-by: Yunxiao Ning --- test/gtest/plugins/gpunetio/README.md | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/test/gtest/plugins/gpunetio/README.md b/test/gtest/plugins/gpunetio/README.md index 11d7cc0534..4acb0c918c 100644 --- a/test/gtest/plugins/gpunetio/README.md +++ b/test/gtest/plugins/gpunetio/README.md @@ -59,6 +59,10 @@ process for each fault trial (`target-fault` plus `source` filtered to the fault case). The harness skips if its role, coordinate directory, GPU count, or source target IPv4 is unavailable. +Payload generation mixes the full epoch/source seed and byte offset, so both +payload bytes and markers change between iterations. Generation and validation +remain outside the measured post-to-completion windows. + ```bash # Terminal/host with two visible target GPUs. NIXL_QP_PROGRESS_ROLE=target \ @@ -99,3 +103,12 @@ Set `NIXL_QP_PROGRESS_CONTROL_BYTES=2097152` to repeat the same control at 2 MiB `RemoteDeregisterReturnsBackendError` requires its fresh fault role and checks bounded `NIXL_ERR_BACKEND` and release after target B memory is deregistered after source metadata consumption and before its post. + +`CpuFatalLatchRejectsQueuedTransfer` and `CpuFatalLatchRejectsQueuedNotification` +use fresh normal `source`/`target` roles. A deliberate CUDA configuration error +rejects B before enqueue while A remains behind a test-only stream delay. +Pending cancellation must return `NIXL_ERR_REPOST_ACTIVE` without losing A's +handle; after GPU error completion, release succeeds. Target sentinel payloads +and epochs must remain unchanged. These are CPU pre-enqueue failure tests, not +arbitrary device-fault injection. Their expected error logs are not benchmark +failures when the assertions and both process exits pass.