support gpu direct rdma by randomkang · Pull Request #3144 · apache/brpc

randomkang · 2025-11-09T02:51:10Z

What problem does this PR solve?

Issue Number: resolve #3102

Problem Summary:

What is changed and the side effects?

Changed:

recv all data on gpu first
the gpu block is alloced from a gpu block pool
brpc header, meta and body will be copied from gpu to cpu to process.
To decrease the d2h counts, we will prefetch 512B to memory

Side effects:

Performance effects:
Breaking backward compatibility:

Check List:

Please make sure your changes are compilable.
When providing us with a new feature, it is best to add related tests.
Please follow Contributor Covenant Code of Conduct.

Copilot

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.

Copilot · 2025-11-09T08:37:24Z

+    clock_gettime(CLOCK_MONOTONIC, &end);
+    double time_us = (end.tv_sec - start.tv_sec) * 1e6 + (end.tv_nsec - start.tv_nsec) / 1e3;


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.

Copilot · 2025-11-09T08:37:25Z

+    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;


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.

Suggested change

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;

Copilot · 2025-11-09T08:37:25Z

+    struct timespec start, end;
+    clock_gettime(CLOCK_MONOTONIC, &start);


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.

Copilot · 2025-11-09T08:37:26Z

+    struct timespec start, end;
+    clock_gettime(CLOCK_MONOTONIC, &start);


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.

yanglimingcn · 2025-11-10T01:25:08Z

+
    char header_buf[12];
-    const size_t n = source->copy_to(header_buf, sizeof(header_buf));
+    size_t n = 0;


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.

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

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.

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?

yanglimingcn · 2025-12-02T07:31:08Z

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.
https://zhuanlan.zhihu.com/p/376989325 This link contains some relevant details.

randomkang · 2025-12-14T12:53:32Z

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. https://zhuanlan.zhihu.com/p/376989325 This link contains some relevant details.

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.

randomkang · 2026-08-04T05:51:55Z

@yanglimingcn @zchuango all comments have been resolved, please review again.

yanglimingcn · 2026-08-25T11:46:03Z

Turn multiple commits into one

randomkang · 2026-08-25T12:59:31Z

Turn multiple commits into one

ok

yanglimingcn · 2026-08-27T11:25:48Z

LGTM

yanglimingcn · 2026-08-28T02:20:34Z

+    return (void *)d_data;
+}
+
+void* get_cpu_mem(int gpu_id, int64_t cpu_mem_size) {


这个入参，应该是gpu_id吗

是的。这里申请的是pinned memory，使用pinned memory和显存进行交换，会比普通的memory快。在使用cudaMallocHost分配pinned memory的时候，建议先执行cudaSetDevice，这样分配的pined memory就会与对应的gpu id关联。

yanglimingcn · 2026-08-28T02:22:48Z

+
+std::atomic<BlockPoolAllocators*> BlockPoolAllocators::instance_{nullptr};
+
+BlockPoolAllocators* BlockPoolAllocators::singleton() {


你看brpc自带的singleton.h是否能满足你的需求。

我看了下，可以用brpc自带的Singleton，已改

yanglimingcn · 2026-08-28T02:31:53Z


+void* IOBuf::get_first_data_ptr() {
+    if (_ref_num() == 0) {
+        return 0;


这里返回nullptr

randomkang · 2026-10-08T03:49:00Z

@yanglimingcn @zchuango Are there any other issues with this pr?

Copilot

🟡 Changes recommended

Critical build, handshake, parsing, allocation, and memory-pool correctness issues remain unresolved.

13 open findings

Bazel GDR build lacks CUDA include and dependency configuration · New
Dispatch GDR connections through the RDMA handshake · New
Stop response processing after GPU buffer fill failure · New
Prevent GPU parse errors from falling through to non-GDR parsers · New
Validate GDR block size before initializing the GPU pool · New
Propagate CUDA allocation failures instead of aborting · New
Synchronize GPU memory region lookup and publication · New
Use value-based BRPC_WITH_GDR guards in iobuf.cpp · New
Allocate GDR attachment on the configured GPU · New
Handle GPU IOBuf metadata allocation failure · New
Select the configured CUDA device in each GPU operation · New
Correct socket mode enum values in documentation · New
Add tests for GPU prefetch and malformed attachment handling · New

🧠 Review effort: Lite

Give feedback about Copilot approvals in this survey to enter a drawing for a $150 gift card.

randomkang · 2026-10-08T07:03:13Z

+    "//bazel/config:brpc_with_gdr": [
+        "-lcuda",
+        "-lcudart",
+    ],


We do not use any cuda kernel in this pr, so CUDA_HOME is not necessary. The cuda headers and libs can be found by C_INCLUDE_PATH and LD_LIBRARY_PATH.

+ParseResult ParseRpcMessageGpu(butil::IOBuf* source, Socket* socket,
+                            bool /*read_eof*/, const void*) {


chenBright · 2026-10-08T04:36:11Z

+#include "butil/raw_pack.h"                      // RawPacker RawUnpacker
+#include "butil/memory/scope_guard.h"
+#include "butil/raw_pack.h"                      // RawPacker RawUnpacker


butil/raw_pack.h is duplicated.

Copilot

Warning

Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.

Copilot review overview

20 open findings

With the current if/else if, GlobalGdrInitializeOrDie() won’t run when both --use_rdma=true… · New
If new(std::nothrow) fails, _rdma_ep stays null but _rdma_state is still set to… · New
Use value-based BRPC_WITH_GDR guards in iobuf.cpp
Synchronize GPU memory region lookup and publication
Propagate CUDA allocation failures instead of aborting
Validate GDR block size before initializing the GPU pool
Prevent GPU parse errors from falling through to non-GDR parsers
Stop response processing after GPU buffer fill failure
Dispatch GDR connections through the RDMA handshake
Bazel GDR build lacks CUDA include and dependency configuration
Using LOG(FATAL) in the protocol parser will terminate the whole process on unexpected… · New
This uses type-punning and potentially unaligned uint32_t loads from header_buf/string literal,… · New
Same issue as the parser: LOG(FATAL) here can crash the server if the message is unexpectedly… · New
This PR changes the default rdma_prepared_qp_cnt from 1024 to 256, which is a… · New
The allocation fast-path always takes two timestamps and computes duration, even in non-DEBUG… · New
The initial cudaStreamSynchronize() before enqueueing copies is redundant if you already… · New
Select the configured CUDA device in each GPU operation
Handle GPU IOBuf metadata allocation failure
Allocate GDR attachment on the configured GPU
Add tests for GPU prefetch and malformed attachment handling

1 resolved since last review

Correct socket mode enum values in documentation

🧠 Review effort: Lite

Give feedback about Copilot approvals in this survey to enter a drawing for a $150 gift card.

randomkang · 2026-10-21T06:28:41Z

+    bool is_gpu_memory = msg->payload.is_gpu_memory();
+    if (!is_gpu_memory) {
+        LOG(FATAL) << "message is not on gpu!!!";
+    }


We expect data is on gpu here, so the fatal is necessary and can help to find errors.

+    std::lock_guard<std::mutex> stream_lock(d2d_locks_[stream_idx]);
+    CHECK_CUDA(cudaStreamSynchronize(d2d_streams_[stream_idx]));
+    for (int i = 0; i < segs; i++) {
+        if (length_list[i] == 0) {
+            continue;
+        }
+        CHECK_CUDA(cudaMemcpyAsync(static_cast<char*>(dst) + offset, src_list[i], length_list[i],
+                    cudaMemcpyDeviceToDevice, d2d_streams_[stream_idx]));
+        offset += length_list[i];
+    }
+    CHECK_CUDA(cudaStreamSynchronize(d2d_streams_[stream_idx]));


wwbmmm · 2026-10-08T12:47:02Z

please resolve merge conflict

randomkang · 2026-10-11T01:59:33Z

please resolve merge conflict

ok

wwbmmm

[replied by brpc-oncall robot] The latest push fixes most earlier findings; three issues remain on the newly changed lines: the reverted socket-mode comment, MLX5_SCATTER_TO_CQE being set after the RDMA context is already open, and an early return that leaves the transport partially initialized.

wwbmmm · 2026-10-12T05:47:44Z

    std::shared_ptr<SocketSSLContext> _ssl_ctx;

-    // Choose to use a certain socket: 0 TCP, 1 RDMA
+    // Choose to use a certain socket: 0 TCP, 1 RDMA, 2 UBRING, 3 GDR


[replied by brpc-oncall robot] The merge-conflict resolution reverted this comment to 2 UBRING, 3 GDR, but the enum is SOCKET_MODE_URMA = 3, SOCKET_MODE_GDR = 4 (see src/brpc/socket_mode.h). Please restore 0 TCP, 1 RDMA, 2 UBRING, 3 URMA, 4 GDR so it stays in sync with the enum.

wwbmmm · 2026-10-12T05:47:44Z

+    // is enabled, small messages are inlined directly into host memory, which
+    // breaks this requirement and causes operations to fail. Force it to 0
+    // regardless of the user-provided value.
+    if (setenv("MLX5_SCATTER_TO_CQE", "0", 1) != 0) {


[replied by brpc-oncall robot] MLX5_SCATTER_TO_CQE is set here, i.e. after the RDMA context has already been created: TransportFactory::ContextInitOrDie calls RdmaTransport::ContextInitOrDie() (which calls GlobalRdmaInitializeOrDie() → ibv_open_device) before GdrContextInitOrDie(), and the example also calls GlobalRdmaInitializeOrDie() first. libmlx5 reads the MLX5_* environment variables when the device context is created, so setting it after ibv_open_device has no effect. Please move it before the RDMA global initialization (or drop it and keep the value as an externally documented prerequisite).

wwbmmm · 2026-10-12T05:47:45Z

+            _rdma_state = RDMA_OFF;
+            _socket = socket;
+            _default_connect = options.app_connect;
+            return;


[replied by brpc-oncall robot] This early return happens before _on_edge_trigger is set and before _tcp_transport is created/initialized, while _rdma_state is set to RDMA_OFF. RdmaTransport::CutFromIOBuf/CutFromIOBufList/WaitEpollOut dereference _tcp_transport unconditionally, so any later call on this transport would crash. Even though the socket was just SetFailed, please leave the transport in a fully initialized state (e.g. fall through to the common tail of DoInit) rather than returning in the middle.

randomkang · 2026-10-20T09:57:15Z

Is there other problem for this pr? @wwbmmm @chenBright

Copilot

Warning

Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.

Copilot review overview

23 open issues

And 3 more that still need to be addressed.

_{Review effort: Lite}

randomkang · 2026-10-21T06:06:36Z

+#if BRPC_WITH_GDR
+    // GPUDirect RDMA requires all received data to be written into the GPU
+    // memory buffer specified at post-receive time. When MLX5_SCATTER_TO_CQE
+    // is enabled, small messages are inlined directly into host memory, which
+    // breaks this requirement and causes operations to fail. libmlx5 reads
+    // MLX5_* environment variables when the device context is created (i.e.
+    // inside ibv_open_device below), so it must be forced to 0 before any
+    // device is opened, regardless of the user-provided value.
+    if (setenv("MLX5_SCATTER_TO_CQE", "0", 1) != 0) {
+        PLOG(ERROR) << "Fail to set MLX5_SCATTER_TO_CQE to 0";
+        ExitWithError();
+    }
+#endif  // if BRPC_WITH_GDR


If user do not want to use gdr, then should not build brpc with gdr, and this environment variable setting will not take effect.

randomkang · 2026-10-21T09:45:51Z

All Copilot issues have been resolved or commented. @chenBright @wwbmmm

wwbmmm · 2026-10-25T09:15:38Z

please resolve merge conflicts

wwbmmm

[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.

wwbmmm · 2026-10-25T09:15:53Z

    }
+#ifdef BRPC_WITH_GDR
+    if (FLAGS_use_gdr) {
+        brpc::rdma::GlobalGdrInitializeOrDie();


[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(); } #endif

Alternatively drop both explicit calls and rely on TransportFactory::ContextInitOrDie, as server.cpp already does.

when gdr is open, rdma must be open.Flag validation has been added.

randomkang · 2026-10-25T13:40:18Z

please resolve merge conflicts

Done

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>

wwbmmm requested review from Copilot and yanglimingcn November 9, 2025 08:34

Copilot AI reviewed Nov 9, 2025

View reviewed changes

yanglimingcn reviewed Nov 10, 2025

View reviewed changes

Comment thread src/brpc/rdma/rdma_endpoint.cpp Outdated

sunce4t force-pushed the gdr branch from 6eedb3e to 1fb9370 Compare November 26, 2025 02:17

randomkang force-pushed the gdr branch 2 times, most recently from 041bef0 to 9e3b2af Compare June 2, 2026 16:41

randomkang force-pushed the gdr branch from 9e3b2af to 051d0b2 Compare August 3, 2026 17:11

yanglimingcn requested a lite review from Copilot August 25, 2026 10:00

Copilot started reviewing on behalf of yanglimingcn August 25, 2026 10:40

View session

Copilot stopped reviewing on behalf of yanglimingcn due to an error August 25, 2026 11:00
Copilot had to stop work due to a timeout.

yanglimingcn reviewed Aug 25, 2026

View reviewed changes

Comment thread docs/cn/gdr.md Outdated

yanglimingcn reviewed Aug 25, 2026

View reviewed changes

Comment thread docs/cn/gdr.md Outdated

yanglimingcn reviewed Aug 25, 2026

View reviewed changes

Comment thread src/brpc/acceptor.h Outdated

yanglimingcn reviewed Aug 25, 2026

View reviewed changes

Comment thread src/butil/gpu/gpu_block_pool.cpp Outdated

randomkang force-pushed the gdr branch from 2d74338 to bf328f5 Compare August 25, 2026 12:57

yanglimingcn reviewed Aug 28, 2026

View reviewed changes

randomkang force-pushed the gdr branch from bf328f5 to 11ee2f3 Compare August 28, 2026 08:50

randomkang force-pushed the gdr branch from 11ee2f3 to 28972ad Compare October 8, 2026 03:42

chenBright removed the request for review from Copilot October 8, 2026 04:39

chenBright requested review from zchuango and a lite review from Copilot and removed request for zchuango October 8, 2026 04:39

Copilot started reviewing on behalf of chenBright October 8, 2026 04:39

View session

Copilot AI reviewed Oct 8, 2026

View reviewed changes

chenBright reviewed Oct 8, 2026

View reviewed changes

wwbmmm requested a lite review from Copilot October 8, 2026 09:57

Copilot AI reviewed Oct 8, 2026

View reviewed changes

Copilot stopped reviewing on behalf of wwbmmm due to an error October 8, 2026 11:00
An unexpected error occurred. For more details, see the detailed logs in GitHub Actions.

randomkang force-pushed the gdr branch from 4ed2a6c to da3556c Compare October 11, 2026 01:56

randomkang force-pushed the gdr branch from f4041e5 to ed9c38c Compare October 11, 2026 15:16

wwbmmm reviewed Oct 12, 2026

View reviewed changes

randomkang force-pushed the gdr branch from b974531 to c71d15f Compare October 12, 2026 13:17

chenBright requested a balanced review from Copilot October 20, 2026 16:18

Copilot AI reviewed Oct 20, 2026

View reviewed changes

Copilot stopped reviewing on behalf of chenBright due to an error October 20, 2026 19:30
Workflow completed but conclusion is not available. This may indicate rate limiting or API issues.

randomkang force-pushed the gdr branch from c71d15f to 806d967 Compare October 21, 2026 09:44

wwbmmm reviewed Oct 25, 2026

View reviewed changes

randomkang force-pushed the gdr branch from 806d967 to 4893910 Compare October 25, 2026 13:40

randomkang force-pushed the gdr branch from 4893910 to 3d0a06e Compare October 25, 2026 15:57

		clock_gettime(CLOCK_MONOTONIC, &end);
		double time_us = (end.tv_sec - start.tv_sec) * 1e6 + (end.tv_nsec - start.tv_nsec) / 1e3;

		struct timespec start, end;
		clock_gettime(CLOCK_MONOTONIC, &start);


		std::atomic<BlockPoolAllocators*> BlockPoolAllocators::instance_{nullptr};

		BlockPoolAllocators* BlockPoolAllocators::singleton() {

		ParseResult ParseRpcMessageGpu(butil::IOBuf* source, Socket* socket,
		bool /read_eof/, const void*) {

Conversation

randomkang commented Nov 9, 2025 • edited by chenBright Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What problem does this PR solve?

What is changed and the side effects?

Check List:

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull Request Overview

Reviewed Changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI Nov 9, 2025

Choose a reason for hiding this comment

Uh oh!

Copilot AI Nov 9, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Copilot AI Nov 9, 2025

Choose a reason for hiding this comment

Uh oh!

Copilot AI Nov 9, 2025

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

randomkang Nov 14, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

yanglimingcn commented Dec 2, 2025

Uh oh!

randomkang commented Dec 14, 2025

Uh oh!

randomkang commented Aug 4, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

yanglimingcn commented Aug 25, 2026

Uh oh!

randomkang commented Aug 25, 2026

Uh oh!

yanglimingcn commented Aug 27, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

randomkang commented Oct 8, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

randomkang commented Nov 9, 2025 •

edited by chenBright

Loading

randomkang Nov 14, 2025 •

edited

Loading