diff --git a/README.md b/README.md
index 071c38d..238e289 100644
--- a/README.md
+++ b/README.md
@@ -355,7 +355,7 @@ export LD_LIBRARY_PATH=${INFINI_INSTALL}/lib:$LD_LIBRARY_PATH
 |---------|---------------|----------------------|---------------|
 | **OpenMPI** | Full | `WITH_OMPI=ON` | The default backend. Requires the OpenMPI development package.|
 | **MPICH** | Full | `WITH_MPICH=ON` | Requires the MPICH development package.|
-| **NCCL** | Partial | `WITH_NCCL=ON` | Requires NVIDIA or Iluvatar NCCL. Currently available when `WITH_NVIDIA=ON` or `WITH_ILUVATAR=ON`.|
+| **NCCL** | Partial | `WITH_NCCL=ON` | Requires NVIDIA or Iluvatar NCCL 2.10 or newer. Currently available when `WITH_NVIDIA=ON` or `WITH_ILUVATAR=ON`.|
 | **MCCL** | Partial | `WITH_MCCL=ON` | Requires MetaX or Moore MCCL. Currently available when `WITH_METAX=ON` or `WITH_MOORE=ON`.|
 
 </details>
diff --git a/examples/ccl/send_recv.cc b/examples/ccl/send_recv.cc
new file mode 100644
index 0000000..3e12afb
--- /dev/null
+++ b/examples/ccl/send_recv.cc
@@ -0,0 +1,254 @@
+/**
+ * InfiniCCL Example: Thread-per-GPU Single-Node Send/Recv
+ *
+ * This example initializes one native CCL communicator rank per GPU thread
+ * and transfers data from rank 0 to rank 1 with `infinicclSend()` and
+ * `infinicclRecv()`. Other ranks participate in communicator initialization
+ * and teardown only.
+ *
+ * Run this example with `--launcher none`; it does not use MPI.
+ */
+
+#include <unistd.h>
+
+#include <array>
+#include <atomic>
+#include <charconv>
+#include <cstdlib>
+#include <cstring>
+#include <iostream>
+#include <limits>
+#include <string>
+#include <system_error>
+#include <thread>
+#include <vector>
+
+#include "backend_manifest.h"
+#include "infiniccl.h"
+#include "utils.h"
+
+namespace ccl = infini::ccl;
+
+namespace {
+
+constexpr int kSenderRank = 0;
+constexpr int kReceiverRank = 1;
+constexpr float kPayloadValue = 7.0f;
+
+struct RunState {
+  std::atomic<int> initialized{0};
+  std::atomic<int> completed{0};
+  bool correct = false;
+  float actual = 0.0f;
+  double elapsed_ms = 0.0;
+};
+
+struct ThreadArgs {
+  int rank;
+  int size;
+  infinicclUniqueId unique_id;
+  size_t num_elements;
+  int warmup_iters;
+  int profile_iters;
+  RunState *state;
+};
+
+template <typename T>
+bool ParseInteger(const char *text, T *value) {
+  if (!text || !value) {
+    return false;
+  }
+
+  const char *end = text + std::strlen(text);
+  auto result = std::from_chars(text, end, *value);
+  return result.ec == std::errc{} && result.ptr == end;
+}
+
+void WaitForAll(std::atomic<int> *counter, int size) {
+  counter->fetch_add(1, std::memory_order_acq_rel);
+  while (counter->load(std::memory_order_acquire) != size) {
+    std::this_thread::yield();
+  }
+}
+
+void RunWorker(ThreadArgs args) {
+  constexpr ccl::Device::Type kDeviceType =
+      ccl::ListGetBest<ccl::DevicePriority>(ccl::EnabledDevices{});
+  using Rt = ccl::Runtime<kDeviceType>;
+
+  CHECK_RT(Rt, Rt::SetDevice(args.rank));
+
+  infinicclComm_t comm = nullptr;
+  CHECK_INFINI(
+      infinicclCommInitRank(&comm, args.size, args.unique_id, args.rank));
+
+  const bool is_sender = args.rank == kSenderRank;
+  const bool is_receiver = args.rank == kReceiverRank;
+  const bool transfers_data = is_sender || is_receiver;
+  const size_t total_bytes = args.num_elements * sizeof(float);
+
+  std::vector<float> host_buffer(transfers_data ? args.num_elements : 0,
+                                 is_sender ? kPayloadValue : 0.0f);
+  float *device_buffer = nullptr;
+
+  if (transfers_data) {
+    CHECK_RT(
+        Rt, Rt::Malloc(reinterpret_cast<void **>(&device_buffer), total_bytes));
+    CHECK_RT(Rt, Rt::Memcpy(device_buffer, host_buffer.data(), total_bytes,
+                            Rt::MemcpyHostToDevice));
+    CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+  }
+
+  WaitForAll(&args.state->initialized, args.size);
+
+  auto transfer = [&]() {
+    if (is_sender) {
+      return infinicclSend(device_buffer, args.num_elements, infinicclFloat32,
+                           kReceiverRank, comm, nullptr);
+    }
+    if (is_receiver) {
+      return infinicclRecv(device_buffer, args.num_elements, infinicclFloat32,
+                           kSenderRank, comm, nullptr);
+    }
+    return infinicclSuccess;
+  };
+
+  for (int i = 0; i < args.warmup_iters; ++i) {
+    CHECK_INFINI(transfer());
+  }
+  CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+
+  Timer timer;
+  for (int i = 0; i < args.profile_iters; ++i) {
+    CHECK_INFINI(transfer());
+  }
+  CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+
+  if (is_sender) {
+    args.state->elapsed_ms =
+        timer.ElapsedMs() / static_cast<double>(args.profile_iters);
+  }
+
+  if (is_receiver) {
+    CHECK_RT(Rt, Rt::Memcpy(host_buffer.data(), device_buffer, total_bytes,
+                            Rt::MemcpyDeviceToHost));
+    args.state->correct = Validator::ValidateResult(
+        host_buffer.data(), args.num_elements, kPayloadValue, args.rank, false,
+        "CCL Send/Recv");
+    args.state->actual = host_buffer.front();
+  }
+
+  WaitForAll(&args.state->completed, args.size);
+
+  if (transfers_data) {
+    CHECK_RT(Rt, Rt::Free(device_buffer));
+  }
+  CHECK_INFINI(infinicclCommDestroy(comm));
+}
+
+void PrintUsage(const char *program) {
+  std::cout << "Usage: " << program << " [options]\n"
+            << "Options:\n"
+            << "  -g <num_gpus>        Number of GPUs (default: 8)\n"
+            << "  -w <warmup_iters>    Warm-up iterations (default: 2)\n"
+            << "  -p <profile_iters>   Profile iterations (default: 20)\n"
+            << "  -n <num_elements>    Number of elements (default: 1048576)\n";
+}
+
+}  // namespace
+
+int main(int argc, char **argv) {
+  int num_gpus = 8;
+  int warmup_iters = 2;
+  int profile_iters = 20;
+  size_t num_elements = 1 << 20;
+
+  int option = 0;
+  while ((option = getopt(argc, argv, "g:w:p:n:h")) != -1) {
+    bool valid = true;
+    switch (option) {
+      case 'g':
+        valid = ParseInteger(optarg, &num_gpus);
+        break;
+      case 'w':
+        valid = ParseInteger(optarg, &warmup_iters);
+        break;
+      case 'p':
+        valid = ParseInteger(optarg, &profile_iters);
+        break;
+      case 'n':
+        valid = ParseInteger(optarg, &num_elements);
+        break;
+      case 'h':
+        PrintUsage(argv[0]);
+        return EXIT_SUCCESS;
+      default:
+        valid = false;
+        break;
+    }
+
+    if (!valid) {
+      std::cerr << "Invalid numeric argument." << std::endl;
+      PrintUsage(argv[0]);
+      return EXIT_FAILURE;
+    }
+  }
+
+  if (optind != argc) {
+    std::cerr << "Unexpected positional argument." << std::endl;
+    PrintUsage(argv[0]);
+    return EXIT_FAILURE;
+  }
+
+  if (num_gpus < 2 || warmup_iters < 0 || profile_iters <= 0 ||
+      num_elements == 0 ||
+      num_elements > std::numeric_limits<size_t>::max() / sizeof(float)) {
+    std::cerr << "Send/Recv requires at least two GPUs, a non-negative warm-up "
+                 "count, a positive profile count, and a positive element "
+                 "count."
+              << std::endl;
+    return EXIT_FAILURE;
+  }
+
+  std::array<char, 256> hostname{};
+  if (gethostname(hostname.data(), hostname.size() - 1) != 0) {
+    std::cerr << "Failed to query the local hostname." << std::endl;
+    return EXIT_FAILURE;
+  }
+  std::cout << "[Main Process] Host: " << hostname.data()
+            << " | Target GPUs: " << num_gpus << " | Sender: " << kSenderRank
+            << " | Receiver: " << kReceiverRank << std::endl;
+
+  infinicclUniqueId unique_id{};
+  CHECK_INFINI(infinicclGetUniqueId(&unique_id));
+
+  RunState state;
+  std::vector<std::thread> workers;
+  workers.reserve(num_gpus);
+
+  for (int rank = 0; rank < num_gpus; ++rank) {
+    workers.emplace_back(RunWorker,
+                         ThreadArgs{rank, num_gpus, unique_id, num_elements,
+                                    warmup_iters, profile_iters, &state});
+  }
+
+  for (auto &worker : workers) {
+    worker.join();
+  }
+
+  const char *color = state.correct ? "\033[32m" : "\033[31m";
+  std::cout << "\n=== CCL Send/Recv Results ===" << std::endl;
+  std::cout << "Correct: " << color << (state.correct ? "YES" : "NO")
+            << "\033[0m" << std::endl;
+  std::cout << "Expect:  " << kPayloadValue << std::endl;
+  std::cout << "Actual:  " << state.actual << std::endl;
+
+  std::cout << "\n=== Single-Node Threaded Send/Recv Results ===" << std::endl;
+  std::cout << "Data size: " << num_elements << " floats ("
+            << num_elements * sizeof(float) / 1024 / 1024 << " MB)"
+            << std::endl;
+  Metrics{state.elapsed_ms, num_elements * sizeof(float), 2}.Print();
+
+  std::cout << "[Main Process] All worker threads joined." << std::endl;
+  return state.correct ? EXIT_SUCCESS : EXIT_FAILURE;
+}
diff --git a/examples/ccl_mpi_hybrid/send_recv.cc b/examples/ccl_mpi_hybrid/send_recv.cc
new file mode 100644
index 0000000..94e1d42
--- /dev/null
+++ b/examples/ccl_mpi_hybrid/send_recv.cc
@@ -0,0 +1,192 @@
+/**
+ * InfiniCCL Example: Send/Recv (OpenMPI + CCL Hybrid)
+ *
+ * OpenMPI launches the ranks and distributes the native CCL unique ID. After
+ * rank-based CCL communicator initialization, rank 0 sends GPU data to rank 1
+ * through the public `infinicclSend()` and `infinicclRecv()` APIs.
+ */
+
+#include <unistd.h>
+
+#include <array>
+#include <charconv>
+#include <cstdint>
+#include <cstdlib>
+#include <cstring>
+#include <iostream>
+#include <string>
+#include <system_error>
+#include <vector>
+
+#include "backend_manifest.h"
+#include "device.h"
+#include "infiniccl.h"
+#include "runtime.h"
+#include "traits.h"
+#include "utils.h"
+
+namespace ccl = infini::ccl;
+
+namespace {
+
+constexpr int kSenderRank = 0;
+constexpr int kReceiverRank = 1;
+constexpr int kRequiredRanks = 2;
+constexpr float kPayloadValue = 7.0f;
+
+bool ParseLocalRank(const char *text, int *local_rank) {
+  if (!text || !local_rank) {
+    return false;
+  }
+
+  const char *end = text + std::strlen(text);
+  auto result = std::from_chars(text, end, *local_rank);
+  return result.ec == std::errc{} && result.ptr == end && *local_rank >= 0;
+}
+
+}  // namespace
+
+int main(int argc, char **argv) {
+  constexpr int kWarmupIters = 2;
+  constexpr int kProfileIters = 20;
+  constexpr size_t kNumElements = 1 << 20;
+
+  constexpr ccl::Device::Type kDeviceType =
+      ccl::ListGetBest<ccl::DevicePriority>(ccl::EnabledDevices{});
+  using Rt = ccl::Runtime<kDeviceType>;
+
+  CHECK_INFINI(infinicclInit(&argc, &argv));
+
+  int rank = 0;
+  int size = 0;
+  CHECK_INFINI(infinicclGetRank(&rank));
+  CHECK_INFINI(infinicclGetSize(&size));
+
+  if (size < kRequiredRanks) {
+    if (rank == kSenderRank) {
+      std::cerr << "Hybrid Send/Recv requires at least two ranks." << std::endl;
+    }
+    CHECK_INFINI(infinicclFinalize());
+    return EXIT_FAILURE;
+  }
+
+  int local_rank = -1;
+  if (!ParseLocalRank(std::getenv("OMPI_COMM_WORLD_LOCAL_RANK"), &local_rank)) {
+    std::cerr << "Rank " << rank
+              << " received an invalid `OMPI_COMM_WORLD_LOCAL_RANK` value."
+              << std::endl;
+    std::exit(EXIT_FAILURE);
+  }
+  CHECK_RT(Rt, Rt::SetDevice(local_rank));
+
+  std::array<char, 256> hostname{};
+  if (gethostname(hostname.data(), hostname.size() - 1) != 0) {
+    std::cerr << "Rank " << rank << " failed to query the local hostname."
+              << std::endl;
+    std::exit(EXIT_FAILURE);
+  }
+  std::cout << "[Rank " << rank << "] Host: " << hostname.data()
+            << " | GPU: " << ccl::Device::StringFromType(kDeviceType)
+            << " | Device " << local_rank << std::endl;
+
+  infinicclComm_t comm = nullptr;
+  CHECK_INFINI(infinicclCommInitAll(&comm, size, nullptr));
+
+  infinicclUniqueId unique_id{};
+  if (rank == kSenderRank) {
+    CHECK_INFINI(infinicclGetUniqueId(&unique_id));
+  }
+  CHECK_INFINI(infinicclBroadcast(&unique_id, &unique_id, sizeof(unique_id),
+                                  infinicclChar, kSenderRank, comm, nullptr));
+  CHECK_INFINI(infinicclCommInitRank(&comm, size, unique_id, rank));
+
+  const bool is_sender = rank == kSenderRank;
+  const bool is_receiver = rank == kReceiverRank;
+  const bool transfers_data = is_sender || is_receiver;
+  const size_t total_bytes = kNumElements * sizeof(float);
+
+  std::vector<float> host_buffer(transfers_data ? kNumElements : 0,
+                                 is_sender ? kPayloadValue : 0.0f);
+  float *device_buffer = nullptr;
+
+  if (transfers_data) {
+    CHECK_RT(
+        Rt, Rt::Malloc(reinterpret_cast<void **>(&device_buffer), total_bytes));
+    CHECK_RT(Rt, Rt::Memcpy(device_buffer, host_buffer.data(), total_bytes,
+                            Rt::MemcpyHostToDevice));
+    CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+  }
+
+  auto transfer = [&]() {
+    if (is_sender) {
+      return infinicclSend(device_buffer, kNumElements, infinicclFloat32,
+                           kReceiverRank, comm, nullptr);
+    }
+    if (is_receiver) {
+      return infinicclRecv(device_buffer, kNumElements, infinicclFloat32,
+                           kSenderRank, comm, nullptr);
+    }
+    return infinicclSuccess;
+  };
+
+  for (int i = 0; i < kWarmupIters; ++i) {
+    CHECK_INFINI(transfer());
+  }
+  CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+
+  Timer timer;
+  for (int i = 0; i < kProfileIters; ++i) {
+    CHECK_INFINI(transfer());
+  }
+  CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+  double elapsed_ms = timer.ElapsedMs() / static_cast<double>(kProfileIters);
+
+  bool correct = true;
+  if (is_receiver) {
+    CHECK_RT(Rt, Rt::Memcpy(host_buffer.data(), device_buffer, total_bytes,
+                            Rt::MemcpyDeviceToHost));
+    correct = Validator::ValidateResult(host_buffer.data(), kNumElements,
+                                        kPayloadValue, rank, false,
+                                        "Hybrid CCL Send/Recv");
+    const char *color = correct ? "\033[32m" : "\033[31m";
+    std::cout << "\n=== Hybrid CCL Send/Recv Results ===" << std::endl;
+    std::cout << "Correct: " << color << (correct ? "YES" : "NO") << "\033[0m"
+              << std::endl;
+    std::cout << "Expect:  " << kPayloadValue << std::endl;
+    std::cout << "Actual:  " << host_buffer.front() << std::endl;
+  }
+
+  // The receiver publishes validation status from device memory. This works
+  // with the current OpenMPI staging path and with native CCL `Broadcast`.
+  std::int32_t completion_token = is_receiver && correct ? 1 : 0;
+  std::int32_t *device_completion = nullptr;
+  CHECK_RT(Rt, Rt::Malloc(reinterpret_cast<void **>(&device_completion),
+                          sizeof(completion_token)));
+  CHECK_RT(Rt, Rt::Memcpy(device_completion, &completion_token,
+                          sizeof(completion_token), Rt::MemcpyHostToDevice));
+  CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+  CHECK_INFINI(infinicclBroadcast(device_completion, device_completion, 1,
+                                  infinicclInt32, kReceiverRank, comm,
+                                  nullptr));
+  CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+  CHECK_RT(Rt, Rt::Memcpy(&completion_token, device_completion,
+                          sizeof(completion_token), Rt::MemcpyDeviceToHost));
+  correct = completion_token == 1;
+
+  if (is_sender) {
+    std::cout << "\n=== OpenMPI-Assisted CCL Send/Recv Results ==="
+              << std::endl;
+    std::cout << "Data size: " << kNumElements << " floats ("
+              << total_bytes / 1024 / 1024 << " MB)" << std::endl;
+    Metrics{elapsed_ms, total_bytes, kRequiredRanks}.Print();
+  }
+
+  if (transfers_data) {
+    CHECK_RT(Rt, Rt::Free(device_buffer));
+  }
+  CHECK_RT(Rt, Rt::Free(device_completion));
+  CHECK_INFINI(infinicclCommDestroy(comm));
+  CHECK_INFINI(infinicclFinalize());
+
+  return correct ? EXIT_SUCCESS : EXIT_FAILURE;
+}
diff --git a/src/backends/ccl/common/impl/point_to_point.h b/src/backends/ccl/common/impl/point_to_point.h
new file mode 100644
index 0000000..ae108e3
--- /dev/null
+++ b/src/backends/ccl/common/impl/point_to_point.h
@@ -0,0 +1,48 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_POINT_TO_POINT_H_
+#define INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_POINT_TO_POINT_H_
+
+#include "backends/ccl/common/api.h"
+#include "backends/ccl/common/comm_instance.h"
+#include "communicator.h"
+#include "data_type_impl.h"
+#include "return_status_impl.h"
+
+namespace infini::ccl {
+
+template <BackendType backend, Device::Type device>
+class CclPointToPoint {
+ public:
+  static bool HasNativeCommunicator(const Communicator *comm) {
+    return comm && comm->intra_comm() &&
+           comm->intra_comm_backend() == backend &&
+           comm->device_type() == device;
+  }
+
+  template <typename Callback>
+  static ReturnStatus Call(DataType data_type, Communicator *comm,
+                           Callback callback) {
+    using Api = CclApi<backend, device>;
+    using TypeMap = CclTypeMap<backend, device>;
+    using CommInstance = CclCommInstance<Api>;
+
+    if (!HasNativeCommunicator(comm)) {
+      return ReturnStatus::kInternalError;
+    }
+
+    auto *instance = static_cast<CommInstance *>(comm->intra_comm());
+    if (!instance->handle) {
+      return ReturnStatus::kInternalError;
+    }
+
+    typename Api::DataType native_type{};
+    if (!TypeMap::ToBackendDataType(data_type, &native_type)) {
+      return ReturnStatus::kNotSupported;
+    }
+
+    return Api::Check(callback(native_type, instance->handle));
+  }
+};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_POINT_TO_POINT_H_
diff --git a/src/backends/ccl/common/impl/recv.h b/src/backends/ccl/common/impl/recv.h
new file mode 100644
index 0000000..b4a8c75
--- /dev/null
+++ b/src/backends/ccl/common/impl/recv.h
@@ -0,0 +1,41 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_RECV_H_
+#define INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_RECV_H_
+
+#include "backends/ccl/common/impl/point_to_point.h"
+#include "base/recv.h"
+
+namespace infini::ccl {
+
+template <BackendType backend, Device::Type device>
+class CclRecvImpl {
+ public:
+  static ReturnStatus Apply(void *recv_buff, size_t count, DataType data_type,
+                            int peer, Communicator *comm, void *stream) {
+    using Api = CclApi<backend, device>;
+    using PointToPoint = CclPointToPoint<backend, device>;
+
+    if (!PointToPoint::HasNativeCommunicator(comm)) {
+      if (!comm || !comm->inter_comm() ||
+          comm->inter_comm_backend() != BackendType::kOmpi) {
+        return ReturnStatus::kInternalError;
+      }
+
+      if constexpr (BackendEnabled<Recv, BackendType::kOmpi>::value) {
+        return RecvImpl<BackendType::kOmpi, device>::Apply(
+            recv_buff, count, data_type, peer, comm, stream);
+      }
+
+      return ReturnStatus::kInternalError;
+    }
+
+    return PointToPoint::Call(
+        data_type, comm, [&](auto native_type, auto native_comm) {
+          return Api::Recv(recv_buff, count, native_type, peer, native_comm,
+                           reinterpret_cast<typename Api::Stream>(stream));
+        });
+  }
+};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_RECV_H_
diff --git a/src/backends/ccl/common/impl/send.h b/src/backends/ccl/common/impl/send.h
new file mode 100644
index 0000000..a681192
--- /dev/null
+++ b/src/backends/ccl/common/impl/send.h
@@ -0,0 +1,42 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_SEND_H_
+#define INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_SEND_H_
+
+#include "backends/ccl/common/impl/point_to_point.h"
+#include "base/send.h"
+
+namespace infini::ccl {
+
+template <BackendType backend, Device::Type device>
+class CclSendImpl {
+ public:
+  static ReturnStatus Apply(const void *send_buff, size_t count,
+                            DataType data_type, int peer, Communicator *comm,
+                            void *stream) {
+    using Api = CclApi<backend, device>;
+    using PointToPoint = CclPointToPoint<backend, device>;
+
+    if (!PointToPoint::HasNativeCommunicator(comm)) {
+      if (!comm || !comm->inter_comm() ||
+          comm->inter_comm_backend() != BackendType::kOmpi) {
+        return ReturnStatus::kInternalError;
+      }
+
+      if constexpr (BackendEnabled<Send, BackendType::kOmpi>::value) {
+        return SendImpl<BackendType::kOmpi, device>::Apply(
+            send_buff, count, data_type, peer, comm, stream);
+      }
+
+      return ReturnStatus::kInternalError;
+    }
+
+    return PointToPoint::Call(
+        data_type, comm, [&](auto native_type, auto native_comm) {
+          return Api::Send(send_buff, count, native_type, peer, native_comm,
+                           reinterpret_cast<typename Api::Stream>(stream));
+        });
+  }
+};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_SEND_H_
diff --git a/src/backends/ccl/mccl/api.h b/src/backends/ccl/mccl/api.h
index a5d2bcd..85baefa 100644
--- a/src/backends/ccl/mccl/api.h
+++ b/src/backends/ccl/mccl/api.h
@@ -49,6 +49,16 @@ struct McclApi {
     return mcclAllReduce(send_buff, recv_buff, count, data_type, op, comm,
                          stream);
   }
+
+  static Result Send(const void *send_buff, size_t count, DataType data_type,
+                     int peer, Comm comm, Stream stream) {
+    return mcclSend(send_buff, count, data_type, peer, comm, stream);
+  }
+
+  static Result Recv(void *recv_buff, size_t count, DataType data_type,
+                     int peer, Comm comm, Stream stream) {
+    return mcclRecv(recv_buff, count, data_type, peer, comm, stream);
+  }
 };
 
 }  // namespace infini::ccl
diff --git a/src/backends/ccl/mccl/impl/recv.h b/src/backends/ccl/mccl/impl/recv.h
new file mode 100644
index 0000000..889a956
--- /dev/null
+++ b/src/backends/ccl/mccl/impl/recv.h
@@ -0,0 +1,17 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_RECV_H_
+#define INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_RECV_H_
+
+#include "backends/ccl/common/impl/recv.h"
+
+namespace infini::ccl {
+
+template <Device::Type device>
+class RecvImpl<BackendType::kMccl, device>
+    : public CclRecvImpl<BackendType::kMccl, device> {};
+
+template <>
+struct BackendEnabled<Recv, BackendType::kMccl> : std::true_type {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_RECV_H_
diff --git a/src/backends/ccl/mccl/impl/send.h b/src/backends/ccl/mccl/impl/send.h
new file mode 100644
index 0000000..2c59a85
--- /dev/null
+++ b/src/backends/ccl/mccl/impl/send.h
@@ -0,0 +1,17 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_SEND_H_
+#define INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_SEND_H_
+
+#include "backends/ccl/common/impl/send.h"
+
+namespace infini::ccl {
+
+template <Device::Type device>
+class SendImpl<BackendType::kMccl, device>
+    : public CclSendImpl<BackendType::kMccl, device> {};
+
+template <>
+struct BackendEnabled<Send, BackendType::kMccl> : std::true_type {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_SEND_H_
diff --git a/src/backends/ccl/nccl/api.h b/src/backends/ccl/nccl/api.h
index e7b6119..81dda6f 100644
--- a/src/backends/ccl/nccl/api.h
+++ b/src/backends/ccl/nccl/api.h
@@ -10,6 +10,12 @@
 #include "return_status_impl.h"
 #include "runtime.h"
 
+#if !defined(NCCL_VERSION_CODE) || !defined(NCCL_VERSION)
+#error "InfiniCCL NCCL support requires NCCL 2.10.0 or newer."
+#elif NCCL_VERSION_CODE < NCCL_VERSION(2, 10, 0)
+#error "InfiniCCL NCCL support requires NCCL 2.10.0 or newer."
+#endif
+
 namespace infini::ccl {
 
 template <Device::Type device>
@@ -46,6 +52,16 @@ struct NcclApi {
     return ncclAllReduce(send_buff, recv_buff, count, data_type, op, comm,
                          stream);
   }
+
+  static Result Send(const void *send_buff, size_t count, DataType data_type,
+                     int peer, Comm comm, Stream stream) {
+    return ncclSend(send_buff, count, data_type, peer, comm, stream);
+  }
+
+  static Result Recv(void *recv_buff, size_t count, DataType data_type,
+                     int peer, Comm comm, Stream stream) {
+    return ncclRecv(recv_buff, count, data_type, peer, comm, stream);
+  }
 };
 
 }  // namespace infini::ccl
diff --git a/src/backends/ccl/nccl/impl/recv.h b/src/backends/ccl/nccl/impl/recv.h
new file mode 100644
index 0000000..dae63b5
--- /dev/null
+++ b/src/backends/ccl/nccl/impl/recv.h
@@ -0,0 +1,17 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_RECV_H_
+#define INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_RECV_H_
+
+#include "backends/ccl/common/impl/recv.h"
+
+namespace infini::ccl {
+
+template <Device::Type device>
+class RecvImpl<BackendType::kNccl, device>
+    : public CclRecvImpl<BackendType::kNccl, device> {};
+
+template <>
+struct BackendEnabled<Recv, BackendType::kNccl> : std::true_type {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_RECV_H_
diff --git a/src/backends/ccl/nccl/impl/send.h b/src/backends/ccl/nccl/impl/send.h
new file mode 100644
index 0000000..d91cb00
--- /dev/null
+++ b/src/backends/ccl/nccl/impl/send.h
@@ -0,0 +1,17 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_SEND_H_
+#define INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_SEND_H_
+
+#include "backends/ccl/common/impl/send.h"
+
+namespace infini::ccl {
+
+template <Device::Type device>
+class SendImpl<BackendType::kNccl, device>
+    : public CclSendImpl<BackendType::kNccl, device> {};
+
+template <>
+struct BackendEnabled<Send, BackendType::kNccl> : std::true_type {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_SEND_H_