diff --git a/CMakeLists.txt b/CMakeLists.txt
index 4ffbc25eb..5f193ba24 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -107,7 +107,7 @@ if(USE_CUDA)
   file(GLOB_RECURSE CUDA_KERNELS ${PROJECT_SOURCE_DIR}/infini_train/src/*.cu)
 
   add_library(infini_train_cuda_kernels STATIC ${CUDA_KERNELS})
-  set_target_properties(infini_train_cuda_kernels PROPERTIES CUDA_ARCHITECTURES "75;80;90")
+  set_target_properties(infini_train_cuda_kernels PROPERTIES CUDA_ARCHITECTURES "75;80;90;120")
 
   target_link_libraries(infini_train_cuda_kernels
     PUBLIC
diff --git a/example/gpt2/main.cc b/example/gpt2/main.cc
index ff79a3775..d028102bc 100644
--- a/example/gpt2/main.cc
+++ b/example/gpt2/main.cc
@@ -16,6 +16,7 @@
 #include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/dataloader.h"
 #include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
 #include "infini_train/include/lr_scheduler.h"
 #include "infini_train/include/nn/lora/lora_utils.h"
 #include "infini_train/include/nn/modules/loss.h"
@@ -216,9 +217,6 @@ void Train(const nn::parallel::Rank &rank) {
     LOG(INFO) << "total desired batch size: " << FLAGS_total_batch_size
               << " => calculated gradient accumulation steps: " << grad_accum_steps;
 
-    // rng / reproducibility
-    // ManualSeed(42);
-
     // init the model, either from scratch or from OpenAI pretrained checkpoint
     nn::TransformerConfig model_config = gpt2::GPT2Config();
     std::shared_ptr<nn::Module> model = nullptr;
@@ -581,6 +579,7 @@ int main(int argc, char *argv[]) {
     nn::parallel::global::InitAllEnv(FLAGS_nthread_per_process, FLAGS_tensor_parallel, FLAGS_sequence_parallel,
                                      FLAGS_pipeline_parallel, FLAGS_virtual_pipeline_parallel);
     utils::PrecisionCheckEnv::Instance().Init(precision_config);
+    infini_train::ManualSeed(42);
 
     LOG(INFO) << nn::parallel::global::ProcessGroupOverview();
 
diff --git a/example/llama3/main.cc b/example/llama3/main.cc
index 344247e2c..b52b7536e 100644
--- a/example/llama3/main.cc
+++ b/example/llama3/main.cc
@@ -15,6 +15,7 @@
 #include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/dataloader.h"
 #include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
 #include "infini_train/include/lr_scheduler.h"
 #include "infini_train/include/nn/lora/lora_utils.h"
 #include "infini_train/include/nn/modules/loss.h"
@@ -206,9 +207,6 @@ void Train(const nn::parallel::Rank &rank) {
                   << " => calculated gradient accumulation steps: " << grad_accum_steps;
     }
 
-    // rng / reproducibility
-    // ManualSeed(42);
-
     nn::TransformerConfig model_config = llama3::LLaMA3Config();
     std::shared_ptr<nn::Module> model = nullptr;
     if (!FLAGS_llmc_filepath.empty()) {
@@ -560,6 +558,7 @@ int main(int argc, char *argv[]) {
     nn::parallel::global::InitAllEnv(FLAGS_nthread_per_process, FLAGS_tensor_parallel, FLAGS_sequence_parallel,
                                      FLAGS_pipeline_parallel, FLAGS_virtual_pipeline_parallel);
     utils::PrecisionCheckEnv::Instance().Init(precision_config);
+    infini_train::ManualSeed(42);
 
     LOG(INFO) << nn::parallel::global::ProcessGroupOverview();
 
diff --git a/example/mixtral/main.cc b/example/mixtral/main.cc
index 553d7a626..2f5fd1d8f 100644
--- a/example/mixtral/main.cc
+++ b/example/mixtral/main.cc
@@ -13,6 +13,7 @@
 #include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/dataloader.h"
 #include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
 #include "infini_train/include/nn/modules/loss.h"
 #include "infini_train/include/nn/modules/transformer/transformer.h"
 #include "infini_train/include/nn/parallel/global.h"
@@ -69,6 +70,7 @@ int main(int argc, char *argv[]) {
         /*sequence_parallel_enabled=*/false,
         /*pipeline_parallel_size=*/1,
         /*virtual_pipeline_parallel_size=*/1);
+    infini_train::ManualSeed(42);
 
     infini_train::nn::TransformerConfig model_config = mixtral::TinyMixtralConfig();
     mixtral::SanitizeTinyMixtralConfig(model_config);
diff --git a/example/mnist/main.cc b/example/mnist/main.cc
index 7744e0947..46095c300 100644
--- a/example/mnist/main.cc
+++ b/example/mnist/main.cc
@@ -11,6 +11,7 @@
 
 #include "infini_train/include/dataloader.h"
 #include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
 #include "infini_train/include/nn/modules/loss.h"
 #include "infini_train/include/optimizer.h"
 
