Repository navigation
support gpu direct rdma - #3144
randomkang wants to merge 9 commits into
Conversation
There was a problem hiding this comment.
Pull Request Overview
This PR adds GPU Direct RDMA (GDR) support to the BRPC framework, enabling efficient data transfer between GPU memory and RDMA-capable network devices. The implementation includes GPU memory pool management, CUDA stream pooling for asynchronous memory operations, and protocol-level changes to handle GPU-resident data.
- Added GPU memory block pool allocator with configurable block sizes
- Implemented CUDA stream pooling for optimized device-to-host and device-to-device transfers
- Integrated GDR support into IOBuf, RDMA endpoint, and RPC protocol parsing
Reviewed Changes
Copilot reviewed 11 out of 11 changed files in this pull request and generated 11 comments.
Show a summary per file
| File | Description |
|---|---|
| src/butil/iobuf.h | Added GPU-specific copy and cut methods with conditional compilation guards |
| src/butil/iobuf.cpp | Implemented cutn_from_gpu and copy_from_gpu methods for GPU memory operations |
| src/butil/gpu/gpu_block_pool.h | New header defining GPU memory pool allocator, stream pool, and region management |
| src/butil/gpu/gpu_block_pool.cpp | Implementation of GPU memory allocation, CUDA stream management, and memory copy operations |
| src/brpc/rdma/rdma_helper.h | Added function declaration for GPU index retrieval |
| src/brpc/rdma/rdma_helper.cpp | Added GPU index configuration and initialization logic for GDR block pool |
| src/brpc/rdma/rdma_endpoint.h | Added remote receive window tracking and GDR-specific receive posting method |
| src/brpc/rdma/rdma_endpoint.cpp | Modified flow control, completion handling, and receive buffer posting for GDR support |
| src/brpc/policy/baidu_rpc_protocol.cpp | Enhanced message parsing to detect and handle GPU memory with optimized prefetch |
| bazel/config/BUILD.bazel | Added build configuration for GDR feature flag |
| BUILD.bazel | Added CUDA dependencies and GPU block pool source to build |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| clock_gettime(CLOCK_MONOTONIC, &end); | ||
| double time_us = (end.tv_sec - start.tv_sec) * 1e6 + (end.tv_nsec - start.tv_nsec) / 1e3; |
There was a problem hiding this comment.
The variables start, end, and time_us are defined and computed but the result is only used in a commented-out log statement (lines 766-767). Consider removing this timing code if it's not being used, or uncomment the logging if it's needed for debugging.
| double time_us = (end.tv_sec - start.tv_sec) * 1e6 + (end.tv_nsec - start.tv_nsec) / 1e3; | ||
| size_t copied_bytes = n - m; | ||
|
|
||
| // LOG(INFO) << "GDRCopy: " << copied_bytes << " bytes, " | ||
| // << time_us << " us" << ", to_gpu " << to_gpu; |
There was a problem hiding this comment.
The variables time_us and copied_bytes are computed but only used in commented-out log statements (lines 1421-1422). Consider removing this unused code or uncommenting the logging if it's needed for debugging.
| double time_us = (end.tv_sec - start.tv_sec) * 1e6 + (end.tv_nsec - start.tv_nsec) / 1e3; | |
| size_t copied_bytes = n - m; | |
| // LOG(INFO) << "GDRCopy: " << copied_bytes << " bytes, " | |
| // << time_us << " us" << ", to_gpu " << to_gpu; |
| struct timespec start, end; | ||
| clock_gettime(CLOCK_MONOTONIC, &start); |
There was a problem hiding this comment.
The timing variables start and end are declared and used but the computed time_us value is only for commented-out logging. If timing is not actively needed, consider removing this instrumentation code.
| struct timespec start, end; | ||
| clock_gettime(CLOCK_MONOTONIC, &start); |
There was a problem hiding this comment.
The timing variables start and end are declared and used but the computed time_us value is only for commented-out logging. If timing is not actively needed, consider removing this instrumentation code.
|
|
||
| char header_buf[12]; | ||
| const size_t n = source->copy_to(header_buf, sizeof(header_buf)); | ||
| size_t n = 0; |
There was a problem hiding this comment.
I think it would be better to put the baidu_rpc_protocol.cpp file containing GPU functionality into a separate file called baidu_rpc_with_gpu_protocol.cpp, or to define a new protocol. Currently, the macro definitions generate too many if-else statements.
There was a problem hiding this comment.
Maybe put GPU functionality into a separate file or define a new protocol is two heavy, we reorg the code to decrease the macro definitions. Is it ok? @yanglimingcn
There was a problem hiding this comment.
I think it's better to separate new protocol file for GDR feature， Besides, I have a idea we can also define a new protocol meta that includes both request meta and response meta carrying the GDR attribute value., the corresponding meta should also contain the GPU address information that both sides need to register with IB. Doing the GDR operations inside processRequest and ProcessRpcResponse method step will probably be more reasonable.
There was a problem hiding this comment.
If the server or client wants to know the GPU memory region of the other side, we can send this information with a separate rpc transport. But i do not see the the benefits of this approach. Can you explain more?
|
I believe the long-term solution is to have users register memory with RDMA, allowing users to customize this memory based on their data organization methods. |
In this pr, brpc will recv data into fragmented gpu blocks and user must call IOBuf::copy_from_gpu to copy these gpu blocks into a continous hbm in order to use it. The d2d copy is time-consuming and not necessary. In feature, we can let the user assign the gpu destination directly with interface like "rdma_memory_pool_user_specified_memory" and recv the data into it with the opcode like IBV_WR_RDMA_READ/IBV_WR_RDMA_WRITE. Then we can skip the d2d copy. Futhermore, we can parse the brpc protocol with gpu kernel and initiate RDMA communication with GPU. Then the control path and the data path both are on gpu like nccl gin. |
041bef0 to
9e3b2af
Compare
|
@yanglimingcn @zchuango all comments have been resolved, please review again. |
|
Turn multiple commits into one |
ok |
|
LGTM |
| return (void *)d_data; | ||
| } | ||
|
|
||
| void* get_cpu_mem(int gpu_id, int64_t cpu_mem_size) { |
There was a problem hiding this comment.
是的。这里申请的是pinned memory，使用pinned memory和显存进行交换，会比普通的memory快。在使用cudaMallocHost分配pinned memory的时候，建议先执行cudaSetDevice，这样分配的pined memory就会与对应的gpu id关联。
|
|
||
| std::atomic<BlockPoolAllocators*> BlockPoolAllocators::instance_{nullptr}; | ||
|
|
||
| BlockPoolAllocators* BlockPoolAllocators::singleton() { |
There was a problem hiding this comment.
你看brpc自带的singleton.h是否能满足你的需求。
There was a problem hiding this comment.
我看了下，可以用brpc自带的Singleton，已改
|
|
||
| void* IOBuf::get_first_data_ptr() { | ||
| if (_ref_num() == 0) { | ||
| return 0; |
|
@yanglimingcn @zchuango Are there any other issues with this pr? |
|
All Copilot issues have been resolved or commented. @chenBright @wwbmmm |
|
please resolve merge conflicts |
wwbmmm
left a comment
wwbmmm
left a comment
There was a problem hiding this comment.
[replied by brpc-oncall robot] Most previous findings are now correctly addressed (prefetch size derived from the CPU pool block size, bool-returning GPU buffer fills with error propagation, cut/append failure checks, DEBUG-only timing, GetGPUIndex in the example). Two items remain: the GDR init ordering in the example client (below) and the still-unaddressed IOBuf::copy_from_gpu() device-memory validation; the Bazel GDR build also still relies on ambient C_INCLUDE_PATH/LD_LIBRARY_PATH rather than declaring the CUDA dependency.
| } | ||
| #ifdef BRPC_WITH_GDR | ||
| if (FLAGS_use_gdr) { | ||
| brpc::rdma::GlobalGdrInitializeOrDie(); |
There was a problem hiding this comment.
[replied by brpc-oncall robot] With the else if turned into a plain if, --use_rdma=false --use_gdr=true now calls GlobalGdrInitializeOrDie() while RDMA is still uninitialized. GlobalGdrInitializeOrDieImpl passes GetRdmaPd() to InitGPUBlockPool, but g_pd is only set by GlobalRdmaInitializeOrDieImpl, so it is still nullptr: ibv_reg_mr(nullptr, ...) fails and BlockPoolAllocator hits its LOG(FATAL). Please make GDR imply the RDMA initialization here, e.g. initialize RDMA when either flag is set:
if (FLAGS_use_rdma || FLAGS_use_gdr) {
brpc::rdma::GlobalRdmaInitializeOrDie();
}
#ifdef BRPC_WITH_GDR
if (FLAGS_use_gdr) {
brpc::rdma::GlobalGdrInitializeOrDie();
}
#endifAlternatively drop both explicit calls and rely on TransportFactory::ContextInitOrDie, as server.cpp already does.
There was a problem hiding this comment.
when gdr is open, rdma must be open.Flag validation has been added.
Done |
4893910 to
3d0a06e
Compare
|
fix code conflicts |
wwbmmm
left a comment
wwbmmm
left a comment
There was a problem hiding this comment.
[replied by brpc-oncall robot] The conflict resolution accidentally reverted the LOG(FATAL) → LOG(ERROR) change that chenBright asked for and you agreed to; please restore it. Apart from that the head matches the previously reviewed version.
| void* prefetch_d2h_data = NULL; | ||
| bool is_gpu_memory = source->is_gpu_memory(); | ||
| if (!is_gpu_memory) { | ||
| LOG(FATAL) << "RpcMessage is not in gpu!!!"; |
There was a problem hiding this comment.
[replied by brpc-oncall robot] This was changed to LOG(ERROR) in 3d0a06e in response to chenBright's comment, but the merge-conflict resolution restored the old version of this file and the LOG(FATAL) is back (the rest of the file, e.g. FillReqBufGpu, kept the LOG(ERROR) form). Please keep LOG(ERROR) here as well so an unexpected non-GPU buffer cannot abort the whole process.
There was a problem hiding this comment.
The LOG(FATAL) → LOG(ERROR) change has been reverted.
| #if BRPC_WITH_GDR | ||
| bool is_gpu_memory = msg->payload.is_gpu_memory(); | ||
| if (is_gpu_memory) { | ||
| if (!FillReqBufGpu(&req_buf, msg.get(), body_without_attachment_size)) { |
There was a problem hiding this comment.
Only the message body is copied back to host memory here; attachments remain in GPU memory. When a request has a non-empty attachment and both CRC32C and checksum_with_attachment are enabled, DeserializeRpcMessage() below passes the GPU memory address directly to CPU CRC32C through IOBufAsZeroCopyInputStream, causing an invalid memory access. The four-byte checksum check does not prevent this path, and the same issue applies on the response side.
There was a problem hiding this comment.
OK，the checksum_with_attachment has been supported when gdr is open. When gdr and checksum_with_attachment both are open, 1) the data will be copied from gpu to cpu and the checksum will be calced in sender 2) the data also wiil be copied from gpu to cpu and checksum will be checked in receiver.
There was a problem hiding this comment.
How about gdr not supporting checksum?
| #ifdef BRPC_WITH_GDR | ||
| #undef BRPC_WITH_GDR | ||
| #endif | ||
| #cmakedefine BRPC_WITH_GDR @WITH_GDR_VAL@ |
There was a problem hiding this comment.
#cmakedefine BRPC_WITH_GDR checks the CMake variable of the same name, but only WITH_GDR_VAL and -DBRPC_WITH_GDR=1 are currently set. As a result, when WITH_GDR=ON is enabled, this generated header first #undefs the macro and leaves it undefined, so the GPU implementation being compiled cannot see the GPU interfaces in subsequent headers. Please generate a numeric macro directly from WITH_GDR_VAL, or also set the variable used by #cmakedefine.
There was a problem hiding this comment.
ok，just as BRPC_WITH_RDMA, the definition of BRPC_WITH_GDR should not appear here and we have removed it.
| saved_errno, "Fail to create RdmaEndpoint: %s", berror(saved_errno)); | ||
| // No endpoint: later code paths must not assume RDMA is active | ||
| // and dereference _rdma_ep. Mark the transport as not using RDMA. | ||
| _rdma_state = RDMA_OFF; |
There was a problem hiding this comment.
When endpoint allocation fails, Socket::OnCreated() has not yet called HoldHCRelatedRef(). For a new Socket with only its creation reference, SetFailed() here synchronously releases the last reference and returns the object to the resource pool, but DoInit() and OnCreated() then continue initializing it. Setting RDMA_OFF only prevents access to a null endpoint; it does not resolve the lifetime issue.
There was a problem hiding this comment.
ok, when rdma ep create is failed, report error log and fallback to tcp, instead of socket SetFailed. The code has been updated.
| } else if (opt.socket_mode == SOCKET_MODE_URMA) { | ||
| buf.append("|urma"); | ||
| } else if (opt.socket_mode == SOCKET_MODE_GDR) { | ||
| buf.append("|gdr"); |
There was a problem hiding this comment.
With the default configuration, ComputeChannelSignature() returns a zero signature at the start and never reaches |gdr here. If a GDR Channel is created for an address first, followed by a regular baidu_std Channel for the same address, SocketMap::Insert() reuses the former's Socket. The regular channel will therefore also receive attachments in GPU memory, and the existing response_attachment().copy_to() will read the device address directly.
There was a problem hiding this comment.
OK，We limit that only SOCKET_MODE_TCP can return ComputeChannelSignature().
| return -1; | ||
| #if BRPC_WITH_GDR | ||
| if (_use_gdr) { | ||
| if (DoPostRecvGDR(_rbuf_data[_rq_received], g_gdr_recv_block_size, lkey) < 0) { |
There was a problem hiding this comment.
When rdma_recv_zerocopy=false, a zero-byte ACK does not enter the GDR zerocopy=true branch above and instead leads to PostRecv(1, false). This reuses the GPU memory address, but lkey remains at its initial value of 0; the next data receive may therefore cause a protection error and disconnect the connection.
There was a problem hiding this comment.
Because the rdma_recv_zerocopy is true at default, this issue has not been exposed. When wc.byte_len is 0, the zerocopy also should be true when gdr is open. Code has been updated.
| auto deleter = [](void* data) { free(data); }; | ||
| out->append_user_data(mem, saved_n, deleter); | ||
| } | ||
| pop_front(saved_n); |
There was a problem hiding this comment.
append_user_data() can return -1 if block metadata allocation fails, but this code still calls pop_front(saved_n) and returns a successful length. The destination receives no data, while the source data has already been consumed and the temporary host memory leaks. The protocol layer's length check also incorrectly reports success.
There was a problem hiding this comment.
ok, the result check of append_user_data has been added.
| // owned by anyone; release it back to the pool instead of leaking. | ||
| LOG(ERROR) << "Fail to append prefetched meta to IOBuf"; | ||
| host_allocator->DeallocateRaw(prefetch_d2h_data); | ||
| return MakeParseError(PARSE_ERROR_NOT_ENOUGH_DATA); |
There was a problem hiding this comment.
Before returning here, the code has already called pop_front(12) and obtained msg. Returning NOT_ENOUGH_DATA at this point makes the caller wait for new input, but the next parse still starts at the PRPC frame header while the source buffer now starts at the metadata. Meanwhile, msg is not returned to the message pool.
There was a problem hiding this comment.
OK，the error has been changed from PARSE_ERROR_NOT_ENOUGH_DATA to PARSE_ERROR_ABSOLUTELY_WRONG.
| meta.attachment_size(), res_size); | ||
| return false; | ||
| } | ||
| int body_without_attachment_size = res_size - meta.attachment_size(); |
There was a problem hiding this comment.
This omits the attachment_size < 0 check present in the CPU response path. If the current payload contains only A and the attachment length is -1, the message body length is calculated as 2. When the prefetch also includes the next frame, the fast path copies out AP, where P comes from the next frame's PRPC. The SerializedResponse path passes it directly to the caller as a successful response.
There was a problem hiding this comment.
OK，we align the check with the CPU response path。
|
当前 GDR 接收路径固定从 GPU block pool 按 gdr_block_size_kb（默认 512KB）分配，导致大消息被切成多个 block。对于接收端已有预分配连续 buffer 的场景（例如推理引擎的固定 shape 显存），应用只能用 copy_from_gpu(to_gpu=true) 再做一次 D2D，抵消了 GDR 的一部分收益。 而 DoPostRecvGDR(void* block, size_t block_size, uint32_t lkey) 本身已经支持任意 buffer。是否可以在 PostRecv 层增加一个可选的 allocator 回调，允许调用方提供自己的显存块？ |
@bzs1118 在当前不改变brpc双边send/recv的模式下，允许调用方自己管理recv的显存，从而避免recv后的一次d2d，我想也是可以的。具体来说，大体可能这样设计： 另一种方法为，将brpc当前的双边send/recv模式，改成单边wirte模式，大体可能这样设计： 辛苦 @yanglimingcn @wwbmmm @chenBright 大佬们看看，这样设计是否合适 |
单边操作更合理。单边操作不会影响双边 send/recv 的数据，不用多一个 QP。 我这边已经实现了单边操作，等这个 PR 合入后，会贡献到社区。 |
| return false; | ||
| } | ||
| res_buf->append((char*)data + meta_size, res_size); | ||
| msg->payload.pop_front(res_size); |
There was a problem hiding this comment.
append() returns -1 when host block allocation fails, but this code still calls pop_front() and returns success, discarding the only copy of the original body. The SerializedResponse path can even pass an empty body to the caller as a successful response. Please consume the payload only after the append succeeds, and fix the corresponding operations in the request fast path and responses with attachments as well.
There was a problem hiding this comment.
Yes，the result of append should be checked.
| // of `source'. Continuing would make the following cutn treat the | ||
| // meta bytes as payload and silently corrupt the message. | ||
| LOG(ERROR) << "Fail to cut meta_size=" << meta_size << " from gpu"; | ||
| return MakeParseError(PARSE_ERROR_NOT_ENOUGH_DATA); |
There was a problem hiding this comment.
When this returns PARSE_ERROR_NOT_ENOUGH_DATA, the complete frame has already arrived, and line 163 has already removed the frame header. After the copy fails, parsing again starts at the metadata and cannot recover the original frame boundary; msg is also not freed. Please destroy the message and return a terminal parse error as in the small-metadata failure branch above, or ensure that the input is not consumed at all on failure.
There was a problem hiding this comment.
OK，When cutn_from_gpu fails， the msg should be destroyed and PARSE_ERROR_ABSOLUTELY_WRONG should be triggered.
| } | ||
| } | ||
| source->cutn(&msg->payload, body_size - meta_size); | ||
| return MakeMessage(msg); |
There was a problem hiding this comment.
The payload returned here still holds GPU blocks, but ProcessRpcRequest passes it directly to rpc_dump before converting the body. RpcDumpContext then passes these addresses to a regular writev, so writing GPU requests fails and may even leave incomplete samples. Please convert sampled data to a host copy at the sampling boundary and handle copy failures.
There was a problem hiding this comment.
Supporting rpcdump within the GDR feature requires a lot of testing. Could we perhaps provide a notification to users—stating that rpcdump is not currently supported when using GDR—and then submit a follow-up PR later to enable rpcdump directly within GDR?
| # gpu_block_pool.cpp is a no-op translation unit without BRPC_WITH_GDR; | ||
| # only link CUDA libs when GDR is actually enabled. | ||
| list(APPEND butil_libraries ${CUDA_LIB} ${CUDART_LIB}) | ||
| endif() |
There was a problem hiding this comment.
This only updates the internal CMake link dependencies and does not update butil_private_libs. As a result, the installed brpc.pc lacks the CUDA dependencies, and external programs linking the GDR version using only pkg-config --static --libs brpc will encounter unresolved CUDA symbols. Please also export the required libraries and search paths, and check a standalone consumer after installation.
| if (cudaMalloc(&gpu_addr, size) != cudaSuccess) { | ||
| LOG(ERROR) << "Fail to alloc gpu memory for attachment"; | ||
| } else if (cntl->request_attachment().copy_from_gpu( | ||
| gpu_addr, size, 0, /*to_gpu=*/true) != size) { |
There was a problem hiding this comment.
A GDR service can still handle ordinary TCP requests, in which case attachments are in host memory. However, copy_from_gpu(..., true) here ultimately always calls cudaMemcpyDeviceToDevice, so it cannot support H2D as the comment above says. Please select the copy direction based on the actual source memory type; the CUDA documentation defines a mismatch between the direction and the pointers as undefined behavior.
There was a problem hiding this comment.
Right, when gdr connection fallback to tcp, the attachment will be on host memory. We should consider this situation.Code has been updated.
| if(CUDA_HOME) | ||
| set(CUDA_INCLUDE_PATH ${CUDA_HOME}/include) | ||
| find_library(CUDA_LIB NAMES cuda PATHS ${CUDA_HOME}/lib64 ${CUDA_HOME}/lib NO_DEFAULT_PATH) | ||
| find_library(CUDART_LIB NAMES cudart PATHS ${CUDA_HOME}/lib64 ${CUDA_HOME}/lib NO_DEFAULT_PATH) |
There was a problem hiding this comment.
CUDA_HOME points to the Toolkit, while libcuda.so may be installed by the driver in a system directory. NO_DEFAULT_PATH here excludes that directory, causing environments with the headers, Runtime, and driver all installed to still report that CUDA cannot be found. Please retain a fallback to system paths for the driver library and update the same logic in the example; see the NVIDIA installation guide for this layout.
1) recv all data on gpu first 2) the gpu block is alloced from a gpu block pool 3) brpc header, meta and body will be copied from gpu to cpu to process. 4) To decrease the d2h counts when processing header, meta and body, we will prefetch 512B to memory Co-authored-by: sunce4t <847480001@qq.com>
|
|
||
| static constexpr int kMaxConcurrent = 32; | ||
| private: | ||
| int gpu_id_ {-1}; |
There was a problem hiding this comment.
gpu_id_ -> _gpu_id;
Please change all variables so that they have an underscore (_) as a prefix.
| "fall back to TCP transport"; | ||
| // No endpoint: later code paths must not assume RDMA is active | ||
| // and dereference _rdma_ep. Mark the transport as not using RDMA. | ||
| _rdma_state = RDMA_OFF; |
There was a problem hiding this comment.
If endpoint allocation fails, setting RDMA_OFF alone is not enough to safely fall back to TCP: the server's socket_mode remains RDMA/GDR, so ExecuteServerHandshake() enters the actual handshake and triggers CHECK(ep != nullptr); the client's Connect() also still returns RdmaConnect, whose StartConnect() likewise checks for a non-null endpoint. This new failure branch will therefore still terminate the process. Please also switch the socket mode to TCP and make Connect() use the TCP connection logic when there is no endpoint; add a regression test for endpoint allocation failure on both sides.
| SocketMode mode = socket->socket_mode(); | ||
| if (mode == SOCKET_MODE_RDMA | ||
| #if BRPC_WITH_GDR | ||
| || mode == SOCKET_MODE_GDR |
There was a problem hiding this comment.
This allows the GDR handshake, but the server does not yet guarantee that the GPU-aware baidu_std parser is registered. For example, with socket_mode=SOCKET_MODE_GDR and enabled_protocols="http", the existing option validation allows startup, and BuildAcceptor() retains the handshake and HTTP but removes the Baidu RPC parser. After the client completes the handshake, GPU data passes through CutInputMessage() to HTTP's ParseFromIOBuf(), which passes the device address directly to http_parser_execute(), causing an invalid memory access. The termination guard for invalid frames in ParseRpcMessageGpu() does not run in this configuration. Please reject configurations that do not enable baidu_std during GDR server initialization, or prevent GPU input from reaching non-GPU parsers at the dispatch boundary.
There was a problem hiding this comment.
yes, we should reject gdr when start server with protocol which is not baidu_std.
| }) + select({ | ||
| # NOTE: `defines` is auto-prefixed with -D by Bazel, so do NOT add -D here. | ||
| # Compare with the rdma/glog/thrift entries above which also omit the prefix. | ||
| "//bazel/config:brpc_with_gdr": ["BRPC_WITH_GDR=1"], |
There was a problem hiding this comment.
When only --define BRPC_WITH_GDR=true is passed, this does not enable BRPC_WITH_RDMA or select the RDMA -libverbs link option. gdr_transport.h declares a class that inherits from RdmaTransport, but the latter is excluded by #if BRPC_WITH_RDMA, so compilation fails even when the CUDA and verbs dependencies are correctly installed. Please make GDR enable the RDMA macro and link dependency as config_brpc.sh does, or explicitly reject this combination during configuration.
| // instead of crashing the whole server. | ||
| LOG(ERROR) << "alloc host data for prefetch d2h failed, size=" | ||
| << prefetch_d2h_size; | ||
| return MakeParseError(PARSE_ERROR_NOT_ENOUGH_DATA); |
There was a problem hiding this comment.
When a complete frame has arrived but the host pool is exhausted or expansion fails, returning PARSE_ERROR_NOT_ENOUGH_DATA here does not schedule a retry: ProcessNewMessage() exits the parsing loop, the RDMA receive path calls it again only when a subsequent receive completion event arrives, and DeallocateRaw() has no wake-up mechanism. Therefore, even if memory becomes available later, the frame may remain in the receive buffer until the RPC times out or another receive event occurs. Consider treating allocation failure as a terminal parsing error and closing the connection, or implementing a resource-recovery retry that does not depend on new input; do not treat insufficient resources as an incomplete frame.
There was a problem hiding this comment.
OK，change the error from PARSE_ERROR_NOT_ENOUGH_DATA to PARSE_ERROR_ABSOLUTELY_WRONG.
| if (host == nullptr) { | ||
| LOG(ERROR) << "Fail to alloc host buffer to checksum gpu memory"; | ||
| // Fail closed: return a value that cannot match the expected crc. | ||
| return crc ^ 0xdeadbeefu; |
There was a problem hiding this comment.
crc ^ 0xdeadbeefu is still a valid CRC value and cannot indicate checksum failure. When GPU attachment checksumming is enabled, if allocation of the temporary host buffer fails on both the sender and receiver, both sides return the same value based on the same body CRC; Crc32cCompute() sends that value, and Crc32cVerify() only compares them for equality, ultimately returning success without checking the attachment at all. Consider having the internal calculation function return the success status and CRC separately: return false immediately on verification failure, and abort normal sending through the Controller/serialization call chain on calculation failure, rather than substituting a fixed XOR value for the error.
| // GetGdrBlockSize() returns 0 on an invalid FLAGS_gdr_block_size_kb; | ||
| // subtracting the header from it would underflow to a huge value. | ||
| const size_t gdr_block_size = butil::gdr::GetGdrBlockSize(); | ||
| if (gdr_block_size <= IOBUF_BLOCK_HEADER_LEN) { |
There was a problem hiding this comment.
This validation still accepts --gdr_block_size_kb=1, but the advertised receive length is only 992 after subtracting the 32-byte header. Both V2's ValidHelloMessage() and V3's ValidRdmaHello() require block_size >= MIN_BLOCK_SIZE (1024), so this configuration passes initialization but causes all connections to fall back to TCP, making GDR unavailable. Please validate the handshake minimum against the length after subtracting the header, explicitly reject this configuration at startup, and cover the 1 KB/2 KB boundary cases.
| #if BRPC_WITH_GDR | ||
| bool is_gpu_memory = msg->payload.is_gpu_memory(); | ||
| if (is_gpu_memory) { | ||
| if (!FillReqBufGpu(&req_buf, msg.get(), body_without_attachment_size)) { |
There was a problem hiding this comment.
The GPU body conversion here only covers the ordinary service branch; the earlier IsBaiduMasterService() branch still directly uses cutn() to move the GPU payload into SerializedRequest::serialized_data(). Therefore, after a GDR request reaches a server configured with baidu_master_service, even without an attachment, if the proxy forwards the SerializedRequest unchanged over a regular TCP Channel, SerializeRpcRequest() still only shares those device memory blocks, and writev() ultimately attempts to send from device addresses and fails with EFAULT; directly parsing the body also causes an invalid memory access. Consider converting the body to a host IOBuf in the master-service branch as well, or moving the conversion to a location shared by both branches; attachments can remain on the GPU.
There was a problem hiding this comment.
It is not recommended to use baidu master service (the proxy-forwarding mode) on a server with GDR enabled. master service forwards each request one more time, and the extra network hop plus proxy scheduling add significant latency. Moreover, when the downstream is a plain TCP Channel, the body must be copied back from GPU memory to host memory (D2H) before forwarding, which completely negates the zero-copy benefit of GDR. GDR targets clients directly connecting to GPU services (e.g. training/inference servers) so that data goes GPU-to-GPU; for scenarios that require proxy forwarding, use a plain TCP/RDMA link instead.
Adopt the Google C++ style underscore prefix for data members throughout the GDR block pool: - BlockPoolAllocator: gpu_id -> _gpu_id, on_gpu -> _on_gpu, pd -> _pd, freeList -> _free_list, poolMutex -> _pool_mutex, g_region_num -> _region_num, g_regions -> _regions, totalAllocated/_deallocated/ peakUsage -> _total_allocated/_total_deallocated/_peak_usage. - GPUStreamPool: suffix style (gpu_id_, d2h_cnt_, ...) to prefix style (_gpu_id, _d2h_cnt, ...). - BlockPoolAllocators: gpu_mem_alloc/cpu_mem_alloc/gpu_stream_pool -> _gpu_mem_alloc/_cpu_mem_alloc/_gpu_stream_pool. Function parameters keep their unprefixed names; the file-scope g_regions/g_region_num in rdma/block_pool.cpp are separate statics of the host RDMA pool and are untouched. Pure rename, no logic change.
The master service forwards the raw SerializedRequest to a backend chosen by GetChannel() without running any user code, and that backend is usually not GDR-capable: forwarding GPU-resident blocks would make writev() send from device addresses and fail with EFAULT, and copying everything down to host memory first would silently negate the zero-copy purpose of GDR. Reject requests arriving on GDR connections with an explicit EREQUEST error instead, and document the combination as unsupported. Non-GDR builds and host-resident requests (e.g. plain TCP clients of the same proxy server) are unaffected.
The handshake advertises the receive SGE length, i.e. the block size minus the 32-byte iobuf block header. Both V2's ValidHelloMessage() and V3's ValidRdmaHello() require block_size >= MIN_BLOCK_SIZE (1024), checked against the advertised length. A small --gdr_block_size_kb (e.g. 1 -> 1024-byte block, 992-byte advertised length) passed the existing >header check, but every GDR handshake would then fail and silently fall back to TCP. Reject such block sizes in GlobalGdrInitialize() against the advertised length (gdr_block_size - IOBUF_BLOCK_HEADER_LEN >= MIN_BLOCK_SIZE), so the misconfiguration surfaces as a startup error instead of silent TCP fallback. GlobalGdrInitialize() failure is fatal via ExitWithError() in rdma_helper.cpp, and the error message suggests the minimum --gdr_block_size_kb. Add RdmaGdrTest.gdr_block_size_boundaries covering the 1 KB (rejected, EINVAL) and 2 KB (accepted, advertised length 2016) boundaries and restoring the original flag afterwards. The test only exercises the size computation, no CUDA/GPU is needed.
When checksum_with_attachment is on and the attachment lives in GPU memory, crc32c computed the CRC over a temporary host copy. If that host buffer allocation failed, the old code substituted "crc ^ 0xdeadbeefu" as a failure marker. That XOR result is itself a valid CRC value: if both sides fail to allocate, both derive the same magic value from the same body CRC, and Crc32cVerify's equality check passes -- the attachment was never checked but the request or response is reported as verified. Separate the status from the CRC value throughout the crc32c path: - ExtendCrc32cFromGpu/ExtendCrc32c/ComputeCrc32c return bool and write the CRC through an out parameter (which ComputeCrc32c initializes to 0 first); failures never touch it. All 0xdeadbeef placeholders are gone. - Crc32cCompute returns bool: on failure it writes no checksum_value, and SerializeRpcMessage fails the controller (EREQUEST) so the send is aborted through the existing call sites; on failure Crc32cVerify returns false immediately. - ComputeDataChecksum/ChecksumHandler::Compute return bool accordingly; VerifyDataChecksum now returns false instead of true when no handler is registered for the requested type. - The crc32c unittest asserts Crc32cCompute succeeds. Verified locally (macOS, non-GDR build): brpc_checksum_unittest 10/10.
When the host block pool is exhausted or expansion fails, the GPU parser returned PARSE_ERROR_NOT_ENOUGH_DATA for the prefetch D2H buffer allocation. That error means "incomplete frame, wait for more input": ProcessNewMessage exits the parsing loop and only re-enters on the next receive completion event, and DeallocateRaw has no mechanism to wake the parser. The already-arrived frame could thus stay stuck in the receive buffer until the RPC times out, even after memory became available again. Treat the allocation failure as what it is -- a resource failure on an already-arrived frame, not an incomplete one -- and return PARSE_ERROR_ABSOLUTELY_WRONG so the connection is closed promptly; the client reconnects and retries. This also matches how the other unrecoverable errors in this parser are handled.
Building with only --define BRPC_WITH_GDR=true left the build broken in two ways: - The defines select emitted only BRPC_WITH_GDR=1. gdr_transport.h inherits from RdmaTransport, which is compiled only under BRPC_WITH_RDMA, so the RDMA symbols were missing. - The link select passed only -lcuda/-lcudart, missing -libverbs for the verbs symbols referenced by the inherited RDMA code. GDR implies RDMA, mirroring CMake (WITH_GDR implies WITH_RDMA) and config_brpc.sh (which appends both -DBRPC_WITH_RDMA and -DBRPC_WITH_GDR): add BRPC_WITH_RDMA=1 to the defines select and -libverbs to the link select. When both GDR and RDMA defines/flags are passed explicitly, the duplicated entries are harmless.
With ServerOptions.socket_mode=SOCKET_MODE_GDR, rdma_handshake still performs the GDR handshake, and BuildAcceptor always serves rdma_handshake itself, but it honors ServerOptions.enabled_protocols for the payload protocols. A whitelist that excludes baidu_std (e.g. enabled_protocols="http") would therefore accept GDR connections whose payloads arrive in GPU memory, then hand the device pointer to a non-GPU parser (e.g. http_parser_execute on a device address), crashing the process; the bad-frame guard in ParseRpcMessageGpu does not run for those protocols. baidu_std is the only parser that understands GPU-resident payloads, so reject the combination at startup: socket_mode=SOCKET_MODE_GDR with a non-empty enabled_protocols that lacks baidu_std now fails Server::Start with a clear error. An empty whitelist (all protocols) and normal GDR servers are unaffected. The check lives under #if BRPC_WITH_GDR, so non-GDR builds are unchanged.
When new(std::nothrow) RdmaEndpoint returned null in RdmaTransport::DoInit(), setting _rdma_state = RDMA_OFF only switched the data path to the TCP transport, but the socket mode stayed RDMA/GDR, so two paths still terminated the process: - Server: ExecuteServerHandshake() dispatches on the socket mode, entered RdmaEndpoint::ExecuteServerHandshake() and hit CHECK(ep != nullptr). - Client: Socket::Connect() still got an RdmaConnect whose StartConnect() CHECKs for a non-null endpoint. Fix both sides of the fallback: - DoInit() additionally resets socket->_socket_mode to SOCKET_MODE_TCP, so the server handshake takes the fallback path. - RdmaTransport::Connect() returns nullptr when there is no endpoint, so Socket::Connect() skips RdmaConnect and uses its own TCP connection logic (AfterAppConnected directly). The allocation failure is simulated in tests through a test-only g_rdma_ep_alloc_fail injection point (mirroring g_skip_rdma_init): RdmaTest.endpoint_alloc_fail_falls_back_to_tcp verifies that an accepted server socket has no endpoint, RDMA_OFF state, TCP socket mode and closes cleanly, and that a client channel in RDMA mode completes an RPC over TCP with a TCP-mode socket. The change compiles in a non-RDMA build (macOS); the test itself needs an RDMA build to run.
What problem does this PR solve?
Issue Number: resolve #3102
Problem Summary:
What is changed and the side effects?
Changed:
Side effects:
Performance effects:
Breaking backward compatibility:
Check List: