diff --git a/examples/ccl/broadcast.cc b/examples/ccl/broadcast.cc
new file mode 100644
index 0000000..5396a93
--- /dev/null
+++ b/examples/ccl/broadcast.cc
@@ -0,0 +1,242 @@
+/**
+ * InfiniCCL Example: Thread-per-GPU Single-Node Broadcast
+ *
+ * This example spawns one CPU thread per GPU and performs native CCL
+ * broadcasts without an MPI launcher.
+ */
+
+#include <unistd.h>
+
+#include <algorithm>
+#include <array>
+#include <atomic>
+#include <charconv>
+#include <cstdlib>
+#include <iostream>
+#include <limits>
+#include <string>
+#include <string_view>
+#include <thread>
+#include <vector>
+
+#include "backend_manifest.h"
+#include "infiniccl.h"
+#include "utils.h"
+
+using namespace infini::ccl;
+
+struct ThreadArgs {
+  int rank;
+  int size;
+  int root;
+  infinicclUniqueId id;
+  size_t num_elements;
+  int warmup_iter;
+  int profile_iter;
+  std::atomic_bool *all_correct;
+};
+
+template <typename T>
+bool ParseIntegerOption(const char *argument, T *value) {
+  if (!argument || !value) {
+    return false;
+  }
+
+  const std::string_view text(argument);
+  if (text.empty()) {
+    return false;
+  }
+
+  T parsed{};
+  const auto [end, error] =
+      std::from_chars(text.data(), text.data() + text.size(), parsed);
+  if (error != std::errc{} || end != text.data() + text.size()) {
+    return false;
+  }
+
+  *value = parsed;
+  return true;
+}
+
+void WorkerThread(ThreadArgs args) {
+  constexpr Device::Type kDevType =
+      ListGetBest<DevicePriority>(EnabledDevices{});
+  using Rt = Runtime<kDevType>;
+
+  CHECK_RT(Rt, Rt::SetDevice(args.rank));
+
+  infinicclComm_t comm = nullptr;
+  CHECK_INFINI(infinicclCommInitRank(&comm, args.size, args.id, args.rank));
+
+  constexpr float kRootValue = 42.0f;
+  constexpr float kSentinelValue = -1.0f;
+  const size_t total_bytes = args.num_elements * sizeof(float);
+
+  std::vector<float> h_send(args.num_elements, kSentinelValue);
+  std::vector<float> h_recv(args.num_elements, kSentinelValue);
+
+  float *d_send = nullptr;
+  float *d_recv = nullptr;
+  CHECK_RT(Rt, Rt::Malloc(reinterpret_cast<void **>(&d_send), total_bytes));
+  CHECK_RT(Rt, Rt::Malloc(reinterpret_cast<void **>(&d_recv), total_bytes));
+
+  auto ResetBuffers = [&]() {
+    std::fill(h_send.begin(), h_send.end(),
+              args.rank == args.root ? kRootValue : kSentinelValue);
+    std::fill(h_recv.begin(), h_recv.end(), kSentinelValue);
+    CHECK_RT(Rt, Rt::Memcpy(d_send, h_send.data(), total_bytes,
+                            Rt::MemcpyHostToDevice));
+    CHECK_RT(Rt, Rt::Memcpy(d_recv, h_recv.data(), total_bytes,
+                            Rt::MemcpyHostToDevice));
+    CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+  };
+
+  auto RunScenario = [&](const std::string &name, float *verify_buff,
+                         auto &&collective_call) {
+    for (int i = 0; i < args.warmup_iter; ++i) {
+      CHECK_INFINI(collective_call());
+    }
+    CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+
+    Timer timer;
+    for (int i = 0; i < args.profile_iter; ++i) {
+      CHECK_INFINI(collective_call());
+    }
+    CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+    const double elapsed =
+        timer.ElapsedMs() / static_cast<double>(args.profile_iter);
+
+    CHECK_RT(Rt, Rt::Memcpy(h_recv.data(), verify_buff, total_bytes,
+                            Rt::MemcpyDeviceToHost));
+
+    const bool correct = Validator::ValidateResult(
+        h_recv.data(), args.num_elements, kRootValue, args.rank, true, name);
+    if (!correct) {
+      args.all_correct->store(false, std::memory_order_relaxed);
+    }
+
+    if (args.rank == 0) {
+      Metrics metrics{elapsed, total_bytes, args.size};
+      metrics.Print();
+    }
+  };
+
+  ResetBuffers();
+  RunScenario("Out-of-Place Broadcast", d_recv, [&]() {
+    return infinicclBroadcast(d_send, d_recv, args.num_elements,
+                              infinicclFloat32, args.root, comm, nullptr);
+  });
+
+  ResetBuffers();
+  float *d_in_place = args.rank == args.root ? d_send : d_recv;
+  RunScenario("In-Place Broadcast", d_in_place, [&]() {
+    return infinicclBroadcast(d_in_place, d_in_place, args.num_elements,
+                              infinicclFloat32, args.root, comm, nullptr);
+  });
+
+  ResetBuffers();
+  d_in_place = args.rank == args.root ? d_send : d_recv;
+  RunScenario("Legacy In-Place Bcast", d_in_place, [&]() {
+    return infinicclBcast(d_in_place, args.num_elements, infinicclFloat32,
+                          args.root, comm, nullptr);
+  });
+
+  CHECK_RT(Rt, Rt::Free(d_send));
+  CHECK_RT(Rt, Rt::Free(d_recv));
+  CHECK_INFINI(infinicclCommDestroy(comm));
+}
+
+int main(int argc, char **argv) {
+  int num_gpus = 8;
+  int warmup_iters = 1;
+  int profile_iters = 20;
+  size_t num_elements = 1 << 25;
+
+  int opt;
+  while ((opt = getopt(argc, argv, "g:w:p:n:h")) != -1) {
+    switch (opt) {
+      case 'g':
+        if (!ParseIntegerOption(optarg, &num_gpus)) {
+          std::cerr << "Invalid value for `-g`." << std::endl;
+          return EXIT_FAILURE;
+        }
+        break;
+      case 'w':
+        if (!ParseIntegerOption(optarg, &warmup_iters)) {
+          std::cerr << "Invalid value for `-w`." << std::endl;
+          return EXIT_FAILURE;
+        }
+        break;
+      case 'p':
+        if (!ParseIntegerOption(optarg, &profile_iters)) {
+          std::cerr << "Invalid value for `-p`." << std::endl;
+          return EXIT_FAILURE;
+        }
+        break;
+      case 'n':
+        if (!ParseIntegerOption(optarg, &num_elements)) {
+          std::cerr << "Invalid value for `-n`." << std::endl;
+          return EXIT_FAILURE;
+        }
+        break;
+      case 'h':
+        std::cout << "Usage: " << argv[0] << " [options]\n"
+                  << "Options:\n"
+                  << "  -g <num_gpus>        Number of GPUs (default: 8)\n"
+                  << "  -w <warmup_iters>    Warmup iterations (default: 1)\n"
+                  << "  -p <profile_iters>   Profile iterations (default: 20)\n"
+                  << "  -n <num_elements>    Number of elements (default: "
+                  << (1 << 25) << ")\n";
+        return EXIT_SUCCESS;
+      default:
+        std::cerr << "Invalid argument. Use `-h` for help." << std::endl;
+        return EXIT_FAILURE;
+    }
+  }
+
+  if (num_gpus <= 0 || warmup_iters < 0 || profile_iters <= 0 ||
+      num_elements == 0 ||
+      num_elements > std::numeric_limits<size_t>::max() / sizeof(float)) {
+    std::cerr << "Invalid execution parameter." << std::endl;
+    return EXIT_FAILURE;
+  }
+
+  std::array<char, 256> hostname{};
+  if (gethostname(hostname.data(), hostname.size()) != 0) {
+    std::cerr << "Failed to query the host name." << std::endl;
+    return EXIT_FAILURE;
+  }
+  hostname.back() = '\0';
+
+  const int root = num_gpus > 1 ? num_gpus - 1 : 0;
+  std::cout << "[Main Process] Host: " << hostname.data()
+            << " | Target GPUs: " << num_gpus << " | Root: " << root
+            << std::endl;
+
+  infinicclUniqueId shared_id;
+  CHECK_INFINI(infinicclGetUniqueId(&shared_id));
+
+  std::atomic_bool all_correct{true};
+  std::vector<std::thread> threads;
+  threads.reserve(num_gpus);
+
+  for (int rank = 0; rank < num_gpus; ++rank) {
+    ThreadArgs args{rank,         num_gpus,     root,          shared_id,
+                    num_elements, warmup_iters, profile_iters, &all_correct};
+    threads.emplace_back(WorkerThread, args);
+  }
+
+  for (auto &thread : threads) {
+    if (thread.joinable()) {
+      thread.join();
+    }
+  }
+
+  if (!all_correct.load(std::memory_order_relaxed)) {
+    std::cerr << "Broadcast validation failed." << std::endl;
+    return EXIT_FAILURE;
+  }
+
+  std::cout << "[Main Process] All broadcast scenarios passed." << std::endl;
+  return EXIT_SUCCESS;
+}
diff --git a/examples/ccl_mpi_hybrid/broadcast.cc b/examples/ccl_mpi_hybrid/broadcast.cc
new file mode 100644
index 0000000..e82fe6d
--- /dev/null
+++ b/examples/ccl_mpi_hybrid/broadcast.cc
@@ -0,0 +1,171 @@
+/**
+ * InfiniCCL Example: Broadcast (OpenMPI + CCL Hybrid)
+ *
+ * This example uses OpenMPI to distribute a CCL unique ID across processes,
+ * initializes one native CCL rank per GPU, and then performs CCL broadcasts.
+ */
+
+#include <unistd.h>
+
+#include <algorithm>
+#include <array>
+#include <cstdlib>
+#include <iostream>
+#include <string>
+#include <vector>
+
+#include "backend_manifest.h"
+#include "device.h"
+#include "infiniccl.h"
+#include "runtime.h"
+#include "traits.h"
+#include "utils.h"
+
+using namespace infini::ccl;
+
+bool RunBroadcastExample(int argc, char **argv, int warmup_iter,
+                         int profile_iter, size_t num_elements) {
+  constexpr Device::Type kDevType =
+      ListGetBest<DevicePriority>(EnabledDevices{});
+  using Rt = Runtime<kDevType>;
+
+  CHECK_INFINI(infinicclInit(&argc, &argv));
+
+  int rank = 0;
+  int size = 0;
+  CHECK_INFINI(infinicclGetRank(&rank));
+  CHECK_INFINI(infinicclGetSize(&size));
+
+  std::array<char, 256> hostname{};
+  if (gethostname(hostname.data(), hostname.size()) != 0) {
+    constexpr char kUnknownHostname[] = "unknown";
+    std::copy_n(kUnknownHostname, sizeof(kUnknownHostname), hostname.begin());
+    std::cerr << "[Rank " << rank
+              << "] Failed to query the host name; using `unknown`."
+              << std::endl;
+  }
+  hostname.back() = '\0';
+
+  const char *local_rank_text = std::getenv("OMPI_COMM_WORLD_LOCAL_RANK");
+  int local_rank = 0;
+  if (local_rank_text) {
+    local_rank = std::atoi(local_rank_text);
+  }
+  CHECK_RT(Rt, Rt::SetDevice(local_rank));
+
+  infinicclComm_t comm = nullptr;
+  CHECK_INFINI(infinicclCommInitAll(&comm, size, nullptr));
+
+  infinicclUniqueId id{};
+  if (rank == 0) {
+    CHECK_INFINI(infinicclGetUniqueId(&id));
+  }
+  CHECK_INFINI(infinicclBroadcast(&id, &id, sizeof(id), infinicclChar, 0, comm,
+                                  nullptr));
+
+  CHECK_INFINI(infinicclCommInitRank(&comm, size, id, rank));
+
+  const int root = size > 1 ? size - 1 : 0;
+  std::cout << "[Rank " << rank << "] Host: " << hostname.data()
+            << " | GPU: " << Device::StringFromType(kDevType) << " | Device "
+            << local_rank << " | Broadcast Root: " << root << std::endl;
+
+  constexpr float kRootValue = 42.0f;
+  constexpr float kSentinelValue = -1.0f;
+  const size_t total_bytes = num_elements * sizeof(float);
+
+  std::vector<float> h_send(num_elements, kSentinelValue);
+  std::vector<float> h_recv(num_elements, kSentinelValue);
+
+  float *d_send = nullptr;
+  float *d_recv = nullptr;
+  CHECK_RT(Rt, Rt::Malloc(reinterpret_cast<void **>(&d_send), total_bytes));
+  CHECK_RT(Rt, Rt::Malloc(reinterpret_cast<void **>(&d_recv), total_bytes));
+
+  auto ResetBuffers = [&]() {
+    std::fill(h_send.begin(), h_send.end(),
+              rank == root ? kRootValue : kSentinelValue);
+    std::fill(h_recv.begin(), h_recv.end(), kSentinelValue);
+    CHECK_RT(Rt, Rt::Memcpy(d_send, h_send.data(), total_bytes,
+                            Rt::MemcpyHostToDevice));
+    CHECK_RT(Rt, Rt::Memcpy(d_recv, h_recv.data(), total_bytes,
+                            Rt::MemcpyHostToDevice));
+    CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+  };
+
+  auto RunScenario = [&](const std::string &name, float *verify_buff,
+                         auto &&collective_call) {
+    for (int i = 0; i < warmup_iter; ++i) {
+      CHECK_INFINI(collective_call());
+    }
+    CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+
+    Timer timer;
+    for (int i = 0; i < profile_iter; ++i) {
+      CHECK_INFINI(collective_call());
+    }
+    CHECK_RT(Rt, Rt::StreamSynchronize(nullptr));
+    const double elapsed =
+        timer.ElapsedMs() / static_cast<double>(profile_iter);
+
+    CHECK_RT(Rt, Rt::Memcpy(h_recv.data(), verify_buff, total_bytes,
+                            Rt::MemcpyDeviceToHost));
+
+    const bool correct = Validator::ValidateResult(
+        h_recv.data(), num_elements, kRootValue, rank, true, name);
+    if (rank == 0) {
+      Metrics metrics{elapsed, total_bytes, size};
+      metrics.Print();
+    }
+
+    return correct;
+  };
+
+  bool all_correct = true;
+
+  ResetBuffers();
+  all_correct &=
+      RunScenario("Out-of-Place Hybrid CCL Broadcast", d_recv, [&]() {
+        return infinicclBroadcast(d_send, d_recv, num_elements,
+                                  infinicclFloat32, root, comm, nullptr);
+      });
+
+  ResetBuffers();
+  float *d_in_place = rank == root ? d_send : d_recv;
+  all_correct &=
+      RunScenario("In-Place Hybrid CCL Broadcast", d_in_place, [&]() {
+        return infinicclBroadcast(d_in_place, d_in_place, num_elements,
+                                  infinicclFloat32, root, comm, nullptr);
+      });
+
+  ResetBuffers();
+  d_in_place = rank == root ? d_send : d_recv;
+  all_correct &=
+      RunScenario("Legacy In-Place Hybrid CCL Bcast", d_in_place, [&]() {
+        return infinicclBcast(d_in_place, num_elements, infinicclFloat32, root,
+                              comm, nullptr);
+      });
+
+  CHECK_RT(Rt, Rt::Free(d_send));
+  CHECK_RT(Rt, Rt::Free(d_recv));
+  CHECK_INFINI(infinicclCommDestroy(comm));
+  CHECK_INFINI(infinicclFinalize());
+
+  if (!all_correct) {
+    std::cerr << "[Rank " << rank << "] Hybrid CCL broadcast validation failed."
+              << std::endl;
+  }
+
+  return all_correct;
+}
+
+int main(int argc, char **argv) {
+  constexpr int kWarmupIterations = 2;
+  constexpr int kProfileIterations = 20;
+  constexpr size_t kNumElements = 1 << 20;
+
+  return RunBroadcastExample(argc, argv, kWarmupIterations, kProfileIterations,
+                             kNumElements)
+             ? EXIT_SUCCESS
+             : EXIT_FAILURE;
+}
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 06df487..eb36f4a 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -16,6 +16,19 @@ set(AUTOGEN_WARNING [[/*
 
 file(GLOB CORE_SRCS "*.cc")
 file(GLOB_RECURSE BASE_IMPL_SRCS "base/*.cc")
+file(GLOB_RECURSE BRIDGE_DEP_HEADERS CONFIGURE_DEPENDS
+    "base/*.h"
+    "backends/*.h"
+    "devices/*.h"
+    "devices/*.cuh"
+)
+string(REPLACE ";" "\n" BRIDGE_DEPENDENCY_CONTENT "${BRIDGE_DEP_HEADERS}")
+set(BRIDGE_DEPENDENCY_MANIFEST
+    "${CMAKE_CURRENT_BINARY_DIR}/bridge_dependencies.txt"
+)
+file(GENERATE OUTPUT "${BRIDGE_DEPENDENCY_MANIFEST}"
+    CONTENT "${BRIDGE_DEPENDENCY_CONTENT}\n"
+)
 
 target_sources(infiniccl PRIVATE
     ${CORE_SRCS}
@@ -280,7 +293,9 @@ add_custom_command(
             "${BACK_STR}"
     DEPENDS "${PROJECT_SOURCE_DIR}/scripts/gen_bridge.py"
             "${PROJECT_SOURCE_DIR}/include/comm.h"
+            "${BRIDGE_DEPENDENCY_MANIFEST}"
             ${BASE_IMPL_SRCS}
+            ${BRIDGE_DEP_HEADERS}
     VERBATIM
     COMMENT "Generating InfiniCCL bridge and manifest files for Devices: [${DEV_STR}] Backends: [${BACK_STR}]..."
 )
diff --git a/src/backends/ccl/common/impl/broadcast.h b/src/backends/ccl/common/impl/broadcast.h
new file mode 100644
index 0000000..c5df9a8
--- /dev/null
+++ b/src/backends/ccl/common/impl/broadcast.h
@@ -0,0 +1,50 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_BROADCAST_H_
+#define INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_BROADCAST_H_
+
+#include "backends/ccl/common/api.h"
+#include "backends/ccl/common/comm_instance.h"
+#include "base/broadcast.h"
+#include "communicator.h"
+
+namespace infini::ccl {
+
+template <BackendType backend, Device::Type device>
+class CclBroadcastImpl {
+ public:
+  static ReturnStatus Apply(const void *send_buff, void *recv_buff,
+                            size_t count, DataType data_type, int root,
+                            Communicator *comm, void *stream) {
+    using Api = CclApi<backend, device>;
+    using TypeMap = CclTypeMap<backend, device>;
+    using CommInstance = CclCommInstance<Api>;
+
+    auto *comm_internal = static_cast<Communicator *>(comm);
+    if (!comm_internal) {
+      return ReturnStatus::kInternalError;
+    }
+
+    if (!comm_internal->intra_comm() ||
+        comm_internal->intra_comm_backend() != backend ||
+        comm_internal->device_type() != device) {
+      return ReturnStatus::kInternalError;
+    }
+
+    auto *intra = static_cast<CommInstance *>(comm_internal->intra_comm());
+    if (!intra->handle) {
+      return ReturnStatus::kInternalError;
+    }
+
+    typename Api::DataType ccl_type{};
+    if (!TypeMap::ToBackendDataType(data_type, &ccl_type)) {
+      return ReturnStatus::kNotSupported;
+    }
+
+    return Api::Check(Api::Broadcast(
+        send_buff, recv_buff, count, ccl_type, root, intra->handle,
+        reinterpret_cast<typename Api::Stream>(stream)));
+  }
+};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_COMMON_IMPL_BROADCAST_H_
diff --git a/src/backends/ccl/mccl/api.h b/src/backends/ccl/mccl/api.h
index a5d2bcd..af65375 100644
--- a/src/backends/ccl/mccl/api.h
+++ b/src/backends/ccl/mccl/api.h
@@ -49,6 +49,13 @@ struct McclApi {
     return mcclAllReduce(send_buff, recv_buff, count, data_type, op, comm,
                          stream);
   }
+
+  static Result Broadcast(const void *send_buff, void *recv_buff, size_t count,
+                          DataType data_type, int root, Comm comm,
+                          Stream stream) {
+    return mcclBroadcast(send_buff, recv_buff, count, data_type, root, comm,
+                         stream);
+  }
 };
 
 }  // namespace infini::ccl
diff --git a/src/backends/ccl/mccl/impl/bcast.h b/src/backends/ccl/mccl/impl/bcast.h
new file mode 100644
index 0000000..79cb470
--- /dev/null
+++ b/src/backends/ccl/mccl/impl/bcast.h
@@ -0,0 +1,13 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_BCAST_H_
+#define INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_BCAST_H_
+
+#include "base/bcast.h"
+
+namespace infini::ccl {
+
+template <>
+struct BackendEnabled<Bcast, BackendType::kMccl> : std::true_type {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_BCAST_H_
diff --git a/src/backends/ccl/mccl/impl/broadcast.h b/src/backends/ccl/mccl/impl/broadcast.h
new file mode 100644
index 0000000..c2c497a
--- /dev/null
+++ b/src/backends/ccl/mccl/impl/broadcast.h
@@ -0,0 +1,17 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_BROADCAST_H_
+#define INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_BROADCAST_H_
+
+#include "backends/ccl/common/impl/broadcast.h"
+
+namespace infini::ccl {
+
+template <Device::Type device>
+class BroadcastImpl<BackendType::kMccl, device>
+    : public CclBroadcastImpl<BackendType::kMccl, device> {};
+
+template <>
+struct BackendEnabled<Broadcast, BackendType::kMccl> : std::true_type {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_MCCL_IMPL_BROADCAST_H_
diff --git a/src/backends/ccl/nccl/api.h b/src/backends/ccl/nccl/api.h
index e7b6119..707a168 100644
--- a/src/backends/ccl/nccl/api.h
+++ b/src/backends/ccl/nccl/api.h
@@ -46,6 +46,13 @@ struct NcclApi {
     return ncclAllReduce(send_buff, recv_buff, count, data_type, op, comm,
                          stream);
   }
+
+  static Result Broadcast(const void *send_buff, void *recv_buff, size_t count,
+                          DataType data_type, int root, Comm comm,
+                          Stream stream) {
+    return ncclBroadcast(send_buff, recv_buff, count, data_type, root, comm,
+                         stream);
+  }
 };
 
 }  // namespace infini::ccl
diff --git a/src/backends/ccl/nccl/impl/bcast.h b/src/backends/ccl/nccl/impl/bcast.h
new file mode 100644
index 0000000..73fe74c
--- /dev/null
+++ b/src/backends/ccl/nccl/impl/bcast.h
@@ -0,0 +1,13 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_BCAST_H_
+#define INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_BCAST_H_
+
+#include "base/bcast.h"
+
+namespace infini::ccl {
+
+template <>
+struct BackendEnabled<Bcast, BackendType::kNccl> : std::true_type {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_BCAST_H_
diff --git a/src/backends/ccl/nccl/impl/broadcast.h b/src/backends/ccl/nccl/impl/broadcast.h
new file mode 100644
index 0000000..88bbcfb
--- /dev/null
+++ b/src/backends/ccl/nccl/impl/broadcast.h
@@ -0,0 +1,17 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_BROADCAST_H_
+#define INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_BROADCAST_H_
+
+#include "backends/ccl/common/impl/broadcast.h"
+
+namespace infini::ccl {
+
+template <Device::Type device>
+class BroadcastImpl<BackendType::kNccl, device>
+    : public CclBroadcastImpl<BackendType::kNccl, device> {};
+
+template <>
+struct BackendEnabled<Broadcast, BackendType::kNccl> : std::true_type {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_NCCL_IMPL_BROADCAST_H_
diff --git a/src/base/broadcast.h b/src/base/broadcast.h
index 0924e78..6f98598 100644
--- a/src/base/broadcast.h
+++ b/src/base/broadcast.h
@@ -31,11 +31,38 @@ class Broadcast : public Operation<Broadcast> {
       return ReturnStatus::kSuccess;
     }
 
+    if (!comm->HasBackend(backend_type) || comm->device_type() != device_type) {
+      BackendType comm_backend = backend_type;
+      if (!comm->HasBackend(comm_backend)) {
+        comm_backend = comm->intra_comm_backend();
+        if (comm_backend == BackendType::kCount) {
+          comm_backend = comm->inter_comm_backend();
+        }
+      }
+      if (comm_backend == BackendType::kCount) {
+        LOG("No initialized backend is available for `Broadcast`.");
+        return ReturnStatus::kInternalError;
+      }
+
+      // Keep the key dependent so provider `BackendEnabled`
+      // specializations are visible when redispatch is instantiated.
+      using DispatchKey =
+          typename BackendDependentType<backend_type, Broadcast>::type;
+      return Operation<DispatchKey>::Call(comm_backend, comm->device_type(),
+                                          send_buff, recv_buff, count, datatype,
+                                          root, comm_handle, stream);
+    }
+
     return BroadcastImpl<backend_type, device_type>::Apply(
         send_buff, recv_buff, count, datatype, root, comm, stream);
   }
 
  private:
+  template <BackendType, typename T>
+  struct BackendDependentType {
+    using type = T;
+  };
+
   static bool HasInvalidArgs(const void *send_buff, void *recv_buff,
                              size_t count, DataType datatype, int root,
                              Communicator *comm) {