@@ -39,6 +40,7 @@ DEFINE_validator(device,
 int main(int argc, char *argv[]) {
     gflags::ParseCommandLineFlags(&argc, &argv, true);
     google::InitGoogleLogging(argv[0]);
+    infini_train::ManualSeed(42);
 
     auto train_dataset = std::make_shared<MNISTDataset>(FLAGS_dataset, true);
     DataLoader train_dataloader(train_dataset, FLAGS_bs);
diff --git a/infini_train/include/autograd/dropout.h b/infini_train/include/autograd/dropout.h
new file mode 100644
index 000000000..a6573e1d5
--- /dev/null
+++ b/infini_train/include/autograd/dropout.h
@@ -0,0 +1,34 @@
+#pragma once
+
+#include <memory>
+#include <optional>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "infini_train/include/autograd/function.h"
+#include "infini_train/include/generator.h"
+
+namespace infini_train {
+class Tensor;
+}
+
+namespace infini_train::autograd {
+
+class Dropout final : public Function {
+public:
+    static constexpr char kType[] = "DropoutFunction";
+
+    Dropout(double p, std::optional<Generator> generator) : Function(kType), p_(p), generator_(std::move(generator)) {}
+
+    std::vector<std::shared_ptr<Tensor>> Forward(const std::vector<std::shared_ptr<Tensor>> &input_tensors) override;
+    void SetupContext(const std::vector<std::shared_ptr<Tensor>> &input_tensors,
+                      const std::vector<std::shared_ptr<Tensor>> &output_tensors) override;
+    std::vector<std::shared_ptr<Tensor>> Backward(const std::vector<std::shared_ptr<Tensor>> &grad_outputs) override;
+
+private:
+    double p_ = 0.0;
+    std::optional<Generator> generator_;
+};
+
+} // namespace infini_train::autograd
diff --git a/infini_train/include/core/runtime/generator_backend.h b/infini_train/include/core/runtime/generator_backend.h
new file mode 100644
index 000000000..71798b1ce
--- /dev/null
+++ b/infini_train/include/core/runtime/generator_backend.h
@@ -0,0 +1,50 @@
+#pragma once
+
+// Internal Generator backend registration interface. Application code should include generator.h.
+
+#include <cstdint>
+#include <memory>
+#include <unordered_map>
+
+#include "infini_train/include/common/common.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
+
+namespace infini_train {
+
+class GeneratorBackend {
+public:
+    virtual ~GeneratorBackend() = default;
+
+    virtual Device::DeviceType Type() const = 0;
+    virtual Generator Create(const Device &device, uint64_t seed) = 0;
+    virtual const Generator &GetDefault(const Device &device) = 0;
+    virtual void ManualSeedAll(uint64_t seed) = 0;
+};
+
+// Registers one backend per device type during static initialization.
+class GeneratorBackendRegistry {
+public:
+    static GeneratorBackendRegistry &Instance();
+
+    void Register(Device::DeviceType type, std::unique_ptr<GeneratorBackend> backend);
+    // Throws if no backend is registered for the requested device type.
+    GeneratorBackend &Get(Device::DeviceType type) const;
+    void ManualSeedAll(uint64_t seed) const;
+
+private:
+    GeneratorBackendRegistry() = default;
+    GeneratorBackendRegistry(const GeneratorBackendRegistry &) = delete;
+    GeneratorBackendRegistry &operator=(const GeneratorBackendRegistry &) = delete;
+
+    std::unordered_map<Device::DeviceType, std::unique_ptr<GeneratorBackend>> backends_;
+};
+
+} // namespace infini_train
+
+// CAT expands __COUNTER__ before token pasting so each registration has a unique name.
+#define INFINI_TRAIN_REGISTER_GENERATOR_BACKEND(device_type, class_impl)                                               \
+    [[maybe_unused]] static const bool CAT(infini_train_generator_backend_registered_, __COUNTER__) = []() {           \
+        infini_train::GeneratorBackendRegistry::Instance().Register(device_type, std::make_unique<class_impl>());      \
+        return true;                                                                                                   \
+    }();
diff --git a/infini_train/include/generator.h b/infini_train/include/generator.h
new file mode 100644
index 000000000..00c508554
--- /dev/null
+++ b/infini_train/include/generator.h
@@ -0,0 +1,53 @@
+#pragma once
+
+#include <cstdint>
+#include <memory>
+
+#include "infini_train/include/device.h"
+
+namespace infini_train {
+
+class Tensor;
+class GeneratorImpl;
+
+// A lightweight handle with shared-copy semantics. Use Clone() for an independent state.
+class Generator {
+public:
+    static constexpr uint64_t kDefaultSeed = 67280421310721;
+
+    Generator() = delete;
+
+    Generator(const Generator &) = default;
+    Generator &operator=(const Generator &) = default;
+    Generator(Generator &&) = default;
+    Generator &operator=(Generator &&) = default;
+
+    ~Generator() = default;
+
+    void set_current_seed(uint64_t seed) const;
+    uint64_t current_seed() const;
+    uint64_t Seed();
+    void set_state(const Tensor &state);
+    std::shared_ptr<Tensor> get_state() const;
+    Device device() const;
+    Generator Clone() const;
+
+    friend bool operator==(const Generator &a, const Generator &b) { return a.impl_ == b.impl_; }
+    friend bool operator!=(const Generator &a, const Generator &b) { return !(a == b); }
+
+private:
+    friend class GeneratorAccessor;
+
+    explicit Generator(std::shared_ptr<GeneratorImpl> impl);
+    std::shared_ptr<GeneratorImpl> impl_;
+};
+
+Generator CreateGenerator(const Device &device, uint64_t seed = Generator::kDefaultSeed);
+
+// Returns the lazily initialized default generator for the requested device.
+const Generator &GetDefaultGenerator(const Device &device);
+
+// Initializes or resets the default generators for all enabled devices with the given seed.
+void ManualSeed(uint64_t seed);
+
+} // namespace infini_train
diff --git a/infini_train/include/generator_impl.h b/infini_train/include/generator_impl.h
new file mode 100644
index 000000000..9a59ee7f6
--- /dev/null
+++ b/infini_train/include/generator_impl.h
@@ -0,0 +1,90 @@
+#pragma once
+
+// Internal Generator implementation interface. Application code should include generator.h.
+
+#include <cstdint>
+#include <memory>
+#include <mutex>
+#include <optional>
+#include <stdexcept>
+#include <type_traits>
+#include <utility>
+
+#include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train {
+
+class Tensor;
+
+class GeneratorImpl {
+public:
+    explicit GeneratorImpl(Device device) : device_(device) {}
+    virtual ~GeneratorImpl() = default;
+
+    GeneratorImpl(const GeneratorImpl &other) = delete;
+    GeneratorImpl(GeneratorImpl &&other) = delete;
+    GeneratorImpl &operator=(const GeneratorImpl &other) = delete;
+    GeneratorImpl &operator=(GeneratorImpl &&other) = delete;
+
+    virtual void set_current_seed(uint64_t seed) = 0;
+    virtual uint64_t current_seed() const = 0;
+    virtual uint64_t Seed() = 0;
+    virtual void set_state(const Tensor &state) = 0;
+    virtual std::shared_ptr<Tensor> get_state() const = 0;
+
+    std::shared_ptr<GeneratorImpl> Clone() const { return std::shared_ptr<GeneratorImpl>(CloneImpl()); }
+    Device device() const { return device_; }
+
+    // Callers must hold this mutex when accessing shared RNG state concurrently.
+    std::mutex mutex_;
+
+protected:
+    Device device_;
+    virtual GeneratorImpl *CloneImpl() const = 0;
+};
+
+class GeneratorAccessor {
+public:
+    static GeneratorImpl &Get(const Generator &generator) {
+        if (!generator.impl_) {
+            throw std::invalid_argument("Undefined Generator");
+        }
+        return *generator.impl_;
+    }
+
+    static std::mutex &Mutex(const Generator &generator) { return Get(generator).mutex_; }
+
+    static Generator FromImpl(std::shared_ptr<GeneratorImpl> impl) { return Generator(std::move(impl)); }
+};
+
+template <typename T> T &CheckedGeneratorImpl(const Generator &generator, const Device &expected_device) {
+    static_assert(std::is_base_of_v<GeneratorImpl, T>);
+    auto &base = GeneratorAccessor::Get(generator);
+    if (base.device() != expected_device) {
+        throw std::invalid_argument("Generator device mismatch");
+    }
+    auto *typed = dynamic_cast<T *>(&base);
+    if (typed == nullptr) {
+        throw std::invalid_argument("Generator backend mismatch");
+    }
+    return *typed;
+}
+
+template <class Impl, class... Args> Generator MakeGenerator(Args &&...args) {
+    return GeneratorAccessor::FromImpl(std::make_shared<Impl>(std::forward<Args>(args)...));
+}
+
+template <typename T>
+T &GetGeneratorOrDefault(const std::optional<Generator> &generator, const Generator &default_generator,
+                         const Device &expected_device) {
+    const Generator &chosen = generator.has_value() ? *generator : default_generator;
+    return CheckedGeneratorImpl<T>(chosen, expected_device);
+}
+
+namespace detail {
+void CheckRngState(const Tensor &state);
+} // namespace detail
+
+} // namespace infini_train
diff --git a/infini_train/include/nn/functional.h b/infini_train/include/nn/functional.h
index e4354fd10..7ed6d047b 100644
--- a/infini_train/include/nn/functional.h
+++ b/infini_train/include/nn/functional.h
@@ -2,8 +2,13 @@
 
 #include <cstdint>
 #include <memory>
+#include <optional>
 #include <vector>
 
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
+
 namespace infini_train {
 class Tensor;
 }
@@ -47,6 +52,19 @@ std::shared_ptr<Tensor> Triu(const std::shared_ptr<Tensor> &input, int64_t diago
 //   A tensor of the given shape filled with the scalar value 1.
 std::shared_ptr<Tensor> Ones(const std::vector<int64_t> size);
 
+// Returns a tensor with uniformly distributed random values in [0, 1).
+std::shared_ptr<Tensor> Rand(const std::vector<int64_t> &size, DataType dtype = DataType::kFLOAT32,
+                             Device device = Device(), std::optional<Generator> generator = std::nullopt,
+                             bool requires_grad = false);
+
+// Returns a tensor with normally distributed random values with mean 0 and standard deviation 1.
+std::shared_ptr<Tensor> Randn(const std::vector<int64_t> &size, DataType dtype = DataType::kFLOAT32,
+                              Device device = Device(), std::optional<Generator> generator = std::nullopt,
+                              bool requires_grad = false);
+
+std::shared_ptr<Tensor> Dropout(const std::shared_ptr<Tensor> &input, double p = 0.5, bool training = true,
+                                std::optional<Generator> generator = std::nullopt);
+
 // Returns a new tensor with the reciprocal of the elements of input.
 //
 // Args:
diff --git a/infini_train/include/nn/init.h b/infini_train/include/nn/init.h
index fc6effec9..bbc76c8ae 100644
--- a/infini_train/include/nn/init.h
+++ b/infini_train/include/nn/init.h
@@ -1,21 +1,21 @@
 #pragma once
 
+#include <cstdint>
 #include <memory>
 #include <optional>
-#include <random>
 #include <utility>
 
 #include "infini_train/include/datatype.h"
 #include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
 
 namespace infini_train {
 class Tensor;
-class Device;
 } // namespace infini_train
 
 namespace infini_train::nn::init {
 std::shared_ptr<Tensor> Normal(const std::shared_ptr<Tensor> &tensor, float mean = 0.0, float std = 1.0,
-                               std::optional<std::mt19937> generator = std::nullopt);
+                               std::optional<Generator> generator = std::nullopt);
 
 std::pair<int64_t, int64_t> CalculateFanInAndFanOut(const std::shared_ptr<Tensor> &tensor);
 
@@ -42,10 +42,10 @@ enum class NonLinearityType : int8_t {
 std::shared_ptr<Tensor> KaimingUniform(const std::shared_ptr<Tensor> &tensor, float a = 0.0f,
                                        KaimingMode mode = KaimingMode::kFanIn,
                                        NonLinearityType non_linearity = NonLinearityType::kLeakyReLU,
-                                       std::optional<std::mt19937> generator = std::nullopt);
+                                       std::optional<Generator> generator = std::nullopt);
 
 std::shared_ptr<Tensor> Uniform(const std::shared_ptr<Tensor> &tensor, float a = 0.0f, float b = 1.0f,
-                                std::optional<std::mt19937> generator = std::nullopt);
+                                std::optional<Generator> generator = std::nullopt);
 
 std::shared_ptr<Tensor> Ones(const std::shared_ptr<Tensor> &tensor);
 
diff --git a/infini_train/include/nn/parallel/ddp/distributed_data_parallel.h b/infini_train/include/nn/parallel/ddp/distributed_data_parallel.h
index 823ae82b5..1b2d3ffa3 100644
--- a/infini_train/include/nn/parallel/ddp/distributed_data_parallel.h
+++ b/infini_train/include/nn/parallel/ddp/distributed_data_parallel.h
@@ -43,6 +43,8 @@ class DistributedDataParallel : public nn::Module {
     void OnGradReady(const std::shared_ptr<Tensor> &param);
 
 private:
+    void SynchronizeModuleState();
+
     std::shared_ptr<Reducer> reducer_ = nullptr;
 
     DistributedDataParallelConfig ddp_config_;
diff --git a/infini_train/include/tensor.h b/infini_train/include/tensor.h
index dcfd8927f..e5ef60ec3 100644
--- a/infini_train/include/tensor.h
+++ b/infini_train/include/tensor.h
@@ -4,7 +4,6 @@
 #include <iostream>
 #include <memory>
 #include <optional>
-#include <random>
 #include <vector>
 
 #include "Eigen/Dense"
@@ -12,6 +11,7 @@
 
 #include "infini_train/include/datatype.h"
 #include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
 #include "infini_train/include/scalar.h"
 
 namespace infini_train {
@@ -74,6 +74,8 @@ class Tensor : public std::enable_shared_from_this<Tensor> {
     void *DataPtr();
     const void *DataPtr() const;
 
+    bool defined() const { return buffer_ != nullptr; }
+
     size_t SizeInBytes() const;
 
     const std::vector<int64_t> &Dims() const;
@@ -151,7 +153,7 @@ class Tensor : public std::enable_shared_from_this<Tensor> {
 
     // distribution
     std::shared_ptr<Tensor> Uniform(float from = 0.0f, float to = 1.0f,
-                                    std::optional<std::mt19937> generator = std::nullopt);
+                                    std::optional<Generator> generator = std::nullopt);
 
     std::shared_ptr<Tensor> Matmul(const std::shared_ptr<Tensor> &other);
     std::shared_ptr<Tensor> Outer(const std::shared_ptr<Tensor> &other);
diff --git a/infini_train/src/autograd/dropout.cc b/infini_train/src/autograd/dropout.cc
new file mode 100644
index 000000000..531710e9f
--- /dev/null
+++ b/infini_train/src/autograd/dropout.cc
@@ -0,0 +1,45 @@
+#include "infini_train/include/autograd/dropout.h"
+
+#include <tuple>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/device.h"
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train::autograd {
+
+std::vector<std::shared_ptr<Tensor>> Dropout::Forward(const std::vector<std::shared_ptr<Tensor>> &input_tensors) {
+    CHECK_EQ(input_tensors.size(), 1);
+    const auto &input = input_tensors[0];
+
+    auto device = input->GetDevice().type();
+    auto outputs = Dispatcher::Instance().Call<std::tuple<std::shared_ptr<Tensor>, std::shared_ptr<Tensor>>>(
+        {device, "DropoutForward"}, input, p_, generator_);
+    // Keep the mask as a non-differentiable output so SetupContext can save it for backward.
+    return {std::get<0>(outputs), std::get<1>(outputs)};
+}
+
+void Dropout::SetupContext(const std::vector<std::shared_ptr<Tensor>> &,
+                           const std::vector<std::shared_ptr<Tensor>> &output_tensors) {
+    CHECK_EQ(output_tensors.size(), 2);
+    ctx_.MarkNonDifferentiable({output_tensors[1]});
+    if (!ctx_.needs_input_grad().empty() && ctx_.needs_input_grad()[0]) {
+        ctx_.SaveForBackward({output_tensors[1]});
+    }
+}
+
+std::vector<std::shared_ptr<Tensor>> Dropout::Backward(const std::vector<std::shared_ptr<Tensor>> &grad_outputs) {
+    // The mask output is non-differentiable, so its gradient slot is empty.
+    CHECK_EQ(grad_outputs.size(), 2);
+    auto saved_tensors = ctx_.GetSavedTensors();
+    CHECK_EQ(saved_tensors.size(), 1);
+    const auto &grad_output = grad_outputs[0];
+    const auto &mask = saved_tensors[0];
+
+    auto device = grad_output->GetDevice().type();
+    return {Dispatcher::Instance().Call<std::shared_ptr<Tensor>>({device, "DropoutBackward"}, grad_output, mask, p_)};
+}
+
+} // namespace infini_train::autograd
diff --git a/infini_train/src/core/runtime/cpu/cpu_generator_impl.cc b/infini_train/src/core/runtime/cpu/cpu_generator_impl.cc
new file mode 100644
index 000000000..cb827533a
--- /dev/null
+++ b/infini_train/src/core/runtime/cpu/cpu_generator_impl.cc
@@ -0,0 +1,263 @@
+#include "infini_train/src/core/runtime/cpu/cpu_generator_impl.h"
+
+#include <cstddef>
+#include <cstring>
+#include <mutex>
+#include <optional>
+#include <random>
+#include <sstream>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/core/runtime/generator_backend.h"
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/generator_impl.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train::core::cpu {
+namespace {
+
+// Backend tag used to reject states from other generator implementations.
+constexpr char kCPUStateMagic[] = "ITRNGCPU";
+constexpr std::size_t kStateMagicSize = sizeof(kCPUStateMagic) - 1;
+
+// Fixed footer after the variable-length mt19937 stream. Cache flags are
+// serialized as uint8_t; absent cache values are serialized as zero.
+constexpr std::size_t kSerializedSeedSize = sizeof(uint64_t);
+constexpr std::size_t kSerializedFloatCacheFlagSize = sizeof(uint8_t);
+constexpr std::size_t kSerializedFloatCacheValueSize = sizeof(float);
+constexpr std::size_t kSerializedDoubleCacheFlagSize = sizeof(uint8_t);
+constexpr std::size_t kSerializedDoubleCacheValueSize = sizeof(double);
+constexpr std::size_t kStateFooterSize = kSerializedSeedSize + kSerializedFloatCacheFlagSize
+                                       + kSerializedFloatCacheValueSize + kSerializedDoubleCacheFlagSize
+                                       + kSerializedDoubleCacheValueSize;
+
+// Seed through seed_seq using both 32-bit halves of the seed. This produces a
+// different sequence from direct std::mt19937(seed) initialization, even when
+// the seed fits in 32 bits.
+std::mt19937 MakeSeededEngine(uint64_t seed) {
+    std::seed_seq seq{static_cast<uint32_t>(seed), static_cast<uint32_t>(seed >> 32)};
+    std::mt19937 engine;
+    engine.seed(seq);
+    return engine;
+}
+
+uint64_t GenerateNonDeterministicSeed() {
+    std::random_device rd;
+    uint64_t val = (static_cast<uint64_t>(rd()) << 32) | rd();
+    return val;
+}
+
+// Copies state bytes and advances offset; data and value must reference valid
+// buffers of data_size and value_size bytes, and offset must be a valid pointer.
+void WriteBytes(uint8_t *data, std::size_t data_size, std::size_t *offset, const void *value, std::size_t value_size) {
+    CHECK_LE(*offset, data_size) << "CPU generator state write offset out of bounds";
+    CHECK_LE(value_size, data_size - *offset) << "CPU generator state write exceeds buffer";
+    if (value_size != 0) {
+        std::memcpy(data + *offset, value, value_size);
+    }
+    *offset += value_size;
+}
+
+void ReadBytes(const uint8_t *data, std::size_t data_size, std::size_t *offset, void *value, std::size_t value_size) {
+    CHECK_LE(*offset, data_size) << "CPU generator state read offset out of bounds";
+    CHECK_LE(value_size, data_size - *offset) << "CPU generator state read exceeds buffer";
+    if (value_size != 0) {
+        std::memcpy(value, data + *offset, value_size);
+    }
+    *offset += value_size;
+}
+
+} // namespace
+
+CPUGeneratorImpl::CPUGeneratorImpl(uint64_t seed)
+    : GeneratorImpl(Device(Device::DeviceType::kCPU, 0)), engine_(MakeSeededEngine(seed)), seed_(seed) {}
+
+void CPUGeneratorImpl::set_current_seed(uint64_t seed) {
+    seed_ = seed;
+    next_float_normal_sample_.reset();
+    next_double_normal_sample_.reset();
+    engine_ = MakeSeededEngine(seed);
+}
+
+uint64_t CPUGeneratorImpl::current_seed() const { return seed_; }
+
+uint64_t CPUGeneratorImpl::Seed() {
+    uint64_t random_seed = GenerateNonDeterministicSeed();
+    set_current_seed(random_seed);
+    return random_seed;
+}
+
+// State layout: magic, serialized engine, seed, and cached float/double normal samples.
+// The std::mt19937 stream format is only portable across compatible standard-library builds.
+
+void CPUGeneratorImpl::set_state(const Tensor &state) {
+    ::infini_train::detail::CheckRngState(state);
+
+    const size_t data_size = state.SizeInBytes();
+    CHECK_GT(data_size, kStateMagicSize + kStateFooterSize) << "CPU generator state is too small";
+
+    const uint8_t *data = static_cast<const uint8_t *>(state.DataPtr());
+    CHECK_EQ(std::memcmp(data, kCPUStateMagic, kStateMagicSize), 0)
+        << "Invalid RNG state: not a CPU generator state (backend magic mismatch)";
+
+    const size_t engine_size = data_size - kStateMagicSize - kStateFooterSize;
+    std::string engine_str(reinterpret_cast<const char *>(data + kStateMagicSize), engine_size);
+
+    std::istringstream iss(engine_str);
+    std::mt19937 restored_engine;
+    iss >> restored_engine;
+    CHECK(!iss.fail()) << "Invalid CPU generator engine state";
+    iss >> std::ws;
+    CHECK(iss.eof()) << "Invalid trailing bytes in CPU generator engine state";
+
+    size_t offset = kStateMagicSize + engine_size;
+    uint64_t restored_seed = 0;
+    ReadBytes(data, data_size, &offset, &restored_seed, kSerializedSeedSize);
+    uint8_t has_float = 0;
+    ReadBytes(data, data_size, &offset, &has_float, kSerializedFloatCacheFlagSize);
+    CHECK_LE(has_float, 1) << "Invalid CPU generator float normal cache flag";
+    float restored_float = 0.0f;
+    ReadBytes(data, data_size, &offset, &restored_float, kSerializedFloatCacheValueSize);
+    uint8_t has_double = 0;
+    ReadBytes(data, data_size, &offset, &has_double, kSerializedDoubleCacheFlagSize);
+    CHECK_LE(has_double, 1) << "Invalid CPU generator double normal cache flag";
+    double restored_double = 0.0;
+    ReadBytes(data, data_size, &offset, &restored_double, kSerializedDoubleCacheValueSize);
+    CHECK_EQ(offset, data_size) << "CPU generator state size mismatch";
+
+    // Do not change the generator until the complete state has been validated.
+    engine_ = restored_engine;
+    seed_ = restored_seed;
+    next_float_normal_sample_ = has_float ? std::optional<float>(restored_float) : std::nullopt;
+    next_double_normal_sample_ = has_double ? std::optional<double>(restored_double) : std::nullopt;
+}
+
+std::shared_ptr<Tensor> CPUGeneratorImpl::get_state() const {
+    std::ostringstream oss;
+    oss << engine_;
+    std::string engine_str = oss.str();
+
+    const size_t engine_size = engine_str.size();
+    const size_t total_size = kStateMagicSize + engine_size + kStateFooterSize;
+
+    auto state_tensor = std::make_shared<Tensor>(std::vector<int64_t>{static_cast<int64_t>(total_size)},
+                                                 DataType::kUINT8, Device(Device::DeviceType::kCPU, 0));
+
+    uint8_t *data = static_cast<uint8_t *>(state_tensor->DataPtr());
+    size_t offset = 0;
+
+    const uint8_t has_float = next_float_normal_sample_.has_value() ? 1 : 0;
+    const float float_val = has_float ? *next_float_normal_sample_ : 0.0f;
+    const uint8_t has_double = next_double_normal_sample_.has_value() ? 1 : 0;
+    const double double_val = has_double ? *next_double_normal_sample_ : 0.0;
+
+    WriteBytes(data, total_size, &offset, kCPUStateMagic, kStateMagicSize);
+    WriteBytes(data, total_size, &offset, engine_str.data(), engine_size);
+    WriteBytes(data, total_size, &offset, &seed_, kSerializedSeedSize);
+    WriteBytes(data, total_size, &offset, &has_float, kSerializedFloatCacheFlagSize);
+    WriteBytes(data, total_size, &offset, &float_val, kSerializedFloatCacheValueSize);
+    WriteBytes(data, total_size, &offset, &has_double, kSerializedDoubleCacheFlagSize);
+    WriteBytes(data, total_size, &offset, &double_val, kSerializedDoubleCacheValueSize);
+    CHECK_EQ(offset, total_size) << "CPU generator state size mismatch";
+
+    return state_tensor;
+}
+
+uint32_t CPUGeneratorImpl::Random() { return engine_(); }
+
+uint64_t CPUGeneratorImpl::Random64() {
+    uint32_t hi = engine_();
+    uint32_t lo = engine_();
+    return (static_cast<uint64_t>(hi) << 32) | lo;
+}
+
+std::optional<float> CPUGeneratorImpl::next_float_normal_sample() const { return next_float_normal_sample_; }
+
+std::optional<double> CPUGeneratorImpl::next_double_normal_sample() const { return next_double_normal_sample_; }
+
+void CPUGeneratorImpl::set_next_float_normal_sample(std::optional<float> randn) { next_float_normal_sample_ = randn; }
+
+void CPUGeneratorImpl::set_next_double_normal_sample(std::optional<double> randn) {
+    next_double_normal_sample_ = randn;
+}
+
+std::shared_ptr<CPUGeneratorImpl> CPUGeneratorImpl::Clone() const {
+    return std::shared_ptr<CPUGeneratorImpl>(CloneImpl());
+}
+
+CPUGeneratorImpl *CPUGeneratorImpl::CloneImpl() const {
+    auto *clone = new CPUGeneratorImpl(seed_);
+    clone->set_engine(engine_);
+    clone->set_next_float_normal_sample(next_float_normal_sample_);
+    clone->set_next_double_normal_sample(next_double_normal_sample_);
+    return clone;
+}
+
+void CPUGeneratorImpl::set_engine(std::mt19937 engine) { engine_ = std::move(engine); }
+
+} // namespace infini_train::core::cpu
+
+namespace infini_train::core::cpu {
+namespace {
+
+struct DefaultCpuGeneratorState {
+    std::once_flag init_once;
+    std::optional<Generator> generator;
+};
+
+DefaultCpuGeneratorState &GetDefaultCpuGeneratorState() {
+    static DefaultCpuGeneratorState state;
+    return state;
+}
+
+} // namespace
+
+const Generator &GetDefaultCpuGenerator() {
+    auto &state = GetDefaultCpuGeneratorState();
+    std::call_once(state.init_once,
+                   [&state] { state.generator.emplace(CreateCpuGenerator(GenerateNonDeterministicSeed())); });
+    return *state.generator;
+}
+
+Generator CreateCpuGenerator(uint64_t seed) { return MakeGenerator<CPUGeneratorImpl>(seed); }
+
+void ManualSeed(uint64_t seed) {
+    auto &state = GetDefaultCpuGeneratorState();
+    bool initialized_here = false;
+    std::call_once(state.init_once, [&state, seed, &initialized_here] {
+        state.generator.emplace(CreateCpuGenerator(seed));
+        initialized_here = true;
+    });
+    if (initialized_here) {
+        return;
+    }
+    const Generator &generator = *state.generator;
+    std::lock_guard<std::mutex> lock(GeneratorAccessor::Mutex(generator));
+    generator.set_current_seed(seed);
+}
+
+} // namespace infini_train::core::cpu
+
+namespace infini_train::core::cpu {
+namespace {
+
+class CPUGeneratorBackend final : public GeneratorBackend {
+public:
+    Device::DeviceType Type() const override { return Device::DeviceType::kCPU; }
+
+    Generator Create(const Device & /*device*/, uint64_t seed) override { return CreateCpuGenerator(seed); }
+
+    const Generator &GetDefault(const Device & /*device*/) override { return GetDefaultCpuGenerator(); }
+
+    void ManualSeedAll(uint64_t seed) override { ::infini_train::core::cpu::ManualSeed(seed); }
+};
+
+INFINI_TRAIN_REGISTER_GENERATOR_BACKEND(Device::DeviceType::kCPU, CPUGeneratorBackend);
+
+} // namespace
+} // namespace infini_train::core::cpu
diff --git a/infini_train/src/core/runtime/cpu/cpu_generator_impl.h b/infini_train/src/core/runtime/cpu/cpu_generator_impl.h
new file mode 100644
index 000000000..738b20c74
--- /dev/null
+++ b/infini_train/src/core/runtime/cpu/cpu_generator_impl.h
@@ -0,0 +1,52 @@
+#pragma once
+
+#include <cstdint>
+#include <memory>
+#include <optional>
+#include <random>
+
+#include "infini_train/include/generator.h"
+#include "infini_train/include/generator_impl.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train::core::cpu {
+
+// CPU generator backed by std::mt19937 with cached Box-Muller samples.
+class CPUGeneratorImpl final : public GeneratorImpl {
+public:
+    explicit CPUGeneratorImpl(uint64_t seed = Generator::kDefaultSeed);
+    ~CPUGeneratorImpl() override = default;
+
+    void set_current_seed(uint64_t seed) override;
+    uint64_t current_seed() const override;
+    uint64_t Seed() override;
+    void set_state(const Tensor &state) override;
+    std::shared_ptr<Tensor> get_state() const override;
+
+    std::shared_ptr<CPUGeneratorImpl> Clone() const;
+
+    uint32_t Random();
+    uint64_t Random64();
+
+    std::optional<float> next_float_normal_sample() const;
+    std::optional<double> next_double_normal_sample() const;
+    void set_next_float_normal_sample(std::optional<float> randn);
+    void set_next_double_normal_sample(std::optional<double> randn);
+
+private:
+    CPUGeneratorImpl *CloneImpl() const override;
+
+    std::mt19937 engine() const { return engine_; }
+    void set_engine(std::mt19937 engine);
+
+    std::mt19937 engine_;
+    uint64_t seed_ = Generator::kDefaultSeed;
+    std::optional<float> next_float_normal_sample_;
+    std::optional<double> next_double_normal_sample_;
+};
+
+const Generator &GetDefaultCpuGenerator();
+Generator CreateCpuGenerator(uint64_t seed);
+void ManualSeed(uint64_t seed);
+
+} // namespace infini_train::core::cpu
diff --git a/infini_train/src/core/runtime/cuda/cuda_generator_impl.cc b/infini_train/src/core/runtime/cuda/cuda_generator_impl.cc
new file mode 100644
index 000000000..9abd97a1f
--- /dev/null
+++ b/infini_train/src/core/runtime/cuda/cuda_generator_impl.cc
@@ -0,0 +1,173 @@
+#include "infini_train/src/core/runtime/cuda/cuda_generator_impl.h"
+
+#include <cstring>
+#include <deque>
+#include <limits>
+#include <memory>
+#include <mutex>
+#include <optional>
+#include <ostream>
+#include <random>
+#include <vector>
+
+#include <cuda_runtime.h>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/common/cuda/common_cuda.h"
+#include "infini_train/include/core/runtime/generator_backend.h"
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/generator_impl.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train::core::cuda {
+namespace {
+
+// Backend tag used to reject states from other generator implementations.
+constexpr char kCUDAStateMagic[] = "ITRNGCUD";
+constexpr size_t kStateMagicSize = sizeof(kCUDAStateMagic) - 1;
+
+// [magic: 8B] + [seed_: 8B] + [next_philox_subsequence_: 8B]
+constexpr size_t kStateSize = kStateMagicSize + sizeof(uint64_t) * 2;
+
+std::once_flag default_generators_init_flag;
+std::vector<std::optional<Generator>> default_generators;
+std::deque<std::once_flag> default_generator_init_flags;
+
+uint64_t GenerateNonDeterministicSeed() {
+    std::random_device random_device;
+    return (static_cast<uint64_t>(random_device()) << 32) | random_device();
+}
+
+void InitDefaultGenerators() {
+    std::call_once(default_generators_init_flag, [] {
+        int device_count = 0;
+        const cudaError_t status = cudaGetDeviceCount(&device_count);
+        if (status == cudaErrorNoDevice) {
+            cudaGetLastError();
+            return;
+        }
+        CHECK_EQ(status, cudaSuccess) << "cudaGetDeviceCount failed: " << cudaGetErrorString(status);
+        default_generators.resize(device_count);
+        default_generator_init_flags.resize(device_count);
+    });
+}
+
+int ResolveDeviceIndex(int8_t device_index) {
+    InitDefaultGenerators();
+    int index = device_index;
+    if (index == -1) {
+        CUDA_CHECK(cudaGetDevice(&index));
+    }
+    int device_count = 0;
+    device_count = static_cast<int>(default_generators.size());
+    CHECK(index >= 0 && index < device_count) << "Invalid CUDA device index " << index;
+    return index;
+}
+
+} // namespace
+
+CUDAGeneratorImpl::CUDAGeneratorImpl(int8_t device_index, uint64_t seed)
+    : GeneratorImpl(Device(Device::DeviceType::kCUDA, device_index)), seed_(seed) {}
+
+void CUDAGeneratorImpl::set_current_seed(uint64_t seed) {
+    seed_ = seed;
+    next_philox_subsequence_ = 0;
+}
+
+uint64_t CUDAGeneratorImpl::current_seed() const { return seed_; }
+
+uint64_t CUDAGeneratorImpl::Seed() {
+    const uint64_t random_seed = GenerateNonDeterministicSeed();
+    set_current_seed(random_seed);
+    return random_seed;
+}
+
+void CUDAGeneratorImpl::set_state(const Tensor &state) {
+    ::infini_train::detail::CheckRngState(state);
+    CHECK_EQ(state.SizeInBytes(), kStateSize);
+
+    const auto *data = static_cast<const uint8_t *>(state.DataPtr());
+    CHECK_EQ(std::memcmp(data, kCUDAStateMagic, kStateMagicSize), 0)
+        << "Invalid RNG state: not a CUDA generator state (backend magic mismatch)";
+    std::memcpy(&seed_, data + kStateMagicSize, sizeof(seed_));
+    std::memcpy(&next_philox_subsequence_, data + kStateMagicSize + sizeof(seed_), sizeof(next_philox_subsequence_));
+}
+
+std::shared_ptr<Tensor> CUDAGeneratorImpl::get_state() const {
+    auto state = std::make_shared<Tensor>(std::vector<int64_t>{static_cast<int64_t>(kStateSize)}, DataType::kUINT8,
+                                          Device(Device::DeviceType::kCPU, 0));
+
+    auto *data = static_cast<uint8_t *>(state->DataPtr());
+    std::memcpy(data, kCUDAStateMagic, kStateMagicSize);
+    std::memcpy(data + kStateMagicSize, &seed_, sizeof(seed_));
+    std::memcpy(data + kStateMagicSize + sizeof(seed_), &next_philox_subsequence_, sizeof(next_philox_subsequence_));
+    return state;
+}
+
+uint64_t CUDAGeneratorImpl::ReservePhiloxSubsequence(uint64_t increment) {
+    CHECK_LE(increment, std::numeric_limits<uint64_t>::max() - next_philox_subsequence_)
+        << "Philox subsequence counter overflow";
+    const uint64_t subsequence = next_philox_subsequence_;
+    next_philox_subsequence_ += increment;
+    return subsequence;
+}
+
+CUDAGeneratorImpl *CUDAGeneratorImpl::CloneImpl() const {
+    auto *clone = new CUDAGeneratorImpl(device().index(), seed_);
+    clone->next_philox_subsequence_ = next_philox_subsequence_;
+    return clone;
+}
+
+const Generator &GetDefaultCudaGenerator(int8_t device_index) {
+    const int index = ResolveDeviceIndex(device_index);
+    std::call_once(default_generator_init_flags[index], [index] {
+        default_generators[index].emplace(
+            CreateCudaGenerator(static_cast<int8_t>(index), GenerateNonDeterministicSeed()));
+    });
+    return *default_generators[index];
+}
+
+Generator CreateCudaGenerator(int8_t device_index, uint64_t seed) {
+    const int index = ResolveDeviceIndex(device_index);
+    return MakeGenerator<CUDAGeneratorImpl>(static_cast<int8_t>(index), seed);
+}
+
+void ManualSeedAll(uint64_t seed) {
+    InitDefaultGenerators();
+    for (size_t index = 0; index < default_generators.size(); ++index) {
+        bool initialized_here = false;
+        std::call_once(default_generator_init_flags[index], [index, seed, &initialized_here] {
+            default_generators[index].emplace(CreateCudaGenerator(static_cast<int8_t>(index), seed));
+            initialized_here = true;
+        });
+        if (initialized_here) {
+            continue;
+        }
+        const Generator &generator = *default_generators[index];
+        std::lock_guard<std::mutex> lock(GeneratorAccessor::Mutex(generator));
+        generator.set_current_seed(seed);
+    }
+}
+
+} // namespace infini_train::core::cuda
+
+namespace infini_train::core::cuda {
+namespace {
+
+class CUDAGeneratorBackend final : public GeneratorBackend {
+public:
+    Device::DeviceType Type() const override { return Device::DeviceType::kCUDA; }
+
+    Generator Create(const Device &device, uint64_t seed) override { return CreateCudaGenerator(device.index(), seed); }
+
+    const Generator &GetDefault(const Device &device) override { return GetDefaultCudaGenerator(device.index()); }
+
+    void ManualSeedAll(uint64_t seed) override { ::infini_train::core::cuda::ManualSeedAll(seed); }
+};
+
+INFINI_TRAIN_REGISTER_GENERATOR_BACKEND(Device::DeviceType::kCUDA, CUDAGeneratorBackend);
+
+} // namespace
+} // namespace infini_train::core::cuda
diff --git a/infini_train/src/core/runtime/cuda/cuda_generator_impl.h b/infini_train/src/core/runtime/cuda/cuda_generator_impl.h
new file mode 100644
index 000000000..ad9577ad0
--- /dev/null
+++ b/infini_train/src/core/runtime/cuda/cuda_generator_impl.h
@@ -0,0 +1,36 @@
+#pragma once
+
+#include <cstdint>
+#include <memory>
+
+#include "infini_train/include/generator.h"
+#include "infini_train/include/generator_impl.h"
+
+namespace infini_train::core::cuda {
+
+class CUDAGeneratorImpl final : public GeneratorImpl {
+public:
+    explicit CUDAGeneratorImpl(int8_t device_index, uint64_t seed = Generator::kDefaultSeed);
+    ~CUDAGeneratorImpl() override = default;
+
+    void set_current_seed(uint64_t seed) override;
+    uint64_t current_seed() const override;
+    uint64_t Seed() override;
+    void set_state(const Tensor &state) override;
+    std::shared_ptr<Tensor> get_state() const override;
+
+    // The caller must hold mutex_ while reserving Philox subsequences.
+    uint64_t ReservePhiloxSubsequence(uint64_t increment);
+
+private:
+    CUDAGeneratorImpl *CloneImpl() const override;
+
+    uint64_t seed_ = Generator::kDefaultSeed;
+    uint64_t next_philox_subsequence_ = 0;
+};
+
+const Generator &GetDefaultCudaGenerator(int8_t device_index = -1);
+Generator CreateCudaGenerator(int8_t device_index, uint64_t seed = Generator::kDefaultSeed);
+void ManualSeedAll(uint64_t seed);
+
+} // namespace infini_train::core::cuda
diff --git a/infini_train/src/core/runtime/generator_backend.cc b/infini_train/src/core/runtime/generator_backend.cc
new file mode 100644
index 000000000..1283f989a
--- /dev/null
+++ b/infini_train/src/core/runtime/generator_backend.cc
@@ -0,0 +1,42 @@
+#include "infini_train/include/core/runtime/generator_backend.h"
+
+#include <ostream>
+#include <stdexcept>
+#include <string>
+#include <utility>
+
+#include "glog/logging.h"
+
+namespace infini_train {
+
+GeneratorBackendRegistry &GeneratorBackendRegistry::Instance() {
+    static GeneratorBackendRegistry instance;
+    return instance;
+}
+
+void GeneratorBackendRegistry::Register(Device::DeviceType type, std::unique_ptr<GeneratorBackend> backend) {
+    CHECK(backend != nullptr) << "Registering a null GeneratorBackend";
+    // DeviceType has no stream insertion operator.
+    CHECK_EQ(static_cast<int>(type), static_cast<int>(backend->Type()))
+        << "GeneratorBackend type mismatch: registered as " << static_cast<int>(type) << " but Type() is "
+        << static_cast<int>(backend->Type());
+    CHECK(!backends_.contains(type)) << "GeneratorBackend for device type " << static_cast<int>(type)
+                                     << " is already registered";
+
+    backends_.emplace(type, std::move(backend));
+}
+
+GeneratorBackend &GeneratorBackendRegistry::Get(Device::DeviceType type) const {
+    auto it = backends_.find(type);
+    if (it == backends_.end()) {
+        throw std::invalid_argument("No GeneratorBackend registered for device type "
+                                    + std::to_string(static_cast<int>(type)));
+    }
+    return *it->second;
+}
+
+void GeneratorBackendRegistry::ManualSeedAll(uint64_t seed) const {
+    for (const auto &entry : backends_) { entry.second->ManualSeedAll(seed); }
+}
+
+} // namespace infini_train
diff --git a/infini_train/src/generator.cc b/infini_train/src/generator.cc
new file mode 100644
index 000000000..f0693264a
--- /dev/null
+++ b/infini_train/src/generator.cc
@@ -0,0 +1,56 @@
+#include "infini_train/include/generator.h"
+#include "infini_train/include/generator_impl.h"
+
+#include <memory>
+#include <ostream>
+#include <utility>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/core/runtime/generator_backend.h"
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train {
+
+Generator::Generator(std::shared_ptr<GeneratorImpl> impl) : impl_(std::move(impl)) {
+    CHECK(impl_) << "GeneratorImpl with nullptr is not supported";
+}
+
+void Generator::set_current_seed(uint64_t seed) const { impl_->set_current_seed(seed); }
+
+uint64_t Generator::current_seed() const { return impl_->current_seed(); }
+
+uint64_t Generator::Seed() { return impl_->Seed(); }
+
+void Generator::set_state(const Tensor &state) {
+    CHECK(state.defined()) << "Undefined tensor is not allowed";
+    impl_->set_state(state);
+}
+
+std::shared_ptr<Tensor> Generator::get_state() const { return impl_->get_state(); }
+
+Device Generator::device() const { return impl_->device(); }
+
+Generator Generator::Clone() const { return Generator(impl_->Clone()); }
+
+namespace detail {
+
+void CheckRngState(const Tensor &state) {
+    CHECK(state.GetDevice().IsCPU()) << "RNG state must be a CPU tensor";
+    CHECK_EQ(static_cast<int>(state.Dtype()), static_cast<int>(DataType::kUINT8)) << "RNG state must be a UINT8 tensor";
+}
+
+} // namespace detail
+
+Generator CreateGenerator(const Device &device, uint64_t seed) {
+    return GeneratorBackendRegistry::Instance().Get(device.type()).Create(device, seed);
+}
+
+const Generator &GetDefaultGenerator(const Device &device) {
+    return GeneratorBackendRegistry::Instance().Get(device.type()).GetDefault(device);
+}
+
+void ManualSeed(uint64_t seed) { GeneratorBackendRegistry::Instance().ManualSeedAll(seed); }
+
+} // namespace infini_train
diff --git a/infini_train/src/kernels/cpu/cross_entropy.cc b/infini_train/src/kernels/cpu/cross_entropy.cc
index f520b2e96..5cf95cf21 100644
--- a/infini_train/src/kernels/cpu/cross_entropy.cc
+++ b/infini_train/src/kernels/cpu/cross_entropy.cc
@@ -2,6 +2,7 @@
 #include <cmath>
 #include <limits>
 #include <memory>
+#include <numeric>
 #include <vector>
 
 #include "glog/logging.h"
diff --git a/infini_train/src/kernels/cpu/distribution.cc b/infini_train/src/kernels/cpu/distribution.cc
new file mode 100644
index 000000000..56e332efb
--- /dev/null
+++ b/infini_train/src/kernels/cpu/distribution.cc
@@ -0,0 +1,93 @@
+#include <cstdint>
+#include <memory>
+#include <mutex>
+#include <optional>
+#include <type_traits>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/generator_impl.h"
+#include "infini_train/include/tensor.h"
+#include "infini_train/src/core/runtime/cpu/cpu_dispatch.h"
+#include "infini_train/src/core/runtime/cpu/cpu_generator_impl.h"
+#include "infini_train/src/kernels/cpu/distributions_helper.h"
+
+namespace infini_train::kernels::cpu {
+namespace {
+
+constexpr int kMaxUniformAttempts = 3;
+
+template <typename StorageT, typename RandomT>
+void UniformImpl(Tensor &tensor, double from, double to, core::cpu::CPUGeneratorImpl *generator) {
+    auto *buf = static_cast<StorageT *>(tensor.DataPtr());
+    common::cpu::UniformRealDistribution<RandomT> dist(static_cast<RandomT>(from), static_cast<RandomT>(to));
+    const StorageT from_value = static_cast<StorageT>(from);
+    if (from == to) {
+        for (int64_t i = 0; i < tensor.NumElements(); ++i) { buf[i] = from_value; }
+        return;
+    }
+    for (int64_t i = 0; i < tensor.NumElements(); ++i) {
+        StorageT value;
+        int attempt = 0;
+        do {
+            value = static_cast<StorageT>(dist(generator));
+            ++attempt;
+        } while (attempt < kMaxUniformAttempts
+                 && (static_cast<double>(value) < from || static_cast<double>(value) >= to));
+        // Bounded retries may still leave out-of-range values in very narrow intervals.
+        buf[i] = static_cast<double>(value) == to ? from_value : value;
+    }
+}
+
+template <typename StorageT, typename RandomT>
+void NormalImpl(Tensor &tensor, double mean, double std, core::cpu::CPUGeneratorImpl *generator) {
+    auto *buf = static_cast<StorageT *>(tensor.DataPtr());
+    common::cpu::NormalDistribution<RandomT> dist(static_cast<RandomT>(mean), static_cast<RandomT>(std));
+    for (int64_t i = 0; i < tensor.NumElements(); ++i) { buf[i] = static_cast<StorageT>(dist(generator)); }
+}
+
+} // namespace
+
+void Uniform(const std::shared_ptr<Tensor> tensor, double from, double to, const std::optional<Generator> gen) {
+    CHECK(tensor->GetDevice().IsCPU());
+    auto &cpu_generator = GetGeneratorOrDefault<core::cpu::CPUGeneratorImpl>(gen, core::cpu::GetDefaultCpuGenerator(),
+                                                                             tensor->GetDevice());
+
+    std::lock_guard<std::mutex> lock(cpu_generator.mutex_);
+    core::cpu::DispatchCpuFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+        tensor->Dtype(),
+        [&]<typename StorageT>() {
+            using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+            UniformImpl<StorageT, RandomT>(*tensor, from, to, &cpu_generator);
+        },
+        "CPU uniform");
+}
+
+void Normal(const std::shared_ptr<Tensor> tensor, double mean, double std, const std::optional<Generator> gen) {
+    CHECK(tensor->GetDevice().IsCPU());
+    auto &cpu_generator = GetGeneratorOrDefault<core::cpu::CPUGeneratorImpl>(gen, core::cpu::GetDefaultCpuGenerator(),
+                                                                             tensor->GetDevice());
+
+    std::lock_guard<std::mutex> lock(cpu_generator.mutex_);
+    core::cpu::DispatchCpuFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+        tensor->Dtype(),
+        [&]<typename StorageT>() {
+            using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+            NormalImpl<StorageT, RandomT>(*tensor, mean, std, &cpu_generator);
+        },
+        "CPU normal");
+}
+
+} // namespace infini_train::kernels::cpu
+
+#define REGISTER_CPU_DISTRIBUTION_KERNEL(kernel_name)                                                                  \
+    REGISTER_KERNEL(infini_train::Device::DeviceType::kCPU, kernel_name, infini_train::kernels::cpu::kernel_name)
+
+REGISTER_CPU_DISTRIBUTION_KERNEL(Uniform)
+REGISTER_CPU_DISTRIBUTION_KERNEL(Normal)
+
+#undef REGISTER_CPU_DISTRIBUTION_KERNEL
diff --git a/infini_train/src/kernels/cpu/distributions_helper.h b/infini_train/src/kernels/cpu/distributions_helper.h
new file mode 100644
index 000000000..6e2c97e65
--- /dev/null
+++ b/infini_train/src/kernels/cpu/distributions_helper.h
@@ -0,0 +1,118 @@
+#pragma once
+
+// Host-side uniform and normal distributions for generators exposing Random()
+// and Random64(). Box-Muller's second sample is cached when supported by the generator.
+
+#include <cmath>
+#include <cstdint>
+#include <limits>
+#include <numbers>
+#include <optional>
+#include <type_traits>
+
+#include "glog/logging.h"
+
+namespace infini_train::common::cpu {
+
+template <typename T> struct UniformRealDistribution {
+    UniformRealDistribution(T from, T to) : from_(from), to_(to) {
+        CHECK_LE(from, to);
+        CHECK_LE(to - from, std::numeric_limits<T>::max());
+    }
+
+    UniformRealDistribution(const UniformRealDistribution &) = default;
+    UniformRealDistribution &operator=(const UniformRealDistribution &) = delete;
+
+    template <typename RNG> T operator()(RNG *generator) const {
+        if constexpr (std::is_same_v<T, double>) {
+            return Transform(generator->Random64());
+        } else {
+            return Transform(generator->Random());
+        }
+    }
+
+private:
+    T from_;
+    T to_;
+
+    template <typename V> T Transform(V val) const {
+        constexpr auto kMask = static_cast<V>((static_cast<uint64_t>(1) << std::numeric_limits<T>::digits) - 1);
+        constexpr auto kDivisor = static_cast<T>(1) / (static_cast<uint64_t>(1) << std::numeric_limits<T>::digits);
+        T x = (val & kMask) * kDivisor;
+        return x * (to_ - from_) + from_;
+    }
+};
+
+template <typename RNG, typename = decltype(&RNG::next_double_normal_sample),
+          typename = decltype(&RNG::set_next_double_normal_sample)>
+bool MaybeGetNextNormalSample(RNG *generator, double *ret) {
+    const auto sample = generator->next_double_normal_sample();
+    if (!sample.has_value()) {
+        return false;
+    }
+    *ret = sample.value();
+    generator->set_next_double_normal_sample(std::nullopt);
+    return true;
+}
+
+template <typename RNG, typename = decltype(&RNG::next_float_normal_sample),
+          typename = decltype(&RNG::set_next_float_normal_sample)>
+bool MaybeGetNextNormalSample(RNG *generator, float *ret) {
+    const auto sample = generator->next_float_normal_sample();
+    if (!sample.has_value()) {
+        return false;
+    }
+    *ret = sample.value();
+    generator->set_next_float_normal_sample(std::nullopt);
+    return true;
+}
+
+// Fallback: RNG without cache support never has a cached sample.
+template <typename RNG> bool MaybeGetNextNormalSample(RNG * /*generator*/, void * /*ret*/) { return false; }
+
+template <typename RNG, typename = decltype(&RNG::set_next_double_normal_sample)>
+void MaybeSetNextNormalSample(RNG *generator, const double *cache) {
+    generator->set_next_double_normal_sample(*cache);
+}
+
+template <typename RNG, typename = decltype(&RNG::set_next_float_normal_sample)>
+void MaybeSetNextNormalSample(RNG *generator, const float *cache) {
+    generator->set_next_float_normal_sample(*cache);
+}
+
+// Fallback: RNG without cache support discards the second sample.
+template <typename RNG> void MaybeSetNextNormalSample(RNG * /*generator*/, const void * /*cache*/) {}
+
+template <typename T> struct NormalDistribution {
+    NormalDistribution(T mean, T stdv) : mean_(mean), stdv_(stdv) { CHECK_GE(stdv, static_cast<T>(0)); }
+
+    NormalDistribution(const NormalDistribution &) = default;
+    NormalDistribution &operator=(const NormalDistribution &) = delete;
+
+    template <typename RNG> T operator()(RNG *generator) const {
+        T ret;
+        if (MaybeGetNextNormalSample(generator, &ret)) {
+            return ret * stdv_ + mean_;
+        }
+
+        UniformRealDistribution<T> uniform(static_cast<T>(0), static_cast<T>(1));
+        const T u1 = uniform(generator);
+        const T u2 = uniform(generator);
+
+        const T r = std::sqrt(static_cast<T>(-2.0) * std::log1p(-u2));
+        constexpr T kTwoPi = static_cast<T>(2.0 * std::numbers::pi_v<double>);
+        const T theta = kTwoPi * u1;
+        const T sample = r * std::sin(theta);
+
+        MaybeSetNextNormalSample(generator, &sample);
+
+        ret = r * std::cos(theta);
+        return ret * stdv_ + mean_;
+    }
+
+private:
+    T mean_;
+    T stdv_;
+};
+
+} // namespace infini_train::common::cpu
diff --git a/infini_train/src/kernels/cpu/dropout.cc b/infini_train/src/kernels/cpu/dropout.cc
new file mode 100644
index 000000000..5ea78574d
--- /dev/null
+++ b/infini_train/src/kernels/cpu/dropout.cc
@@ -0,0 +1,130 @@
+#include <cstdint>
+#include <memory>
+#include <mutex>
+#include <optional>
+#include <ostream>
+#include <tuple>
+#include <type_traits>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/generator_impl.h"
+#include "infini_train/include/tensor.h"
+#include "infini_train/src/core/runtime/cpu/cpu_dispatch.h"
+#include "infini_train/src/core/runtime/cpu/cpu_generator_impl.h"
+#include "infini_train/src/kernels/cpu/distributions_helper.h"
+
+namespace infini_train::kernels::cpu {
+namespace {
+
+template <typename StorageT, typename RandomT>
+void DropoutForwardImpl(Tensor &output, Tensor &mask, const Tensor &input, double p,
+                        core::cpu::CPUGeneratorImpl *generator) {
+    auto *output_data = static_cast<StorageT *>(output.DataPtr());
+    auto *mask_data = static_cast<uint8_t *>(mask.DataPtr());
+    const auto *input_data = static_cast<const StorageT *>(input.DataPtr());
+    const int64_t n = input.NumElements();
+
+    if (p == 0.0) {
+        for (int64_t index = 0; index < n; ++index) {
+            mask_data[index] = 1;
+            output_data[index] = input_data[index];
+        }
+        return;
+    }
+    if (p == 1.0) {
+        for (int64_t index = 0; index < n; ++index) {
+            mask_data[index] = 0;
+            output_data[index] = static_cast<StorageT>(0.0);
+        }
+        return;
+    }
+
+    const RandomT scale = static_cast<RandomT>(1.0 / (1.0 - p));
+    common::cpu::UniformRealDistribution<RandomT> distribution(static_cast<RandomT>(0), static_cast<RandomT>(1));
+    for (int64_t index = 0; index < n; ++index) {
+        const bool keep = distribution(generator) >= static_cast<RandomT>(p);
+        mask_data[index] = keep ? 1 : 0;
+        output_data[index] = keep ? static_cast<StorageT>(static_cast<RandomT>(input_data[index]) * scale)
+                                  : static_cast<StorageT>(0.0);
+    }
+}
+
+template <typename StorageT, typename RandomT>
+void DropoutBackwardImpl(Tensor &grad_input, const Tensor &grad_output, const Tensor &mask, double p) {
+    auto *grad_input_data = static_cast<StorageT *>(grad_input.DataPtr());
+    const auto *grad_output_data = static_cast<const StorageT *>(grad_output.DataPtr());
+    const auto *mask_data = static_cast<const uint8_t *>(mask.DataPtr());
+    const RandomT scale = p == 1.0 ? static_cast<RandomT>(0) : static_cast<RandomT>(1.0 / (1.0 - p));
+
+    for (int64_t index = 0; index < grad_output.NumElements(); ++index) {
+        grad_input_data[index] = mask_data[index]
+                                   ? static_cast<StorageT>(static_cast<RandomT>(grad_output_data[index]) * scale)
+                                   : static_cast<StorageT>(0.0);
+    }
+}
+
+} // namespace
+
+std::tuple<std::shared_ptr<Tensor>, std::shared_ptr<Tensor>>
+DropoutForward(const std::shared_ptr<Tensor> input, double p, const std::optional<Generator> gen) {
+    CHECK(input->GetDevice().IsCPU());
+    CHECK(IsFloatingPointDType(input->Dtype())) << "Dropout supports floating-point tensors only";
+    CHECK_GE(p, 0.0) << "dropout probability has to be between 0 and 1, but got " << p;
+    CHECK_LE(p, 1.0) << "dropout probability has to be between 0 and 1, but got " << p;
+
+    auto output = std::make_shared<Tensor>(input->Dims(), input->Dtype(), input->GetDevice());
+    auto mask = std::make_shared<Tensor>(input->Dims(), DataType::kUINT8, input->GetDevice());
+
+    if (p == 0.0 || p == 1.0) {
+        core::cpu::DispatchCpuFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+            input->Dtype(),
+            [&]<typename StorageT>() {
+                using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+                DropoutForwardImpl<StorageT, RandomT>(*output, *mask, *input, p, nullptr);
+            },
+            "CPU dropout forward");
+        return {output, mask};
+    }
+
+    auto &cpu_generator = GetGeneratorOrDefault<core::cpu::CPUGeneratorImpl>(gen, core::cpu::GetDefaultCpuGenerator(),
+                                                                             input->GetDevice());
+    std::lock_guard<std::mutex> lock(cpu_generator.mutex_);
+    core::cpu::DispatchCpuFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+        input->Dtype(),
+        [&]<typename StorageT>() {
+            using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+            DropoutForwardImpl<StorageT, RandomT>(*output, *mask, *input, p, &cpu_generator);
+        },
+        "CPU dropout forward");
+    return {output, mask};
+}
+
+std::shared_ptr<Tensor> DropoutBackward(const std::shared_ptr<Tensor> grad_output, const std::shared_ptr<Tensor> mask,
+                                        double p) {
+    CHECK(grad_output->GetDevice().IsCPU());
+    auto grad_input = std::make_shared<Tensor>(grad_output->Dims(), grad_output->Dtype(), grad_output->GetDevice());
+
+    core::cpu::DispatchCpuFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+        grad_output->Dtype(),
+        [&]<typename StorageT>() {
+            using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+            DropoutBackwardImpl<StorageT, RandomT>(*grad_input, *grad_output, *mask, p);
+        },
+        "CPU dropout backward");
+    return grad_input;
+}
+
+} // namespace infini_train::kernels::cpu
+
+#define REGISTER_CPU_DROPOUT_KERNEL(kernel_name)                                                                       \
+    REGISTER_KERNEL(infini_train::Device::DeviceType::kCPU, kernel_name, infini_train::kernels::cpu::kernel_name)
+
+REGISTER_CPU_DROPOUT_KERNEL(DropoutForward)
+REGISTER_CPU_DROPOUT_KERNEL(DropoutBackward)
+
+#undef REGISTER_CPU_DROPOUT_KERNEL
diff --git a/infini_train/src/kernels/cpu/transform.cc b/infini_train/src/kernels/cpu/transform.cc
index 48063c7a0..a76102b19 100644
--- a/infini_train/src/kernels/cpu/transform.cc
+++ b/infini_train/src/kernels/cpu/transform.cc
@@ -2,6 +2,7 @@
 #include <cstddef>
 #include <cstring>
 #include <memory>
+#include <numeric>
 
 #include "glog/logging.h"
 
diff --git a/infini_train/src/kernels/cuda/distribution.cu b/infini_train/src/kernels/cuda/distribution.cu
new file mode 100644
index 000000000..9d17d7ad4
--- /dev/null
+++ b/infini_train/src/kernels/cuda/distribution.cu
@@ -0,0 +1,169 @@
+#include <cstdint>
+#include <memory>
+#include <mutex>
+#include <optional>
+#include <type_traits>
+
+#include <curand_kernel.h>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/common/cuda/common_cuda.h"
+#include "infini_train/include/common/cuda/kernel_helper.cuh"
+#include "infini_train/include/core/runtime/device_guard.h"
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/generator_impl.h"
+#include "infini_train/include/tensor.h"
+#include "infini_train/src/core/runtime/cuda/cuda_dispatch.h"
+#include "infini_train/src/core/runtime/cuda/cuda_generator_impl.h"
+#include "infini_train/src/core/runtime/cuda/cuda_runtime_common.h"
+
+namespace infini_train::kernels::cuda {
+namespace {
+
+constexpr int kThreadsPerBlock = 256;
+constexpr int kMaxUniformAttempts = 3;
+
+template <typename RandomT> __device__ RandomT UniformSample(curandStatePhilox4_32_10_t *state) {
+    if constexpr (std::is_same_v<RandomT, double>) {
+        double sample;
+        do {
+            const double2 values = curand_uniform2_double(state);
+            sample = 1.0 - values.x;
+            // Subtraction can round a tiny positive draw to 1.0.
+        } while (sample >= 1.0);
+        return sample;
+    } else {
+        return static_cast<float>(curand(state)) * 0x1p-32f;
+    }
+}
+
+template <typename RandomT> __device__ RandomT NormalSample(curandStatePhilox4_32_10_t *state) {
+    if constexpr (std::is_same_v<RandomT, double>) {
+        return curand_normal_double(state);
+    } else {
+        return curand_normal(state);
+    }
+}
+
+template <typename StorageT, typename RandomT>
+__global__ void UniformKernel(StorageT *data, int64_t n, double from, double to, uint64_t seed, uint64_t subsequence) {
+    const int64_t index = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    if (index >= n) {
+        return;
+    }
+
+    const StorageT from_value = common::cuda::Cast<StorageT>(from);
+    if (from == to) {
+        data[index] = from_value;
+        return;
+    }
+    curandStatePhilox4_32_10_t state;
+    curand_init(seed, subsequence + static_cast<uint64_t>(index), 0, &state);
+    const RandomT sample_from = static_cast<RandomT>(from);
+    const RandomT sample_to = static_cast<RandomT>(to);
+    StorageT value;
+    int attempt = 0;
+    do {
+        value = common::cuda::Cast<StorageT>(sample_from + UniformSample<RandomT>(&state) * (sample_to - sample_from));
+        ++attempt;
+    } while (attempt < kMaxUniformAttempts
+             && (common::cuda::Cast<double>(value) < from || common::cuda::Cast<double>(value) >= to));
+    // Bounded retries may still leave out-of-range values in very narrow intervals.
+    data[index] = common::cuda::Cast<double>(value) == to ? from_value : value;
+}
+
+template <typename StorageT, typename RandomT>
+__global__ void NormalKernel(StorageT *data, int64_t n, RandomT mean, RandomT std, uint64_t seed,
+                             uint64_t subsequence) {
+    const int64_t index = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    if (index >= n) {
+        return;
+    }
+
+    curandStatePhilox4_32_10_t state;
+    curand_init(seed, subsequence + static_cast<uint64_t>(index), 0, &state);
+    data[index] = common::cuda::Cast<StorageT>(mean + NormalSample<RandomT>(&state) * std);
+}
+
+const core::cuda::CudaStream *GetCudaStream(const Device &device) {
+    return dynamic_cast<const core::cuda::CudaStream *>(core::GetDeviceGuardImpl(device.type())->GetStream(device));
+}
+
+} // namespace
+
+void Uniform(const std::shared_ptr<Tensor> tensor, double from, double to, const std::optional<Generator> gen) {
+    const Device device = tensor->GetDevice();
+    CHECK(device.IsCUDA());
+    const int64_t n = tensor->NumElements();
+    if (n == 0) {
+        return;
+    }
+    core::DeviceGuard guard(device);
+    auto &cuda_generator = GetGeneratorOrDefault<core::cuda::CUDAGeneratorImpl>(
+        gen, core::cuda::GetDefaultCudaGenerator(device.index()), device);
+
+    uint64_t seed = 0;
+    uint64_t subsequence = 0;
+    if (from != to) {
+        std::lock_guard<std::mutex> lock(cuda_generator.mutex_);
+        seed = cuda_generator.current_seed();
+        subsequence = cuda_generator.ReservePhiloxSubsequence(static_cast<uint64_t>(n));
+    }
+
+    const int blocks = static_cast<int>((n + kThreadsPerBlock - 1) / kThreadsPerBlock);
+    const auto *stream = GetCudaStream(device);
+    core::cuda::DispatchCudaFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+        tensor->Dtype(),
+        [&]<typename StorageT>() {
+            using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+            UniformKernel<StorageT, RandomT><<<blocks, kThreadsPerBlock, 0, stream->cuda_stream()>>>(
+                static_cast<StorageT *>(tensor->DataPtr()), n, from, to, seed, subsequence);
+        },
+        "CUDA uniform");
+    CUDA_CHECK(cudaGetLastError());
+}
+
+void Normal(const std::shared_ptr<Tensor> tensor, double mean, double std, const std::optional<Generator> gen) {
+    const Device device = tensor->GetDevice();
+    CHECK(device.IsCUDA());
+    const int64_t n = tensor->NumElements();
+    if (n == 0) {
+        return;
+    }
+    core::DeviceGuard guard(device);
+    auto &cuda_generator = GetGeneratorOrDefault<core::cuda::CUDAGeneratorImpl>(
+        gen, core::cuda::GetDefaultCudaGenerator(device.index()), device);
+
+    uint64_t seed = 0;
+    uint64_t subsequence = 0;
+    {
+        std::lock_guard<std::mutex> lock(cuda_generator.mutex_);
+        seed = cuda_generator.current_seed();
+        subsequence = cuda_generator.ReservePhiloxSubsequence(static_cast<uint64_t>(n));
+    }
+
+    const int blocks = static_cast<int>((n + kThreadsPerBlock - 1) / kThreadsPerBlock);
+    const auto *stream = GetCudaStream(device);
+    core::cuda::DispatchCudaFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+        tensor->Dtype(),
+        [&]<typename StorageT>() {
+            using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+            NormalKernel<StorageT, RandomT><<<blocks, kThreadsPerBlock, 0, stream->cuda_stream()>>>(
+                static_cast<StorageT *>(tensor->DataPtr()), n, static_cast<RandomT>(mean), static_cast<RandomT>(std),
+                seed, subsequence);
+        },
+        "CUDA normal");
+    CUDA_CHECK(cudaGetLastError());
+}
+
+} // namespace infini_train::kernels::cuda
+
+#define REGISTER_CUDA_DISTRIBUTION_KERNEL(kernel_name)                                                                 \
+    REGISTER_KERNEL(infini_train::Device::DeviceType::kCUDA, kernel_name, infini_train::kernels::cuda::kernel_name)
+
+REGISTER_CUDA_DISTRIBUTION_KERNEL(Uniform)
+REGISTER_CUDA_DISTRIBUTION_KERNEL(Normal)
+
+#undef REGISTER_CUDA_DISTRIBUTION_KERNEL
diff --git a/infini_train/src/kernels/cuda/dropout.cu b/infini_train/src/kernels/cuda/dropout.cu
new file mode 100644
index 000000000..ad212267a
--- /dev/null
+++ b/infini_train/src/kernels/cuda/dropout.cu
@@ -0,0 +1,168 @@
+#include <cstdint>
+#include <memory>
+#include <mutex>
+#include <optional>
+#include <tuple>
+#include <type_traits>
+
+#include <curand_kernel.h>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/common/cuda/common_cuda.h"
+#include "infini_train/include/common/cuda/kernel_helper.cuh"
+#include "infini_train/include/core/runtime/device_guard.h"
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/generator_impl.h"
+#include "infini_train/include/tensor.h"
+#include "infini_train/src/core/runtime/cuda/cuda_dispatch.h"
+#include "infini_train/src/core/runtime/cuda/cuda_generator_impl.h"
+#include "infini_train/src/core/runtime/cuda/cuda_runtime_common.h"
+
+namespace infini_train::kernels::cuda {
+namespace {
+
+constexpr int kThreadsPerBlock = 256;
+
+template <typename RandomT> __device__ RandomT UniformSample(curandStatePhilox4_32_10_t *state) {
+    if constexpr (std::is_same_v<RandomT, double>) {
+        double sample;
+        do {
+            const double2 values = curand_uniform2_double(state);
+            sample = 1.0 - values.x;
+            // Subtraction can round a tiny positive draw to 1.0.
+        } while (sample >= 1.0);
+        return sample;
+    } else {
+        return static_cast<float>(curand(state)) * 0x1p-32f;
+    }
+}
+
+template <typename StorageT, typename RandomT>
+__global__ void DropoutForwardKernel(StorageT *output, uint8_t *mask, const StorageT *input, int64_t n, RandomT p,
+                                     uint64_t seed, uint64_t subsequence) {
+    const int64_t index = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    if (index >= n) {
+        return;
+    }
+    if (p == static_cast<RandomT>(0)) {
+        mask[index] = 1;
+        output[index] = input[index];
+        return;
+    }
+    if (p == static_cast<RandomT>(1)) {
+        mask[index] = 0;
+        output[index] = common::cuda::Cast<StorageT>(0.0f);
+        return;
+    }
+
+    curandStatePhilox4_32_10_t state;
+    curand_init(seed, subsequence + static_cast<uint64_t>(index), 0, &state);
+    const bool keep = UniformSample<RandomT>(&state) >= p;
+    const RandomT scale = static_cast<RandomT>(1) / (static_cast<RandomT>(1) - p);
+    mask[index] = keep ? 1 : 0;
+    output[index] = keep ? common::cuda::Cast<StorageT>(common::cuda::Cast<RandomT>(input[index]) * scale)
+                         : common::cuda::Cast<StorageT>(0.0f);
+}
+
+template <typename StorageT, typename RandomT>
+__global__ void DropoutBackwardKernel(StorageT *grad_input, const StorageT *grad_output, const uint8_t *mask, int64_t n,
+                                      RandomT p) {
+    const int64_t index = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    if (index >= n) {
+        return;
+    }
+    const RandomT scale = p == static_cast<RandomT>(1) ? static_cast<RandomT>(0)
+                                                       : static_cast<RandomT>(1) / (static_cast<RandomT>(1) - p);
+    grad_input[index] = mask[index]
+                          ? common::cuda::Cast<StorageT>(common::cuda::Cast<RandomT>(grad_output[index]) * scale)
+                          : common::cuda::Cast<StorageT>(0.0f);
+}
+
+const core::cuda::CudaStream *GetCudaStream(const Device &device) {
+    return dynamic_cast<const core::cuda::CudaStream *>(core::GetDeviceGuardImpl(device.type())->GetStream(device));
+}
+
+} // namespace
+
+std::tuple<std::shared_ptr<Tensor>, std::shared_ptr<Tensor>>
+DropoutForward(const std::shared_ptr<Tensor> input, double p, const std::optional<Generator> gen) {
+    const Device device = input->GetDevice();
+    CHECK(device.IsCUDA());
+    CHECK(IsFloatingPointDType(input->Dtype())) << "Dropout supports floating-point tensors only";
+    CHECK_GE(p, 0.0) << "dropout probability has to be between 0 and 1, but got " << p;
+    CHECK_LE(p, 1.0) << "dropout probability has to be between 0 and 1, but got " << p;
+    core::DeviceGuard guard(device);
+
+    auto output = std::make_shared<Tensor>(input->Dims(), input->Dtype(), device);
+    auto mask = std::make_shared<Tensor>(input->Dims(), DataType::kUINT8, device);
+
+    const int64_t n = input->NumElements();
+    if (n == 0) {
+        return {output, mask};
+    }
+
+    uint64_t seed = 0;
+    uint64_t subsequence = 0;
+    if (p > 0.0 && p < 1.0) {
+        auto &cuda_generator = GetGeneratorOrDefault<core::cuda::CUDAGeneratorImpl>(
+            gen, core::cuda::GetDefaultCudaGenerator(device.index()), device);
+        std::lock_guard<std::mutex> lock(cuda_generator.mutex_);
+        seed = cuda_generator.current_seed();
+        subsequence = cuda_generator.ReservePhiloxSubsequence(static_cast<uint64_t>(n));
+    }
+
+    const int blocks = static_cast<int>((n + kThreadsPerBlock - 1) / kThreadsPerBlock);
+    const auto *stream = GetCudaStream(device);
+    core::cuda::DispatchCudaFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+        input->Dtype(),
+        [&]<typename StorageT>() {
+            using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+            DropoutForwardKernel<StorageT, RandomT><<<blocks, kThreadsPerBlock, 0, stream->cuda_stream()>>>(
+                static_cast<StorageT *>(output->DataPtr()), static_cast<uint8_t *>(mask->DataPtr()),
+                static_cast<const StorageT *>(input->DataPtr()), n, static_cast<RandomT>(p), seed, subsequence);
+        },
+        "CUDA dropout forward");
+    CUDA_CHECK(cudaGetLastError());
+    return {output, mask};
+}
+
+std::shared_ptr<Tensor> DropoutBackward(const std::shared_ptr<Tensor> grad_output, const std::shared_ptr<Tensor> mask,
+                                        double p) {
+    const Device device = grad_output->GetDevice();
+    CHECK(device.IsCUDA());
+    core::DeviceGuard guard(device);
+
+    auto grad_input = std::make_shared<Tensor>(grad_output->Dims(), grad_output->Dtype(), device);
+
+    const int64_t n = grad_output->NumElements();
+    if (n == 0) {
+        return grad_input;
+    }
+
+    const int blocks = static_cast<int>((n + kThreadsPerBlock - 1) / kThreadsPerBlock);
+    const auto *stream = GetCudaStream(device);
+    core::cuda::DispatchCudaFunc<DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64>(
+        grad_output->Dtype(),
+        [&]<typename StorageT>() {
+            using RandomT = std::conditional_t<std::is_same_v<StorageT, double>, double, float>;
+            DropoutBackwardKernel<StorageT, RandomT><<<blocks, kThreadsPerBlock, 0, stream->cuda_stream()>>>(
+                static_cast<StorageT *>(grad_input->DataPtr()), static_cast<const StorageT *>(grad_output->DataPtr()),
+                static_cast<const uint8_t *>(mask->DataPtr()), n, static_cast<RandomT>(p));
+        },
+        "CUDA dropout backward");
+    CUDA_CHECK(cudaGetLastError());
+    return grad_input;
+}
+
+} // namespace infini_train::kernels::cuda
+
+#define REGISTER_CUDA_DROPOUT_KERNEL(kernel_name)                                                                      \
+    REGISTER_KERNEL(infini_train::Device::DeviceType::kCUDA, kernel_name, infini_train::kernels::cuda::kernel_name)
+
+REGISTER_CUDA_DROPOUT_KERNEL(DropoutForward)
+REGISTER_CUDA_DROPOUT_KERNEL(DropoutBackward)
+
+#undef REGISTER_CUDA_DROPOUT_KERNEL
diff --git a/infini_train/src/kernels/cuda/elementwise.cu b/infini_train/src/kernels/cuda/elementwise.cu
index fc423b35f..a4b6dae83 100644
--- a/infini_train/src/kernels/cuda/elementwise.cu
+++ b/infini_train/src/kernels/cuda/elementwise.cu
@@ -1,4 +1,5 @@
 #include <cstddef>
+#include <numeric>
 
 #include <cub/warp/warp_reduce.cuh>
 
diff --git a/infini_train/src/kernels/cuda/gather.cu b/infini_train/src/kernels/cuda/gather.cu
index d8b0cffa5..7b3862ac5 100644
--- a/infini_train/src/kernels/cuda/gather.cu
+++ b/infini_train/src/kernels/cuda/gather.cu
@@ -1,3 +1,5 @@
+#include <numeric>
+
 #include "glog/logging.h"
 
 #include "infini_train/include/common/common.h"
diff --git a/infini_train/src/kernels/cuda/no_op.cu b/infini_train/src/kernels/cuda/no_op.cu
index ef2c95668..d26024b0d 100644
--- a/infini_train/src/kernels/cuda/no_op.cu
+++ b/infini_train/src/kernels/cuda/no_op.cu
@@ -1,3 +1,5 @@
+#include <numeric>
+
 #include "glog/logging.h"
 
 #include "infini_train/include/dispatcher.h"
diff --git a/infini_train/src/kernels/cuda/reduction.cu b/infini_train/src/kernels/cuda/reduction.cu
index b9fb57e0d..494e7acb2 100644
--- a/infini_train/src/kernels/cuda/reduction.cu
+++ b/infini_train/src/kernels/cuda/reduction.cu
@@ -1,3 +1,5 @@
+#include <numeric>
+
 #include <cub/cub.cuh>
 
 #include "infini_train/include/common/cuda/common_cuda.h"
diff --git a/infini_train/src/nn/functional.cc b/infini_train/src/nn/functional.cc
index c33e23684..a28c46c1d 100644
--- a/infini_train/src/nn/functional.cc
+++ b/infini_train/src/nn/functional.cc
@@ -2,9 +2,13 @@
 
 #include <cstdint>
 #include <memory>
+#include <utility>
 #include <vector>
 
+#include "glog/logging.h"
+
 #include "infini_train/include/autograd/activations.h"
+#include "infini_train/include/autograd/dropout.h"
 #include "infini_train/include/autograd/elementwise.h"
 #include "infini_train/include/autograd/reduction.h"
 #include "infini_train/include/autograd/softmax.h"
@@ -26,6 +30,29 @@ std::shared_ptr<Tensor> Ones(const std::vector<int64_t> size) {
     return init::Ones(ones);
 }
 
+std::shared_ptr<Tensor> Rand(const std::vector<int64_t> &size, DataType dtype, Device device,
+                             std::optional<Generator> generator, bool requires_grad) {
+    auto result = std::make_shared<Tensor>(size, dtype, device, requires_grad);
+    return init::Uniform(result, 0.0f, 1.0f, generator);
+}
+
+std::shared_ptr<Tensor> Randn(const std::vector<int64_t> &size, DataType dtype, Device device,
+                              std::optional<Generator> generator, bool requires_grad) {
+    auto result = std::make_shared<Tensor>(size, dtype, device, requires_grad);
+    return init::Normal(result, 0.0f, 1.0f, generator);
+}
+
+std::shared_ptr<Tensor> Dropout(const std::shared_ptr<Tensor> &input, double p, bool training,
+                                std::optional<Generator> generator) {
+    CHECK(input != nullptr);
+    CHECK_GE(p, 0.0) << "dropout probability has to be between 0 and 1, but got " << p;
+    CHECK_LE(p, 1.0) << "dropout probability has to be between 0 and 1, but got " << p;
+    if (!training || p == 0.0 || input->NumElements() == 0) {
+        return input;
+    }
+    return std::make_shared<autograd::Dropout>(p, std::move(generator))->Apply({input})[0];
+}
+
 std::shared_ptr<Tensor> Reciprocal(const std::shared_ptr<Tensor> &input) { return input->Reciprocal(); }
 
 std::shared_ptr<Tensor> Sin(const std::shared_ptr<Tensor> &input) { return input->Sin(); }
diff --git a/infini_train/src/nn/init.cc b/infini_train/src/nn/init.cc
index 79b4b48b1..e64abfa42 100644
--- a/infini_train/src/nn/init.cc
+++ b/infini_train/src/nn/init.cc
@@ -1,67 +1,83 @@
 #include "infini_train/include/nn/init.h"
 
-#include <cstring>
+#include <cmath>
 #include <functional>
+#include <limits>
 #include <memory>
-#include <optional>
+#include <numeric>
+#include <ostream>
 #include <random>
+#include <string>
+#include <unordered_map>
 #include <unordered_set>
-
-#ifdef USE_OMP
-#include <omp.h>
-#endif
+#include <vector>
 
 #include "glog/logging.h"
 
 #include "infini_train/include/core/runtime/device_guard.h"
+#include "infini_train/include/datatype.h"
 #include "infini_train/include/device.h"
+#include "infini_train/include/dispatcher.h"
 #include "infini_train/include/tensor.h"
 
 namespace infini_train::nn::init {
 namespace {
-constexpr int kRandomSeed = 42;
-
-// FIXME: RNG design is incomplete.
-//
-// Current implementation lacks:
-//   - unified Generator abstraction
-//   - global default generator and seed control
-//   - reproducible / clonable RNG state
-//
-// TODO:
-//   - introduce Generator interface and backend impl
-//   - add default generator management (per device)
-//   - refactor random ops to consume Generator
-static std::mt19937 gen(kRandomSeed);
-} // namespace
 
-std::shared_ptr<Tensor> Normal(const std::shared_ptr<Tensor> &tensor, float mean, float std,
-                               std::optional<std::mt19937> generator) {
-    const int64_t num_elements = tensor->NumElements();
-    std::vector<float> buffer(num_elements);
-
-#ifdef USE_OMP
-#pragma omp parallel
-    {
-        std::mt19937 local_gen(kRandomSeed + omp_get_thread_num());
-        std::normal_distribution<float> local_dis(mean, std);
-#pragma omp for
-        for (int i = 0; i < buffer.size(); ++i) {
-            buffer[i] = generator ? local_dis(generator.value()) : local_dis(local_gen);
-        }
+struct DistributionBounds {
+    double lowest;
+    double max;
+};
+
+DistributionBounds GetDistributionBounds(DataType dtype) {
+    switch (dtype) {
+    case DataType::kFLOAT16: {
+        const double max = static_cast<float>(FP16(static_cast<uint16_t>(0x7bff), FP16::from_bits()));
+        return {-max, max};
+    }
+    case DataType::kBFLOAT16: {
+        const double max = static_cast<float>(BF16(static_cast<uint16_t>(0x7f7f), BF16::from_bits()));
+        return {-max, max};
+    }
+    case DataType::kFLOAT32:
+        return {-std::numeric_limits<float>::max(), std::numeric_limits<float>::max()};
+    case DataType::kFLOAT64:
+        return {-std::numeric_limits<double>::max(), std::numeric_limits<double>::max()};
+    default:
+        LOG(FATAL) << "Unsupported distribution dtype: " << kDataTypeToDesc.at(dtype);
+        return {};
     }
-#else
-    std::normal_distribution<float> dis(mean, std);
-    std::generate(buffer.begin(), buffer.end(), [&]() { return generator ? dis(generator.value()) : dis(gen); });
-#endif
+}
+
+void CheckDistributionTensor(const Tensor &tensor) {
+    CHECK(IsFloatingPointDType(tensor.Dtype()))
+        << "Uniform and Normal initialization support floating-point tensors only";
+}
+
+void CheckUniformParameters(const Tensor &tensor, double from, double to) {
+    const auto bounds = GetDistributionBounds(tensor.Dtype());
+    CHECK_GE(from, bounds.lowest) << "uniform expects from to be within the range of "
+                                  << kDataTypeToDesc.at(tensor.Dtype());
+    CHECK_LE(from, bounds.max) << "uniform expects from to be within the range of "
+                               << kDataTypeToDesc.at(tensor.Dtype());
+    CHECK_GE(to, bounds.lowest) << "uniform expects to to be within the range of "
+                                << kDataTypeToDesc.at(tensor.Dtype());
+    CHECK_LE(to, bounds.max) << "uniform expects to to be within the range of " << kDataTypeToDesc.at(tensor.Dtype());
+    CHECK_LE(from, to) << "uniform expects a [from, to) range, but found from=" << from << " > to=" << to;
+    CHECK_LE(to - from, bounds.max) << "uniform expects to - from to fit in " << kDataTypeToDesc.at(tensor.Dtype());
+}
+
+void CheckNormalParameters(double std) { CHECK_GE(std, 0.0) << "normal expects std >= 0.0, but found std=" << std; }
+
+} // namespace
 
+std::shared_ptr<Tensor> Normal(const std::shared_ptr<Tensor> &tensor, float mean, float std,
+                               std::optional<Generator> generator) {
+    CheckDistributionTensor(*tensor);
+    CheckNormalParameters(std);
     auto device = tensor->GetDevice();
     core::DeviceGuard guard(device);
-    auto impl = core::GetDeviceGuardImpl(device.type());
-
-    impl->MemcpyAsync(tensor->DataPtr(), buffer.data(), num_elements * sizeof(float),
-                      device.type() == Device::DeviceType::kCPU ? core::MemcpyKind::kD2D : core::MemcpyKind::kH2D,
-                      impl->GetStream(device));
+    Dispatcher::Instance().Call<void>({device.type(), "Normal"}, tensor, static_cast<double>(mean),
+                                      static_cast<double>(std), generator);
     return tensor;
 }
 
@@ -113,7 +129,7 @@ float CalculateGain(NonLinearityType nonlinearity, std::optional<float> param =
 } // namespace
 
 std::shared_ptr<Tensor> KaimingUniform(const std::shared_ptr<Tensor> &tensor, float a, KaimingMode mode,
-                                       NonLinearityType nonlinearity, std::optional<std::mt19937> generator) {
+                                       NonLinearityType nonlinearity, std::optional<Generator> generator) {
     for (const auto dim : tensor->Dims()) {
         if (dim == 0) {
             LOG(WARNING) << "Initializing zero-element tensors is a no-op";
@@ -128,34 +144,13 @@ std::shared_ptr<Tensor> KaimingUniform(const std::shared_ptr<Tensor> &tensor, fl
 }
 
 std::shared_ptr<Tensor> Uniform(const std::shared_ptr<Tensor> &tensor, float a, float b,
-                                std::optional<std::mt19937> generator) {
-    const int64_t num_elements = tensor->NumElements();
-    std::vector<float> buffer(num_elements);
-
-#ifdef USE_OMP
-#pragma omp parallel
-    {
-        std::mt19937 local_gen(kRandomSeed + omp_get_thread_num());
-        std::uniform_real_distribution<float> local_dis(a, b);
-#pragma omp for
-        for (int i = 0; i < buffer.size(); ++i) {
-            buffer[i] = generator ? local_dis(generator.value()) : local_dis(local_gen);
-        }
-    }
-#else
-    std::uniform_real_distribution<float> dis(a, b);
-    std::generate(buffer.begin(), buffer.end(), [&]() { return generator ? dis(generator.value()) : dis(gen); });
-#endif
-
+                                std::optional<Generator> generator) {
+    CheckDistributionTensor(*tensor);
+    CheckUniformParameters(*tensor, a, b);
     auto device = tensor->GetDevice();
-
     core::DeviceGuard guard(device);
-    auto impl = core::GetDeviceGuardImpl(device.type());
-
-    impl->MemcpyAsync(tensor->DataPtr(), buffer.data(), num_elements * sizeof(float),
-                      device.type() == Device::DeviceType::kCPU ? core::MemcpyKind::kD2D : core::MemcpyKind::kH2D,
-                      impl->GetStream(device));
-
+    Dispatcher::Instance().Call<void>({device.type(), "Uniform"}, tensor, static_cast<double>(a),
+                                      static_cast<double>(b), generator);
     return tensor;
 }
 
diff --git a/infini_train/src/nn/parallel/ddp/distributed_data_parallel.cc b/infini_train/src/nn/parallel/ddp/distributed_data_parallel.cc
index 19361e960..141f6851e 100644
--- a/infini_train/src/nn/parallel/ddp/distributed_data_parallel.cc
+++ b/infini_train/src/nn/parallel/ddp/distributed_data_parallel.cc
@@ -1,8 +1,10 @@
 #include "infini_train/include/nn/parallel/ddp/distributed_data_parallel.h"
 
+#include <algorithm>
 #include <functional>
 #include <map>
 #include <memory>
+#include <string>
 #include <utility>
 #include <vector>
 
@@ -53,6 +55,7 @@ DistributedDataParallel::DistributedDataParallel(std::shared_ptr<nn::Module> mod
     }
     for (auto &buffer : module->Buffers()) { validate_device(buffer->GetDevice(), "buffer"); }
     modules_[kModuleName] = std::move(module);
+    SynchronizeModuleState();
 
     if (ddp_config.zero_stage >= 1) {
         BuildParamAndGradBuffers();
@@ -70,6 +73,30 @@ DistributedDataParallel::DistributedDataParallel(std::shared_ptr<nn::Module> mod
     }
 }
 
+void DistributedDataParallel::SynchronizeModuleState() {
+    auto state = modules_.at(kModuleName)->StateDict();
+    std::vector<std::string> names;
+    names.reserve(state.size());
+    for (const auto &[name, _] : state) { names.push_back(name); }
+    std::sort(names.begin(), names.end());
+
+    std::vector<std::shared_ptr<Tensor>> tensors;
+    tensors.reserve(names.size());
+    for (const auto &name : names) {
+        auto tensor = state.at(name);
+        CHECK(tensor != nullptr) << "DDP module state tensor '" << name << "' is null";
+        if (!tensors.empty()) {
+            CHECK_EQ(tensor->GetDevice(), tensors.front()->GetDevice())
+                << "DDP module state tensors must be on one device";
+        }
+        tensors.push_back(std::move(tensor));
+    }
+
+    if (!tensors.empty()) {
+        ddp_pg_->Broadcast(tensors, 0);
+    }
+}
+
 void DistributedDataParallel::BuildParamAndGradBuffers() {
     // (param_dtype, grad_dtype)
     using DTypePair = std::pair<DataType, DataType>;
diff --git a/infini_train/src/nn/parallel/ddp/param_and_grad_buffer.cc b/infini_train/src/nn/parallel/ddp/param_and_grad_buffer.cc
index ab3a80024..b5afd48ba 100644
--- a/infini_train/src/nn/parallel/ddp/param_and_grad_buffer.cc
+++ b/infini_train/src/nn/parallel/ddp/param_and_grad_buffer.cc
@@ -3,6 +3,7 @@
 #include <algorithm>
 #include <cstddef>
 #include <iostream>
+#include <numeric>
 
 #include "glog/logging.h"
 
diff --git a/infini_train/src/tensor.cc b/infini_train/src/tensor.cc
index 4e61e221b..fe958f2ba 100644
--- a/infini_train/src/tensor.cc
+++ b/infini_train/src/tensor.cc
@@ -3,6 +3,7 @@
 #include <cstring>
 #include <fstream>
 #include <functional>
+#include <iomanip>
 #include <memory>
 #include <numeric>
 
@@ -473,7 +474,7 @@ std::shared_ptr<Tensor> Tensor::Outer(const std::shared_ptr<Tensor> &other) {
 }
 
 // distribution
-std::shared_ptr<Tensor> Tensor::Uniform(float from, float to, std::optional<std::mt19937> generator) {
+std::shared_ptr<Tensor> Tensor::Uniform(float from, float to, std::optional<Generator> generator) {
     return nn::init::Uniform(shared_from_this(), from, to, generator);
 }
 
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 43bef143c..efb6962b1 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -36,5 +36,8 @@ add_subdirectory(dtype)
 # Transformer architecture tests
 add_subdirectory(transformer)
 
+# Generator tests
+add_subdirectory(generator)
+
 # Checkpoint tests
 add_subdirectory(checkpoint)
diff --git a/tests/common/test_utils.h b/tests/common/test_utils.h
index 0774622a0..2095ba8fa 100644
--- a/tests/common/test_utils.h
+++ b/tests/common/test_utils.h
@@ -1,5 +1,6 @@
 #pragma once
 
+#include <cstdint>
 #include <memory>
 #include <vector>
 
@@ -7,6 +8,7 @@
 #include <cuda_runtime_api.h>
 #endif
 
+#include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/device.h"
 #include "infini_train/include/tensor.h"
 #include "gtest/gtest.h"
@@ -105,6 +107,16 @@ inline void ExpectTensorNear(const std::shared_ptr<Tensor> &val1, float val2, fl
         }                                                                                                              \
     } while (0)
 
+// Returns tensor bytes on the host for exact comparisons.
+// Synchronizes the source device before reading the copied data.
+inline std::vector<std::uint8_t> TensorBytes(const std::shared_ptr<Tensor> &tensor) {
+    const Tensor host = tensor->To(Device());
+    const Device source_device = tensor->GetDevice();
+    core::GetDeviceGuardImpl(source_device.type())->SynchronizeDevice(source_device);
+    const auto *data = static_cast<const std::uint8_t *>(host.DataPtr());
+    return {data, data + host.SizeInBytes()};
+}
+
 class InfiniTrainTest : public ::testing::TestWithParam<Device::DeviceType> {
 protected:
     Device GetDevice() const { return Device(GetParam(), 0); }
diff --git a/tests/generator/CMakeLists.txt b/tests/generator/CMakeLists.txt
new file mode 100644
index 000000000..ce48367fc
--- /dev/null
+++ b/tests/generator/CMakeLists.txt
@@ -0,0 +1,39 @@
+# Shared CPU/CUDA tests; DDP is built separately below.
+file(GLOB GENERATOR_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/test_*.cc)
+set(GENERATOR_DDP_SOURCE ${CMAKE_CURRENT_SOURCE_DIR}/test_generator_ddp.cc)
+list(REMOVE_ITEM GENERATOR_SOURCES ${GENERATOR_DDP_SOURCE})
+
+infini_train_add_test_suite(test_generator
+  SOURCES ${GENERATOR_SOURCES}
+  TEST_TIMEOUT 40
+)
+
+file(GLOB GENERATOR_CPU_ONLY_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/cpu_only/test_*.cc)
+if(GENERATOR_CPU_ONLY_SOURCES)
+  infini_train_add_test(test_generator_cpu_only
+    SOURCES ${GENERATOR_CPU_ONLY_SOURCES}
+    LABELS cpu
+    TEST_TIMEOUT 40
+  )
+endif()
+
+if(USE_CUDA)
+  file(GLOB GENERATOR_CUDA_ONLY_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/cuda_only/test_*.cc)
+  if(GENERATOR_CUDA_ONLY_SOURCES)
+    infini_train_add_test(test_generator_cuda_only
+      SOURCES ${GENERATOR_CUDA_ONLY_SOURCES}
+      LABELS cuda
+      TEST_TIMEOUT 40
+    )
+  endif()
+
+  # DDP uses a separate executable with its own main.
+  add_executable(test_generator_ddp ${GENERATOR_DDP_SOURCE})
+  target_compile_options(test_generator_ddp PRIVATE -Wno-error)
+  target_link_libraries(test_generator_ddp PRIVATE GTest::gtest)
+  target_include_directories(test_generator_ddp PRIVATE ${glog_SOURCE_DIR}/src)
+  link_infini_train_exe(test_generator_ddp)
+
+  add_test(NAME test_generator_ddp COMMAND $<TARGET_FILE:test_generator_ddp>)
+  set_tests_properties(test_generator_ddp PROPERTIES LABELS "cuda;distributed" TIMEOUT 40)
+endif()
diff --git a/tests/generator/cpu_only/test_generator_cpu.cc b/tests/generator/cpu_only/test_generator_cpu.cc
new file mode 100644
index 000000000..5a080ee41
--- /dev/null
+++ b/tests/generator/cpu_only/test_generator_cpu.cc
@@ -0,0 +1,56 @@
+#include <cstdint>
+#include <memory>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/nn/functional.h"
+
+#include "tests/common/test_utils.h"
+
+using namespace infini_train;
+using infini_train::test::TensorBytes;
+
+class GeneratorCpuTest : public ::testing::Test {};
+
+TEST_F(GeneratorCpuTest, HighBitsOfCPUSeedAffectTheSequence) {
+    const Device cpu;
+    constexpr uint64_t low_seed = 17;
+    constexpr uint64_t high_seed = low_seed + (uint64_t{1} << 32);
+    EXPECT_NE(TensorBytes(nn::function::Rand({256}, DataType::kFLOAT32, cpu, CreateGenerator(cpu, low_seed))),
+              TensorBytes(nn::function::Rand({256}, DataType::kFLOAT32, cpu, CreateGenerator(cpu, high_seed))));
+}
+
+TEST_F(GeneratorCpuTest, StateRestoresCachedNormalSequence) {
+    const Device cpu;
+    auto generator = CreateGenerator(cpu, 4567);
+    nn::function::Randn({3}, DataType::kFLOAT32, cpu, generator);
+    nn::function::Randn({3}, DataType::kFLOAT64, cpu, generator);
+    const auto state_after_odd_normal = generator.get_state();
+    const auto normal_float = TensorBytes(nn::function::Randn({127}, DataType::kFLOAT32, cpu, generator));
+    const auto normal_double = TensorBytes(nn::function::Randn({127}, DataType::kFLOAT64, cpu, generator));
+    const auto state_after_replay = TensorBytes(generator.get_state());
+
+    generator.set_state(*state_after_odd_normal);
+    EXPECT_EQ(normal_float, TensorBytes(nn::function::Randn({127}, DataType::kFLOAT32, cpu, generator)));
+    EXPECT_EQ(normal_double, TensorBytes(nn::function::Randn({127}, DataType::kFLOAT64, cpu, generator)));
+    EXPECT_EQ(state_after_replay, TensorBytes(generator.get_state()));
+}
+
+TEST_F(GeneratorCpuTest, ReseedingClearsBothNormalCaches) {
+    const Device cpu;
+    auto generator = CreateGenerator(cpu, 4567);
+    nn::function::Randn({3}, DataType::kFLOAT32, cpu, generator);
+    nn::function::Randn({3}, DataType::kFLOAT64, cpu, generator);
+
+    generator.set_current_seed(5678);
+    auto fresh = CreateGenerator(cpu, 5678);
+    EXPECT_EQ(TensorBytes(generator.get_state()), TensorBytes(fresh.get_state()));
+    EXPECT_EQ(TensorBytes(nn::function::Randn({127}, DataType::kFLOAT32, cpu, generator)),
+              TensorBytes(nn::function::Randn({127}, DataType::kFLOAT32, cpu, fresh)));
+    EXPECT_EQ(TensorBytes(nn::function::Randn({127}, DataType::kFLOAT64, cpu, generator)),
+              TensorBytes(nn::function::Randn({127}, DataType::kFLOAT64, cpu, fresh)));
+    EXPECT_EQ(TensorBytes(generator.get_state()), TensorBytes(fresh.get_state()));
+}
diff --git a/tests/generator/cuda_only/test_generator_cuda.cc b/tests/generator/cuda_only/test_generator_cuda.cc
new file mode 100644
index 000000000..19672ef4a
--- /dev/null
+++ b/tests/generator/cuda_only/test_generator_cuda.cc
@@ -0,0 +1,110 @@
+#include <cstdint>
+#include <memory>
+#include <stdexcept>
+#include <vector>
+
+#include <cuda_runtime_api.h>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/nn/functional.h"
+#include "infini_train/include/nn/init.h"
+#include "infini_train/include/tensor.h"
+
+#include "tests/common/test_utils.h"
+
+using namespace infini_train;
+using infini_train::test::TensorBytes;
+
+class GeneratorCudaTest : public ::testing::Test {
+protected:
+    const Device cuda0_{Device::DeviceType::kCUDA, 0};
+};
+
+TEST_F(GeneratorCudaTest, StateRejectsCudaAndForeignInput) {
+    ::testing::GTEST_FLAG(death_test_style) = "threadsafe";
+    auto generator = CreateGenerator(cuda0_, 4601);
+    const auto valid_state = generator.get_state();
+
+    auto wrong_device = std::make_shared<Tensor>(std::vector<int64_t>{1}, DataType::kUINT8, cuda0_);
+    EXPECT_DEATH(generator.set_state(*wrong_device), "CPU");
+
+    const auto cpu_state = CreateGenerator(Device(), 4601).get_state();
+    EXPECT_DEATH(generator.set_state(*cpu_state), "Check failed");
+
+    auto cpu_generator = CreateGenerator(Device(), 4601);
+    EXPECT_DEATH(cpu_generator.set_state(*valid_state), "Check failed");
+}
+
+TEST_F(GeneratorCudaTest, CrossBackendGeneratorsAreRejected) {
+    const Device cpu;
+    const auto cpu_generator = CreateGenerator(cpu, 8001);
+    const auto cuda_generator = CreateGenerator(cuda0_, 8002);
+
+    auto cpu_input = std::make_shared<Tensor>(std::vector<int64_t>{8}, DataType::kFLOAT32, cpu);
+    auto cuda_input = std::make_shared<Tensor>(std::vector<int64_t>{8}, DataType::kFLOAT32, cuda0_);
+    cpu_input->Fill(1.0f);
+    cuda_input->Fill(1.0f);
+
+    EXPECT_THROW(nn::function::Rand({8}, DataType::kFLOAT32, cpu, cuda_generator), std::invalid_argument);
+    EXPECT_THROW(nn::function::Rand({8}, DataType::kFLOAT32, cuda0_, cpu_generator), std::invalid_argument);
+    EXPECT_THROW(nn::function::Dropout(cpu_input, 0.25, true, cuda_generator), std::invalid_argument);
+    EXPECT_THROW(nn::function::Dropout(cuda_input, 0.25, true, cpu_generator), std::invalid_argument);
+}
+
+// Regression for the device-index check: a generator bound to cuda:0 must not
+// drive a cuda:1 tensor, and the rejected calls must not advance its state.
+TEST_F(GeneratorCudaTest, CUDAGeneratorRejectsAnotherCUDADevice) {
+    REQUIRE_MIN_DEVICES(2);
+    const Device cuda1(Device::DeviceType::kCUDA, 1);
+    auto generator = CreateGenerator(cuda0_, 8401);
+    const auto state_before = TensorBytes(generator.get_state());
+
+    auto input = std::make_shared<Tensor>(std::vector<int64_t>{257}, DataType::kFLOAT32, cuda1);
+    input->Fill(1.0f);
+
+    EXPECT_THROW(nn::function::Rand({257}, DataType::kFLOAT32, cuda1, generator), std::invalid_argument);
+    EXPECT_THROW(nn::function::Randn({257}, DataType::kFLOAT32, cuda1, generator), std::invalid_argument);
+    EXPECT_THROW(nn::function::Dropout(input, 0.25, true, generator), std::invalid_argument);
+
+    EXPECT_EQ(state_before, TensorBytes(generator.get_state()));
+}
+
+TEST_F(GeneratorCudaTest, DefaultCUDAGeneratorsAreIndependentAcrossAllDevices) {
+    REQUIRE_MIN_DEVICES(2);
+    int device_count = 0;
+    ASSERT_EQ(cudaGetDeviceCount(&device_count), cudaSuccess);
+
+    ManualSeed(8501);
+    std::vector<Device> devices;
+    devices.reserve(device_count);
+    for (int index = 0; index < device_count; ++index) {
+        const Device device(Device::DeviceType::kCUDA, index);
+        devices.push_back(device);
+        EXPECT_EQ(GetDefaultGenerator(device).device(), device);
+        EXPECT_EQ(GetDefaultGenerator(device).current_seed(), 8501U);
+    }
+
+    for (int active_index = 0; active_index < device_count; ++active_index) {
+        std::vector<std::vector<uint8_t>> states_before;
+        states_before.reserve(device_count);
+        for (const auto &device : devices) {
+            states_before.push_back(TensorBytes(GetDefaultGenerator(device).get_state()));
+        }
+
+        auto tensor = std::make_shared<Tensor>(std::vector<int64_t>{64}, DataType::kFLOAT32, devices[active_index]);
+        nn::init::Uniform(tensor);
+
+        for (int observed_index = 0; observed_index < device_count; ++observed_index) {
+            const auto state_after = TensorBytes(GetDefaultGenerator(devices[observed_index]).get_state());
+            if (observed_index == active_index) {
+                EXPECT_NE(states_before[observed_index], state_after);
+            } else {
+                EXPECT_EQ(states_before[observed_index], state_after);
+            }
+        }
+    }
+}
diff --git a/tests/generator/test_generator_core.cc b/tests/generator/test_generator_core.cc
new file mode 100644
index 000000000..87b0a4862
--- /dev/null
+++ b/tests/generator/test_generator_core.cc
@@ -0,0 +1,127 @@
+#include <cstdint>
+#include <memory>
+#include <optional>
+#include <vector>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/nn/functional.h"
+#include "infini_train/include/nn/init.h"
+#include "infini_train/include/tensor.h"
+
+#include "infini_train/include/device.h"
+#include "tests/common/test_utils.h"
+
+using namespace infini_train;
+using infini_train::test::TensorBytes;
+
+class GeneratorCoreTest : public infini_train::test::InfiniTrainTest {};
+
+TEST_P(GeneratorCoreTest, PublicInterfaceCopyCloneAndSeed) {
+    const Device device = GetDevice();
+    auto generator = CreateGenerator(device);
+    EXPECT_EQ(generator.device(), device);
+    EXPECT_EQ(generator.current_seed(), Generator::kDefaultSeed);
+
+    generator.set_current_seed(1234);
+    EXPECT_EQ(generator.current_seed(), 1234U);
+    const auto first = TensorBytes(nn::function::Rand({128}, DataType::kFLOAT32, device, generator));
+    generator.set_current_seed(1234);
+    EXPECT_EQ(first, TensorBytes(nn::function::Rand({128}, DataType::kFLOAT32, device, generator)));
+
+    generator.set_current_seed(2345);
+    Generator alias = generator;
+    EXPECT_EQ(alias, generator);
+    const auto state_before = TensorBytes(alias.get_state());
+    nn::function::Rand({32}, DataType::kFLOAT32, device, generator);
+    EXPECT_NE(state_before, TensorBytes(alias.get_state()));
+
+    auto clone = generator.Clone();
+    EXPECT_NE(clone, generator);
+    EXPECT_EQ(TensorBytes(clone.get_state()), TensorBytes(generator.get_state()));
+    clone.set_current_seed(3456);
+    EXPECT_EQ(generator.current_seed(), 2345U);
+    EXPECT_EQ(clone.current_seed(), 3456U);
+
+    const uint64_t generated_seed = clone.Seed();
+    EXPECT_EQ(clone.current_seed(), generated_seed);
+}
+
+TEST_P(GeneratorCoreTest, StateRestoresUniformSequence) {
+    const Device device = GetDevice();
+    auto generator = CreateGenerator(device, 4567);
+    const auto state = generator.get_state();
+    EXPECT_EQ(state->Dtype(), DataType::kUINT8);
+    EXPECT_EQ(state->GetDevice(), Device());
+
+    const auto uniform = TensorBytes(nn::function::Rand({129}, DataType::kFLOAT32, device, generator));
+    generator.set_state(*state);
+    EXPECT_EQ(uniform, TensorBytes(nn::function::Rand({129}, DataType::kFLOAT32, device, generator)));
+}
+
+TEST_P(GeneratorCoreTest, StateRejectsMalformedInput) {
+    ::testing::GTEST_FLAG(death_test_style) = "threadsafe";
+    const Device device = GetDevice();
+    auto generator = CreateGenerator(device, 4601);
+    auto wrong_dtype = std::make_shared<Tensor>(std::vector<int64_t>{1}, DataType::kFLOAT32, Device());
+    EXPECT_DEATH(generator.set_state(*wrong_dtype), "UINT8");
+
+    const auto valid_state = generator.get_state();
+    auto truncated = std::make_shared<Tensor>(
+        std::vector<int64_t>{static_cast<int64_t>(valid_state->SizeInBytes() - 1)}, DataType::kUINT8, Device());
+    EXPECT_DEATH(generator.set_state(*truncated), "Check failed");
+}
+
+TEST_P(GeneratorCoreTest, DefaultExplicitAndUndefinedPathsHaveExpectedState) {
+    const Device device = GetDevice();
+    constexpr uint64_t seed = 7890;
+
+    ManualSeed(seed);
+    const auto *first_default = &GetDefaultGenerator(device);
+    const auto *second_default = &GetDefaultGenerator(device);
+    EXPECT_EQ(first_default, second_default);
+    EXPECT_EQ(first_default->device(), device);
+
+    const auto default_state = TensorBytes(first_default->get_state());
+    const auto explicit_result
+        = TensorBytes(nn::function::Rand({256}, DataType::kFLOAT32, device, CreateGenerator(device, seed)));
+    EXPECT_EQ(default_state, TensorBytes(first_default->get_state()));
+    EXPECT_EQ(explicit_result, TensorBytes(nn::function::Rand({256}, DataType::kFLOAT32, device)));
+
+    ManualSeed(seed + 1);
+    const auto default_result = TensorBytes(nn::function::Rand({256}, DataType::kFLOAT32, device));
+    ManualSeed(seed + 1);
+    std::optional<Generator> undefined = std::nullopt;
+    EXPECT_EQ(default_result, TensorBytes(nn::function::Rand({256}, DataType::kFLOAT32, device, undefined)));
+}
+
+TEST_P(GeneratorCoreTest, InitializerEntryPointsUseExplicitGenerator) {
+    const Device device = GetDevice();
+    auto make_tensor = [&] {
+        return std::make_shared<Tensor>(std::vector<int64_t>{16, 16}, DataType::kFLOAT32, device);
+    };
+
+    auto normal_a = make_tensor();
+    auto normal_b = make_tensor();
+    nn::init::Normal(normal_a, 1.5f, 0.5f, CreateGenerator(device, 8101));
+    nn::init::Normal(normal_b, 1.5f, 0.5f, CreateGenerator(device, 8101));
+    EXPECT_EQ(TensorBytes(normal_a), TensorBytes(normal_b));
+
+    auto uniform_a = make_tensor();
+    auto uniform_b = make_tensor();
+    nn::init::Uniform(uniform_a, -2.0f, 3.0f, CreateGenerator(device, 8102));
+    uniform_b->Uniform(-2.0f, 3.0f, CreateGenerator(device, 8102));
+    EXPECT_EQ(TensorBytes(uniform_a), TensorBytes(uniform_b));
+
+    auto kaiming_a = make_tensor();
+    auto kaiming_b = make_tensor();
+    nn::init::KaimingUniform(kaiming_a, 0.1f, nn::init::KaimingMode::kFanIn, nn::init::NonLinearityType::kLeakyReLU,
+                             CreateGenerator(device, 8103));
+    nn::init::KaimingUniform(kaiming_b, 0.1f, nn::init::KaimingMode::kFanIn, nn::init::NonLinearityType::kLeakyReLU,
+                             CreateGenerator(device, 8103));
+    EXPECT_EQ(TensorBytes(kaiming_a), TensorBytes(kaiming_b));
+}
+
+INFINI_TRAIN_REGISTER_TEST(GeneratorCoreTest);
diff --git a/tests/generator/test_generator_ddp.cc b/tests/generator/test_generator_ddp.cc
new file mode 100644
index 000000000..51ee2a4b5
--- /dev/null
+++ b/tests/generator/test_generator_ddp.cc
@@ -0,0 +1,91 @@
+#include <cstddef>
+#include <memory>
+#include <thread>
+#include <vector>
+
+#include "gtest/gtest.h"
+#include <cuda_runtime_api.h>
+
+#include "infini_train/include/core/runtime/device_guard.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/nn/modules/linear.h"
+#include "infini_train/include/nn/parallel/ddp/distributed_data_parallel.h"
+#include "infini_train/include/nn/parallel/ddp/distributed_data_parallel_config.h"
+#include "infini_train/include/nn/parallel/global.h"
+#include "infini_train/include/nn/parallel/process_group.h"
+#include "infini_train/include/nn/parallel/rank.h"
+#include "infini_train/include/nn/parallel/utils.h"
+#include "infini_train/include/tensor.h"
+
+using namespace infini_train;
+
+namespace {
+
+std::vector<float> CopyValuesToCPU(const std::shared_ptr<Tensor> &tensor) {
+    Tensor host(tensor->Dims(), tensor->Dtype(), Device());
+    host.CopyFrom(*tensor);
+    core::GetDeviceGuardImpl(tensor->GetDevice().type())->SynchronizeDevice(tensor->GetDevice());
+
+    const auto *values = static_cast<const float *>(host.DataPtr());
+    return {values, values + host.NumElements()};
+}
+
+TEST(DistributedDataParallelTest, SynchronizesParametersFromRankZeroAcrossAvailableGPUs) {
+    int device_count = 0;
+    ASSERT_EQ(cudaGetDeviceCount(&device_count), cudaSuccess);
+    if (device_count < 2) {
+        GTEST_SKIP() << "requires at least 2 GPUs (found " << device_count << ")";
+    }
+
+    const int world_size = nn::parallel::global::GetDataParallelSize();
+    ASSERT_EQ(world_size, device_count);
+
+    auto *factory = nn::parallel::ProcessGroupFactory::Instance(Device::DeviceType::kCUDA);
+    factory->GetOrCreate(nn::parallel::GetDataParallelProcessGroupName(0), nn::parallel::GetDataParallelGroupRanks(0));
+
+    std::vector<std::vector<float>> synchronized_values(static_cast<size_t>(world_size));
+    std::vector<std::thread> workers;
+    workers.reserve(static_cast<size_t>(world_size));
+    for (int thread_rank = 0; thread_rank < world_size; ++thread_rank) {
+        workers.emplace_back([&, thread_rank] {
+            nn::parallel::global::thread_global_rank = thread_rank;
+            const Device device(Device::DeviceType::kCUDA, thread_rank);
+            core::DeviceGuard guard(device);
+
+            auto module = std::make_shared<nn::Linear>(8, 4);
+            module->To(device);
+            for (const auto &parameter : module->Parameters()) { parameter->Fill(static_cast<float>(thread_rank + 1)); }
+
+            const nn::parallel::Rank rank(0, thread_rank, 1, world_size);
+            auto ddp = std::make_shared<nn::parallel::DistributedDataParallel>(
+                std::move(module), rank, nn::parallel::DistributedDataParallelConfig{});
+
+            for (const auto &parameter : ddp->module()->Parameters()) {
+                auto values = CopyValuesToCPU(parameter);
+                synchronized_values[static_cast<size_t>(thread_rank)].insert(
+                    synchronized_values[static_cast<size_t>(thread_rank)].end(), values.begin(), values.end());
+            }
+        });
+    }
+    for (auto &worker : workers) { worker.join(); }
+
+    ASSERT_FALSE(synchronized_values.front().empty());
+    for (const auto &rank_values : synchronized_values) {
+        EXPECT_EQ(rank_values, synchronized_values.front());
+        for (float value : rank_values) { EXPECT_FLOAT_EQ(value, 1.0f); }
+    }
+}
+
+} // namespace
+
+int main(int argc, char **argv) {
+    ::testing::InitGoogleTest(&argc, argv);
+
+    int device_count = 0;
+    if (cudaGetDeviceCount(&device_count) != cudaSuccess || device_count < 1) {
+        device_count = 1;
+    }
+    nn::parallel::global::GlobalEnv::Instance().Init(device_count, 1, false, 1, 1);
+
+    return RUN_ALL_TESTS();
+}
diff --git a/tests/generator/test_generator_random_ops.cc b/tests/generator/test_generator_random_ops.cc
new file mode 100644
index 000000000..9cff02523
--- /dev/null
+++ b/tests/generator/test_generator_random_ops.cc
@@ -0,0 +1,395 @@
+#include <cmath>
+#include <cstddef>
+#include <cstdint>
+#include <initializer_list>
+#include <memory>
+#include <optional>
+#include <vector>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/autograd/dropout.h"
+#include "infini_train/include/core/runtime/device_guard.h"
+#include "infini_train/include/datatype.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/nn/functional.h"
+#include "infini_train/include/nn/init.h"
+#include "infini_train/include/tensor.h"
+
+#include "infini_train/include/device.h"
+#include "tests/common/test_utils.h"
+
+using namespace infini_train;
+using infini_train::test::TensorBytes;
+
+namespace {
+
+template <typename T> void ExpectUniformRangeAndFinite(const Tensor &tensor) {
+    const auto *data = static_cast<const T *>(tensor.DataPtr());
+    for (int64_t index = 0; index < tensor.NumElements(); ++index) {
+        const double value = static_cast<double>(data[index]);
+        EXPECT_TRUE(std::isfinite(value));
+        EXPECT_GE(value, 0.0);
+        EXPECT_LT(value, 1.0);
+    }
+}
+
+template <typename T> void ExpectFinite(const Tensor &tensor) {
+    const auto *data = static_cast<const T *>(tensor.DataPtr());
+    for (int64_t index = 0; index < tensor.NumElements(); ++index) {
+        EXPECT_TRUE(std::isfinite(static_cast<double>(data[index])));
+    }
+}
+
+void CheckUniformRangeAndFinite(const std::shared_ptr<Tensor> &tensor) {
+    const Tensor host = tensor->To(Device());
+    const Device source_device = tensor->GetDevice();
+    core::GetDeviceGuardImpl(source_device.type())->SynchronizeDevice(source_device);
+    switch (host.Dtype()) {
+    case DataType::kFLOAT16:
+        ExpectUniformRangeAndFinite<FP16>(host);
+        break;
+    case DataType::kBFLOAT16:
+        ExpectUniformRangeAndFinite<BF16>(host);
+        break;
+    case DataType::kFLOAT32:
+        ExpectUniformRangeAndFinite<float>(host);
+        break;
+    case DataType::kFLOAT64:
+        ExpectUniformRangeAndFinite<double>(host);
+        break;
+    default:
+        FAIL() << "Unexpected dtype";
+    }
+}
+
+void CheckFinite(const std::shared_ptr<Tensor> &tensor) {
+    const Tensor host = tensor->To(Device());
+    const Device source_device = tensor->GetDevice();
+    core::GetDeviceGuardImpl(source_device.type())->SynchronizeDevice(source_device);
+    switch (host.Dtype()) {
+    case DataType::kFLOAT16:
+        ExpectFinite<FP16>(host);
+        break;
+    case DataType::kBFLOAT16:
+        ExpectFinite<BF16>(host);
+        break;
+    case DataType::kFLOAT32:
+        ExpectFinite<float>(host);
+        break;
+    case DataType::kFLOAT64:
+        ExpectFinite<double>(host);
+        break;
+    default:
+        FAIL() << "Unexpected dtype";
+    }
+}
+
+std::vector<std::vector<uint8_t>> RunRandomScript(const Device &device, std::optional<Generator> generator) {
+    auto input = std::make_shared<Tensor>(std::vector<int64_t>{257}, DataType::kFLOAT32, device);
+    input->Fill(1.0f);
+    return {
+        TensorBytes(nn::function::Rand({257}, DataType::kFLOAT32, device, generator)),
+        TensorBytes(nn::function::Randn({257}, DataType::kFLOAT32, device, generator)),
+        TensorBytes(nn::function::Dropout(input, 0.25, true, generator)),
+        TensorBytes(nn::function::Rand({257}, DataType::kFLOAT32, device, generator)),
+    };
+}
+
+void ExpectMaskValues(const std::shared_ptr<Tensor> &mask) {
+    EXPECT_EQ(mask->Dtype(), DataType::kUINT8);
+    const Tensor host = mask->To(Device());
+    const Device source_device = mask->GetDevice();
+    core::GetDeviceGuardImpl(source_device.type())->SynchronizeDevice(source_device);
+    const auto *data = static_cast<const uint8_t *>(host.DataPtr());
+    bool saw_zero = false;
+    bool saw_one = false;
+    for (int64_t index = 0; index < host.NumElements(); ++index) {
+        EXPECT_TRUE(data[index] == 0 || data[index] == 1);
+        saw_zero |= data[index] == 0;
+        saw_one |= data[index] == 1;
+    }
+    EXPECT_TRUE(saw_zero);
+    EXPECT_TRUE(saw_one);
+}
+
+template <typename T> void ExpectUnitInputDropoutValues(const Tensor &output, const Tensor &mask, double p) {
+    const auto *output_data = static_cast<const T *>(output.DataPtr());
+    const auto *mask_data = static_cast<const uint8_t *>(mask.DataPtr());
+    const double kept_value = 1.0 / (1.0 - p);
+    for (int64_t index = 0; index < output.NumElements(); ++index) {
+        const double expected = mask_data[index] ? kept_value : 0.0;
+        EXPECT_NEAR(static_cast<double>(output_data[index]), expected, 0.02);
+    }
+}
+
+void ExpectUnitInputDropoutValues(const std::shared_ptr<Tensor> &output, const std::shared_ptr<Tensor> &mask,
+                                  double p) {
+    const Tensor host_output = output->To(Device());
+    const Tensor host_mask = mask->To(Device());
+    // Each async D2H copy must complete before its host buffer is read.
+    const Device output_device = output->GetDevice();
+    core::GetDeviceGuardImpl(output_device.type())->SynchronizeDevice(output_device);
+    const Device mask_source_device = mask->GetDevice();
+    core::GetDeviceGuardImpl(mask_source_device.type())->SynchronizeDevice(mask_source_device);
+    switch (host_output.Dtype()) {
+    case DataType::kFLOAT16:
+        ExpectUnitInputDropoutValues<FP16>(host_output, host_mask, p);
+        break;
+    case DataType::kBFLOAT16:
+        ExpectUnitInputDropoutValues<BF16>(host_output, host_mask, p);
+        break;
+    case DataType::kFLOAT32:
+        ExpectUnitInputDropoutValues<float>(host_output, host_mask, p);
+        break;
+    case DataType::kFLOAT64:
+        ExpectUnitInputDropoutValues<double>(host_output, host_mask, p);
+        break;
+    default:
+        FAIL() << "Unexpected dtype";
+    }
+}
+
+} // namespace
+
+class GeneratorRandomOpsTest : public infini_train::test::InfiniTrainTest {};
+
+TEST_P(GeneratorRandomOpsTest, RandAndRandnSupportAllFloatingDtypes) {
+    const std::vector<DataType> dtypes
+        = {DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64};
+    const Device device = GetDevice();
+
+    for (const auto dtype : dtypes) {
+        auto uniform_first = CreateGenerator(device, 1001);
+        auto uniform_replay = CreateGenerator(device, 1001);
+        auto uniform = nn::function::Rand({4097}, dtype, device, uniform_first);
+        EXPECT_EQ(TensorBytes(uniform), TensorBytes(nn::function::Rand({4097}, dtype, device, uniform_replay)));
+        CheckUniformRangeAndFinite(uniform);
+
+        auto normal_first = CreateGenerator(device, 2002);
+        auto normal_replay = CreateGenerator(device, 2002);
+        auto normal = nn::function::Randn({4097}, dtype, device, normal_first);
+        EXPECT_EQ(TensorBytes(normal), TensorBytes(nn::function::Randn({4097}, dtype, device, normal_replay)));
+        CheckFinite(normal);
+    }
+}
+
+TEST_P(GeneratorRandomOpsTest, UniformConstantDoesNotAdvanceGenerator) {
+    const Device device = GetDevice();
+    for (const auto dtype : {DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64}) {
+        SCOPED_TRACE(static_cast<int>(dtype));
+        auto generator = CreateGenerator(device, 2101);
+        nn::function::Rand({11}, dtype, device, generator);
+        const auto state_before = TensorBytes(generator.get_state());
+        auto replay = generator.Clone();
+        auto tensor = std::make_shared<Tensor>(std::vector<int64_t>{257}, dtype, device);
+        auto expected = std::make_shared<Tensor>(tensor->Dims(), dtype, Device());
+        expected->Fill(0.1f);
+
+        nn::init::Uniform(tensor, 0.1f, 0.1f, generator);
+
+        EXPECT_EQ(TensorBytes(tensor), TensorBytes(expected));
+        EXPECT_EQ(TensorBytes(generator.get_state()), state_before);
+        EXPECT_EQ(TensorBytes(nn::function::Rand({129}, dtype, device, generator)),
+                  TensorBytes(nn::function::Rand({129}, dtype, device, replay)));
+    }
+}
+
+TEST_P(GeneratorRandomOpsTest, UniformEmptyTensorDoesNotAdvanceGenerator) {
+    const Device device = GetDevice();
+    for (const auto dtype : {DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64}) {
+        SCOPED_TRACE(static_cast<int>(dtype));
+        auto generator = CreateGenerator(device, 2102);
+        nn::function::Rand({11}, dtype, device, generator);
+        const auto state_before = TensorBytes(generator.get_state());
+        auto replay = generator.Clone();
+        auto tensor = std::make_shared<Tensor>(std::vector<int64_t>{2, 0, 3}, dtype, device);
+
+        nn::init::Uniform(tensor, -2.0f, 3.0f, generator);
+
+        EXPECT_EQ(tensor->NumElements(), 0);
+        EXPECT_EQ(TensorBytes(generator.get_state()), state_before);
+        EXPECT_EQ(TensorBytes(nn::function::Rand({129}, dtype, device, generator)),
+                  TensorBytes(nn::function::Rand({129}, dtype, device, replay)));
+    }
+}
+
+TEST_P(GeneratorRandomOpsTest, DropoutSupportsAllFloatingDtypesWithExplicitGenerator) {
+    const std::vector<DataType> dtypes
+        = {DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64};
+    const Device device = GetDevice();
+
+    for (const auto dtype : dtypes) {
+        auto input = nn::function::Rand({4097}, dtype, device, CreateGenerator(device, 3003));
+        ManualSeed(4004);
+        const auto default_state_before = TensorBytes(GetDefaultGenerator(device).get_state());
+
+        auto dropout_first = CreateGenerator(device, 5005);
+        auto dropout_replay = CreateGenerator(device, 5005);
+        auto output = nn::function::Dropout(input, 0.25, true, dropout_first);
+        EXPECT_EQ(TensorBytes(output), TensorBytes(nn::function::Dropout(input, 0.25, true, dropout_replay)));
+        CheckFinite(output);
+        EXPECT_EQ(default_state_before, TensorBytes(GetDefaultGenerator(device).get_state()));
+    }
+}
+
+TEST_P(GeneratorRandomOpsTest, SameSeedReplaysRandRandnAndDropoutScript) {
+    const Device device = GetDevice();
+    EXPECT_EQ(RunRandomScript(device, CreateGenerator(device, 5678)),
+              RunRandomScript(device, CreateGenerator(device, 5678)));
+    EXPECT_NE(RunRandomScript(device, CreateGenerator(device, 5678)),
+              RunRandomScript(device, CreateGenerator(device, 5679)));
+
+    auto generator = CreateGenerator(device, 6789);
+    const auto first = TensorBytes(nn::function::Rand({256}, DataType::kFLOAT32, device, generator));
+    const auto second = TensorBytes(nn::function::Rand({256}, DataType::kFLOAT32, device, generator));
+    EXPECT_NE(first, second);
+}
+
+TEST_P(GeneratorRandomOpsTest, DefaultGeneratorReplaysRandRandnAndDropoutScript) {
+    const Device device = GetDevice();
+
+    ManualSeed(7001);
+    const auto expected = RunRandomScript(device, std::nullopt);
+    ManualSeed(7001);
+    EXPECT_EQ(expected, RunRandomScript(device, std::nullopt));
+
+    ManualSeed(7001);
+    std::optional<Generator> undefined = std::nullopt;
+    EXPECT_EQ(expected, RunRandomScript(device, undefined));
+}
+
+TEST_P(GeneratorRandomOpsTest, ExplicitRandomGeneratorsDoNotAdvanceDefaultGenerator) {
+    const Device device = GetDevice();
+    ManualSeed(7101);
+    const auto default_state_before = TensorBytes(GetDefaultGenerator(device).get_state());
+
+    RunRandomScript(device, CreateGenerator(device, 7102));
+
+    EXPECT_EQ(default_state_before, TensorBytes(GetDefaultGenerator(device).get_state()));
+}
+
+TEST_P(GeneratorRandomOpsTest, OffsetViewsUseTheirOwnDataPointers) {
+    const Device device = GetDevice();
+    constexpr int64_t kStorageElements = 12;
+    constexpr int64_t kOffsetElements = 4;
+    constexpr int64_t kViewElements = 4;
+    constexpr size_t kOffsetBytes = kOffsetElements * sizeof(float);
+
+    auto storage = std::make_shared<Tensor>(std::vector<int64_t>{kStorageElements}, DataType::kFLOAT32, device);
+    storage->Fill(-1.0f);
+    auto uniform_view = std::make_shared<Tensor>(*storage, kOffsetBytes, std::vector<int64_t>{kViewElements});
+    nn::init::Uniform(uniform_view, 0.0f, 1.0f, CreateGenerator(device, 7201));
+
+    const Tensor uniform_storage = storage->To(Device());
+    const Device storage_device = storage->GetDevice();
+    core::GetDeviceGuardImpl(storage_device.type())->SynchronizeDevice(storage_device);
+    const auto *uniform_data = static_cast<const float *>(uniform_storage.DataPtr());
+    for (int64_t index = 0; index < kOffsetElements; ++index) { EXPECT_FLOAT_EQ(uniform_data[index], -1.0f); }
+    for (int64_t index = kOffsetElements; index < kOffsetElements + kViewElements; ++index) {
+        EXPECT_GE(uniform_data[index], 0.0f);
+        EXPECT_LT(uniform_data[index], 1.0f);
+    }
+    for (int64_t index = kOffsetElements + kViewElements; index < kStorageElements; ++index) {
+        EXPECT_FLOAT_EQ(uniform_data[index], -1.0f);
+    }
+
+    auto input_storage = std::make_shared<Tensor>(std::vector<int64_t>{kStorageElements}, DataType::kFLOAT32, device);
+    input_storage->Fill(-3.0f);
+    auto input_view = std::make_shared<Tensor>(*input_storage, kOffsetBytes, std::vector<int64_t>{kViewElements});
+    input_view->Fill(1.0f);
+    const auto output = nn::function::Dropout(input_view, 0.5, true, CreateGenerator(device, 7202));
+    const Tensor host_output = output->To(Device());
+    const Device output_device = output->GetDevice();
+    core::GetDeviceGuardImpl(output_device.type())->SynchronizeDevice(output_device);
+    const auto *output_data = static_cast<const float *>(host_output.DataPtr());
+    for (int64_t index = 0; index < kViewElements; ++index) {
+        EXPECT_TRUE(output_data[index] == 0.0f || std::abs(output_data[index] - 2.0f) < 1e-6f);
+    }
+
+    const Tensor input_storage_host = input_storage->To(Device());
+    const Device input_storage_device = input_storage->GetDevice();
+    core::GetDeviceGuardImpl(input_storage_device.type())->SynchronizeDevice(input_storage_device);
+    const auto *input_storage_data = static_cast<const float *>(input_storage_host.DataPtr());
+    for (int64_t index = 0; index < kOffsetElements; ++index) { EXPECT_FLOAT_EQ(input_storage_data[index], -3.0f); }
+    for (int64_t index = kOffsetElements + kViewElements; index < kStorageElements; ++index) {
+        EXPECT_FLOAT_EQ(input_storage_data[index], -3.0f);
+    }
+}
+
+TEST_P(GeneratorRandomOpsTest, DropoutBackwardSupportsAllFloatingDtypes) {
+    const Device device = GetDevice();
+    const std::vector<DataType> dtypes
+        = {DataType::kFLOAT16, DataType::kBFLOAT16, DataType::kFLOAT32, DataType::kFLOAT64};
+
+    for (const auto dtype : dtypes) {
+        auto input = std::make_shared<Tensor>(std::vector<int64_t>{4097}, dtype, device, true);
+        input->Fill(1.0f);
+        auto dropout = std::make_shared<autograd::Dropout>(0.25, CreateGenerator(device, 8201));
+        const auto outputs = dropout->Apply({input});
+        ASSERT_EQ(outputs.size(), 2U);
+        ExpectMaskValues(outputs[1]);
+        ExpectUnitInputDropoutValues(outputs[0], outputs[1], 0.25);
+
+        auto grad_output = std::make_shared<Tensor>(input->Dims(), dtype, device);
+        grad_output->Fill(1.0f);
+        const auto grad_inputs = dropout->Backward({grad_output, nullptr});
+        ASSERT_EQ(grad_inputs.size(), 1U);
+        EXPECT_EQ(TensorBytes(outputs[0]), TensorBytes(grad_inputs[0]));
+    }
+}
+
+TEST_P(GeneratorRandomOpsTest, DropoutIdentityBoundariesDoNotAdvanceGenerator) {
+    const Device device = GetDevice();
+    auto input = std::make_shared<Tensor>(std::vector<int64_t>{64}, DataType::kFLOAT32, device, true);
+    input->Fill(1.0f);
+    auto generator = CreateGenerator(device, 8301);
+
+    auto expect_unchanged = [&](const auto &operation) {
+        const auto before = TensorBytes(generator.get_state());
+        operation();
+        EXPECT_EQ(before, TensorBytes(generator.get_state()));
+    };
+
+    expect_unchanged([&] { EXPECT_EQ(nn::function::Dropout(input, 0.5, false, generator), input); });
+    expect_unchanged([&] { EXPECT_EQ(nn::function::Dropout(input, 0.0, true, generator), input); });
+}
+
+TEST_P(GeneratorRandomOpsTest, DropoutP1ProducesZeroForwardAndBackwardWithoutAdvancing) {
+    const Device device = GetDevice();
+    auto input = std::make_shared<Tensor>(std::vector<int64_t>{64}, DataType::kFLOAT32, device, true);
+    input->Fill(1.0f);
+    auto generator = CreateGenerator(device, 8301);
+    const auto before = TensorBytes(generator.get_state());
+
+    auto all_dropped = nn::function::Dropout(input, 1.0, true, generator);
+    EXPECT_EQ(before, TensorBytes(generator.get_state()));
+    auto zeros = std::make_shared<Tensor>(input->Dims(), input->Dtype(), device);
+    zeros->Fill(0.0f);
+    EXPECT_EQ(TensorBytes(all_dropped), TensorBytes(zeros));
+
+    auto dropout = std::make_shared<autograd::Dropout>(1.0, generator);
+    const auto outputs = dropout->Apply({input});
+    auto grad_output = std::make_shared<Tensor>(input->Dims(), input->Dtype(), device);
+    grad_output->Fill(1.0f);
+    const auto grad_inputs = dropout->Backward({grad_output, nullptr});
+    EXPECT_EQ(TensorBytes(grad_inputs[0]), TensorBytes(zeros));
+}
+
+TEST_P(GeneratorRandomOpsTest, InvalidDistributionParametersAreRejected) {
+    ::testing::GTEST_FLAG(death_test_style) = "threadsafe";
+    const Device device = GetDevice();
+    auto tensor = std::make_shared<Tensor>(std::vector<int64_t>{8}, DataType::kFLOAT32, device);
+    auto integer_tensor = std::make_shared<Tensor>(std::vector<int64_t>{8}, DataType::kINT32, device);
+    auto generator = CreateGenerator(device, 9001);
+
+    EXPECT_DEATH(nn::init::Uniform(tensor, 2.0f, 1.0f, generator), "uniform expects a");
+    EXPECT_DEATH(nn::init::Normal(tensor, 0.0f, -1.0f, generator), "std >= 0.0");
+    EXPECT_DEATH(nn::function::Dropout(tensor, -0.1, true, generator), "between 0 and 1");
+    EXPECT_DEATH(nn::function::Dropout(tensor, 1.1, true, generator), "between 0 and 1");
+    EXPECT_DEATH(nn::function::Rand({8}, DataType::kINT32, device, generator), "floating-point");
+    EXPECT_DEATH(nn::init::Uniform(integer_tensor, 0.0f, 1.0f, generator), "floating-point");
+    EXPECT_DEATH(nn::init::Normal(integer_tensor, 0.0f, 1.0f, generator), "floating-point");
+}
+
+INFINI_TRAIN_REGISTER_TEST(GeneratorRandomOpsTest);