diff --git a/infini_train/include/generator.h b/infini_train/include/generator.h
new file mode 100644
index 00000000..789d2358
--- /dev/null
+++ b/infini_train/include/generator.h
@@ -0,0 +1,46 @@
+#pragma once
+
+#include <cstdint>
+#include <memory>
+#include <vector>
+
+#include "infini_train/include/device.h"
+
+namespace infini_train {
+
+class GeneratorImpl;
+namespace detail {
+class GeneratorAccessor;
+}
+
+class Generator {
+public:
+    void ManualSeed(uint64_t seed);
+    uint64_t Seed();
+    uint64_t InitialSeed() const;
+    std::vector<uint8_t> GetState() const;
+    void SetState(const std::vector<uint8_t> &state);
+    // The CUDA index is the generator's home index. An explicit CUDA generator
+    // may be consumed by tensors on another CUDA index.
+    Device GetDevice() const;
+
+private:
+    explicit Generator(std::shared_ptr<GeneratorImpl> impl);
+
+    std::shared_ptr<GeneratorImpl> impl_;
+
+    friend class detail::GeneratorAccessor;
+    friend std::shared_ptr<Generator> MakeCPUGenerator(uint64_t seed);
+    friend std::shared_ptr<Generator> MakeCUDAGenerator(int8_t device_index, uint64_t seed);
+};
+
+std::shared_ptr<Generator> MakeCPUGenerator(uint64_t seed = 42);
+std::shared_ptr<Generator> MakeCUDAGenerator(int8_t device_index, uint64_t seed = 42);
+std::shared_ptr<Generator> GetDefaultCPUGenerator();
+std::shared_ptr<Generator> GetDefaultCUDAGenerator(int8_t device_index);
+std::shared_ptr<Generator> GetDefaultGenerator(const Device &device);
+// ManualSeed is an alias for ManualSeedAll.
+void ManualSeed(uint64_t seed);
+void ManualSeedAll(uint64_t seed);
+
+} // namespace infini_train
diff --git a/infini_train/include/nn/functional.h b/infini_train/include/nn/functional.h
index e4354fd1..dddb297c 100644
--- a/infini_train/include/nn/functional.h
+++ b/infini_train/include/nn/functional.h
@@ -4,9 +4,12 @@
 #include <memory>
 #include <vector>
 
+#include "infini_train/include/device.h"
+
 namespace infini_train {
+class Generator;
 class Tensor;
-}
+} // namespace infini_train
 
 namespace infini_train::nn::function {
 
@@ -47,6 +50,12 @@ 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);
 
+std::shared_ptr<Tensor> Rand(const std::vector<int64_t> &size, Device device = Device(),
+                             std::shared_ptr<Generator> generator = nullptr);
+
+std::shared_ptr<Tensor> Randn(const std::vector<int64_t> &size, Device device = Device(),
+                              std::shared_ptr<Generator> generator = nullptr);
+
 // 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 fc6effec..f655563d 100644
--- a/infini_train/include/nn/init.h
+++ b/infini_train/include/nn/init.h
@@ -1,21 +1,20 @@
 #pragma once
 
+#include <cstdint>
 #include <memory>
-#include <optional>
-#include <random>
 #include <utility>
 
 #include "infini_train/include/datatype.h"
 #include "infini_train/include/device.h"
 
 namespace infini_train {
+class Generator;
 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::shared_ptr<Tensor> Normal(const std::shared_ptr<Tensor> &tensor, float mean = 0.0f, float stddev = 1.0f,
+                               std::shared_ptr<Generator> generator = nullptr);
 
 std::pair<int64_t, int64_t> CalculateFanInAndFanOut(const std::shared_ptr<Tensor> &tensor);
 
@@ -42,10 +41,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::shared_ptr<Generator> generator = nullptr);
 
 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::shared_ptr<Generator> generator = nullptr);
 
 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 823ae82b..e9b84e63 100644
--- a/infini_train/include/nn/parallel/ddp/distributed_data_parallel.h
+++ b/infini_train/include/nn/parallel/ddp/distributed_data_parallel.h
@@ -38,6 +38,7 @@ class DistributedDataParallel : public nn::Module {
     const std::vector<std::shared_ptr<ParamAndGradBucketGroup>> &bucket_groups() const { return bucket_groups_; }
 
 private:
+    void SynchronizeModuleState();
     void BuildParamAndGradBuffers();
     void RegisterBackwardHooks();
     void OnGradReady(const std::shared_ptr<Tensor> &param);
diff --git a/infini_train/include/nn/parallel/process_group.h b/infini_train/include/nn/parallel/process_group.h
index 2c439f40..4db1b873 100644
--- a/infini_train/include/nn/parallel/process_group.h
+++ b/infini_train/include/nn/parallel/process_group.h
@@ -134,25 +134,26 @@ class ProcessGroupFactory {
     template <typename Creator, typename = std::enable_if_t<std::is_invocable_v<Creator>>>
     const ProcessGroup *GetOrCreate(const std::string &name, Creator &&creator) {
         std::unique_lock<std::mutex> lock(mutex_);
-        auto [it, inserted] = name_to_group_.emplace(name, nullptr);
+        const bool inserted = name_to_group_.emplace(name, nullptr).second;
         if (!inserted) {
-            while (it->second == nullptr) { cond_.wait(lock); }
-            return it->second.get();
+            cond_.wait(lock, [this, &name]() { return name_to_group_.at(name) != nullptr; });
+            return name_to_group_.at(name).get();
         }
 
         lock.unlock();
         auto new_group = creator();
         lock.lock();
 
-        it->second = std::move(new_group);
+        auto &group = name_to_group_.at(name);
+        group = std::move(new_group);
         cond_.notify_all();
-        return it->second.get();
+        return group.get();
     }
 
 private:
     // TODO(dcj): maybe RWLock later?
     mutable std::mutex mutex_;
-    std::condition_variable cond_;
+    mutable std::condition_variable cond_;
     std::unordered_map<std::string, std::unique_ptr<ProcessGroup>> name_to_group_;
     Device::DeviceType backend_ = Device::DeviceType::kInvalid;
 };
diff --git a/infini_train/include/tensor.h b/infini_train/include/tensor.h
index dcfd8927..fe066ed3 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"
@@ -15,6 +14,7 @@
 #include "infini_train/include/scalar.h"
 
 namespace infini_train {
+class Generator;
 namespace autograd {
 class Function;
 class AccumulateGrad;
@@ -150,8 +150,7 @@ class Tensor : public std::enable_shared_from_this<Tensor> {
     std::shared_ptr<Tensor> Unsqueeze(int64_t dim);
 
     // distribution
-    std::shared_ptr<Tensor> Uniform(float from = 0.0f, float to = 1.0f,
-                                    std::optional<std::mt19937> generator = std::nullopt);
+    std::shared_ptr<Tensor> Uniform(float from = 0.0f, float to = 1.0f, std::shared_ptr<Generator> generator = nullptr);
 
     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/generator.cc b/infini_train/src/generator.cc
new file mode 100644
index 00000000..3ffdcd41
--- /dev/null
+++ b/infini_train/src/generator.cc
@@ -0,0 +1,404 @@
+#include "infini_train/include/generator.h"
+#include "infini_train/src/generator_internal.h"
+#include "infini_train/src/random_utils.h"
+
+#include <bit>
+#include <charconv>
+#include <cmath>
+#include <limits>
+#include <locale>
+#include <mutex>
+#include <random>
+#include <sstream>
+#include <string>
+#include <unordered_map>
+#include <utility>
+
+#include "glog/logging.h"
+
+namespace infini_train {
+namespace {
+constexpr char kCPUStateHeader[] = "InfiniTrain CPU Generator";
+constexpr char kCUDAStateHeader[] = "InfiniTrain CUDA Generator";
+constexpr float kTwoPi = 6.283185307179586476925286766559f;
+
+// Keep the raw word-to-unit-interval mapping explicit: std::uniform_real_distribution does not specify an identical
+// value sequence across standard library implementations.
+float UintToUnitFloat(uint32_t value) { return static_cast<float>(value >> 8) * 0x1.0p-24f; }
+
+std::pair<float, float> BoxMullerPair(uint32_t first, uint32_t second) {
+    const float uniform1 = 1.0f - UintToUnitFloat(first);
+    const float uniform2 = UintToUnitFloat(second);
+    const float radius = std::sqrt(-2.0f * std::log(uniform1));
+    const float angle = kTwoPi * uniform2;
+    return {radius * std::cos(angle), radius * std::sin(angle)};
+}
+
+void SeedEngine(std::mt19937 &engine, uint64_t seed) {
+    std::seed_seq seed_seq{static_cast<uint32_t>(seed), static_cast<uint32_t>(seed >> 32)};
+    engine.seed(seed_seq);
+}
+
+uint64_t MakeRandomSeed() {
+    std::random_device rd;
+    return (static_cast<uint64_t>(rd()) << 32) ^ rd();
+}
+
+std::string ToString(const std::vector<uint8_t> &state) { return std::string(state.begin(), state.end()); }
+
+std::vector<uint8_t> ToBytes(const std::string &state) { return std::vector<uint8_t>(state.begin(), state.end()); }
+
+template <typename T> T ParseStateIntegerToken(const std::string &token, const char *error_message) {
+    T value = 0;
+    const auto [ptr, error] = std::from_chars(token.data(), token.data() + token.size(), value);
+    CHECK(error == std::errc{} && ptr == token.data() + token.size()) << error_message;
+    return value;
+}
+
+template <typename T> T ParseStateIntegerLine(std::istringstream &iss, const char *error_message) {
+    std::string line;
+    CHECK(std::getline(iss, line)) << error_message;
+    return ParseStateIntegerToken<T>(line, error_message);
+}
+
+void ValidateCPUEngineState(const std::string &serialized) {
+    std::istringstream iss(serialized);
+    iss.imbue(std::locale::classic());
+    bool any_nonzero_word = false;
+    for (size_t i = 0; i < std::mt19937::state_size; ++i) {
+        std::string token;
+        CHECK(iss >> token) << "Invalid CPU generator engine state: missing state word";
+        const uint64_t word = ParseStateIntegerToken<uint64_t>(token, "Invalid CPU generator engine word in state");
+        CHECK_LE(word, std::numeric_limits<uint32_t>::max()) << "Invalid CPU generator engine word in state";
+        any_nonzero_word = any_nonzero_word || word != 0;
+    }
+
+    std::string position_token;
+    CHECK(iss >> position_token) << "Invalid CPU generator engine state: missing position";
+    const uint64_t position
+        = ParseStateIntegerToken<uint64_t>(position_token, "Invalid CPU generator engine position in state");
+    CHECK_LE(position, std::mt19937::state_size) << "Invalid CPU generator engine position in state";
+
+    std::string trailing_token;
+    CHECK(!(iss >> trailing_token)) << "Invalid CPU generator state: trailing data";
+    CHECK(any_nonzero_word) << "Invalid CPU generator engine state: all-zero state";
+}
+
+void CheckNoTrailingStateData(std::istringstream &iss, const char *generator_name) {
+    CHECK_EQ(iss.peek(), std::char_traits<char>::eof())
+        << "Invalid " << generator_name << " generator state: trailing data";
+}
+
+std::mutex &DefaultGeneratorMutex() {
+    static std::mutex mutex;
+    return mutex;
+}
+
+uint64_t ProcessDefaultSeed() {
+    static const uint64_t seed = MakeRandomSeed();
+    return seed;
+}
+
+uint64_t &DefaultGeneratorSeed() {
+    static uint64_t seed = ProcessDefaultSeed();
+    return seed;
+}
+
+std::unordered_map<int8_t, std::shared_ptr<Generator>> &DefaultCUDAGenerators() {
+    static std::unordered_map<int8_t, std::shared_ptr<Generator>> generators;
+    return generators;
+}
+} // namespace
+
+class CPUGeneratorImpl : public GeneratorImpl {
+public:
+    explicit CPUGeneratorImpl(uint64_t seed);
+
+    void ManualSeed(uint64_t seed) override;
+    uint64_t Seed() override;
+    uint64_t InitialSeed() const override;
+    std::vector<uint8_t> GetState() const override;
+    void SetState(const std::vector<uint8_t> &state) override;
+    Device GetDevice() const override;
+
+    void FillUniform(float *data, size_t num_elements, float from, float to);
+    void FillNormal(float *data, size_t num_elements, float mean, float stddev);
+
+private:
+    mutable std::mutex mutex_;
+    uint64_t initial_seed_ = 0;
+    std::mt19937 engine_;
+    // Box-Muller produces two samples at a time. Preserve the second so the
+    // stream does not depend on how callers partition output tensors.
+    bool has_next_normal_sample_ = false;
+    float next_normal_sample_ = 0.0f;
+};
+
+class CUDAGeneratorImpl : public GeneratorImpl {
+public:
+    CUDAGeneratorImpl(int8_t device_index, uint64_t seed);
+
+    void ManualSeed(uint64_t seed) override;
+    uint64_t Seed() override;
+    uint64_t InitialSeed() const override;
+    std::vector<uint8_t> GetState() const override;
+    void SetState(const std::vector<uint8_t> &state) override;
+    Device GetDevice() const override;
+    std::pair<uint64_t, uint64_t> ReserveRandomOffset(uint64_t increment);
+
+private:
+    mutable std::mutex mutex_;
+    Device device_;
+    uint64_t initial_seed_ = 0;
+    uint64_t offset_ = 0;
+};
+
+Generator::Generator(std::shared_ptr<GeneratorImpl> impl) : impl_(std::move(impl)) { CHECK(impl_ != nullptr); }
+
+void Generator::ManualSeed(uint64_t seed) { impl_->ManualSeed(seed); }
+
+uint64_t Generator::Seed() { return impl_->Seed(); }
+
+uint64_t Generator::InitialSeed() const { return impl_->InitialSeed(); }
+
+std::vector<uint8_t> Generator::GetState() const { return impl_->GetState(); }
+
+void Generator::SetState(const std::vector<uint8_t> &state) { impl_->SetState(state); }
+
+Device Generator::GetDevice() const { return impl_->GetDevice(); }
+
+std::pair<uint64_t, uint64_t>
+detail::GeneratorAccessor::ReserveCUDARandomOffset(const std::shared_ptr<Generator> &generator, uint64_t increment) {
+    CHECK(generator != nullptr);
+    auto *cuda_generator = dynamic_cast<CUDAGeneratorImpl *>(generator->impl_.get());
+    CHECK(cuda_generator != nullptr) << "CUDA random offset requires a CUDA generator";
+    return cuda_generator->ReserveRandomOffset(increment);
+}
+
+void detail::GeneratorAccessor::FillCPUUniform(const std::shared_ptr<Generator> &generator, float *data,
+                                               size_t num_elements, float from, float to) {
+    CHECK(generator != nullptr);
+    auto *cpu_generator = dynamic_cast<CPUGeneratorImpl *>(generator->impl_.get());
+    CHECK(cpu_generator != nullptr) << "CPU random fill requires a CPU generator";
+    cpu_generator->FillUniform(data, num_elements, from, to);
+}
+
+void detail::GeneratorAccessor::FillCPUNormal(const std::shared_ptr<Generator> &generator, float *data,
+                                              size_t num_elements, float mean, float stddev) {
+    CHECK(generator != nullptr);
+    auto *cpu_generator = dynamic_cast<CPUGeneratorImpl *>(generator->impl_.get());
+    CHECK(cpu_generator != nullptr) << "CPU random fill requires a CPU generator";
+    cpu_generator->FillNormal(data, num_elements, mean, stddev);
+}
+
+CPUGeneratorImpl::CPUGeneratorImpl(uint64_t seed) { ManualSeed(seed); }
+
+void CPUGeneratorImpl::ManualSeed(uint64_t seed) {
+    std::lock_guard<std::mutex> lock(mutex_);
+    initial_seed_ = seed;
+    SeedEngine(engine_, seed);
+    has_next_normal_sample_ = false;
+    next_normal_sample_ = 0.0f;
+}
+
+uint64_t CPUGeneratorImpl::Seed() {
+    const uint64_t seed = MakeRandomSeed();
+    ManualSeed(seed);
+    return seed;
+}
+
+uint64_t CPUGeneratorImpl::InitialSeed() const {
+    std::lock_guard<std::mutex> lock(mutex_);
+    return initial_seed_;
+}
+
+std::vector<uint8_t> CPUGeneratorImpl::GetState() const {
+    std::lock_guard<std::mutex> lock(mutex_);
+    std::ostringstream oss;
+    oss.imbue(std::locale::classic());
+    oss << kCPUStateHeader << "\n"
+        << initial_seed_ << "\n"
+        << has_next_normal_sample_ << "\n"
+        << std::bit_cast<uint32_t>(next_normal_sample_) << "\n"
+        << engine_;
+    return ToBytes(oss.str());
+}
+
+void CPUGeneratorImpl::SetState(const std::vector<uint8_t> &state) {
+    const std::string serialized = ToString(state);
+    std::istringstream iss(serialized);
+    iss.imbue(std::locale::classic());
+
+    std::string header;
+    std::mt19937 engine;
+    std::getline(iss, header);
+    CHECK_EQ(header, kCPUStateHeader) << "Invalid CPU generator state header";
+    const uint64_t seed = ParseStateIntegerLine<uint64_t>(iss, "Invalid CPU generator seed in state");
+    const uint32_t has_next_normal
+        = ParseStateIntegerLine<uint32_t>(iss, "Invalid CPU generator normal cache flag in state");
+    CHECK_LE(has_next_normal, 1U) << "Invalid CPU generator normal cache flag in state";
+    const uint32_t next_normal_bits
+        = ParseStateIntegerLine<uint32_t>(iss, "Invalid CPU generator normal cache value in state");
+    CHECK(has_next_normal != 0 || next_normal_bits == 0) << "Invalid unused CPU generator normal cache value";
+    const float next_normal = std::bit_cast<float>(next_normal_bits);
+    CHECK(has_next_normal == 0 || std::isfinite(next_normal)) << "Invalid CPU generator normal cache value in state";
+    std::ostringstream engine_state_stream;
+    engine_state_stream << iss.rdbuf();
+    const std::string engine_state = engine_state_stream.str();
+    ValidateCPUEngineState(engine_state);
+    std::istringstream engine_iss(engine_state);
+    engine_iss.imbue(std::locale::classic());
+    CHECK(engine_iss >> engine) << "Invalid CPU generator engine state";
+    CheckNoTrailingStateData(engine_iss, "CPU");
+
+    std::lock_guard<std::mutex> lock(mutex_);
+    initial_seed_ = seed;
+    engine_ = engine;
+    has_next_normal_sample_ = has_next_normal != 0;
+    next_normal_sample_ = next_normal;
+}
+
+Device CPUGeneratorImpl::GetDevice() const { return Device(); }
+
+void CPUGeneratorImpl::FillUniform(float *data, size_t num_elements, float from, float to) {
+    CHECK(data != nullptr || num_elements == 0);
+    detail::CheckUniformBounds(from, to);
+    std::lock_guard<std::mutex> lock(mutex_);
+
+    const float range = to - from;
+    for (size_t i = 0; i < num_elements; ++i) {
+        const float value = from + range * UintToUnitFloat(engine_());
+        // Preserve the half-open interval if the final float rounding reaches the upper bound.
+        data[i] = value == to ? from : value;
+    }
+}
+
+void CPUGeneratorImpl::FillNormal(float *data, size_t num_elements, float mean, float stddev) {
+    CHECK(data != nullptr || num_elements == 0);
+    CHECK_GE(stddev, 0.0f);
+    std::lock_guard<std::mutex> lock(mutex_);
+    for (size_t i = 0; i < num_elements; ++i) {
+        float normal;
+        if (has_next_normal_sample_) {
+            normal = next_normal_sample_;
+            has_next_normal_sample_ = false;
+            next_normal_sample_ = 0.0f;
+        } else {
+            const uint32_t first_random = engine_();
+            const uint32_t second_random = engine_();
+            const auto [first, second] = BoxMullerPair(first_random, second_random);
+            normal = first;
+            has_next_normal_sample_ = true;
+            next_normal_sample_ = second;
+        }
+        data[i] = mean + stddev * normal;
+    }
+}
+
+CUDAGeneratorImpl::CUDAGeneratorImpl(int8_t device_index, uint64_t seed)
+    : device_(Device::DeviceType::kCUDA, device_index) {
+    CHECK_GE(device_index, 0);
+    ManualSeed(seed);
+}
+
+void CUDAGeneratorImpl::ManualSeed(uint64_t seed) {
+    std::lock_guard<std::mutex> lock(mutex_);
+    initial_seed_ = seed;
+    offset_ = 0;
+}
+
+uint64_t CUDAGeneratorImpl::Seed() {
+    const uint64_t seed = MakeRandomSeed();
+    ManualSeed(seed);
+    return seed;
+}
+
+uint64_t CUDAGeneratorImpl::InitialSeed() const {
+    std::lock_guard<std::mutex> lock(mutex_);
+    return initial_seed_;
+}
+
+std::vector<uint8_t> CUDAGeneratorImpl::GetState() const {
+    std::lock_guard<std::mutex> lock(mutex_);
+    std::ostringstream oss;
+    oss.imbue(std::locale::classic());
+    oss << kCUDAStateHeader << "\n" << initial_seed_ << "\n" << offset_ << "\n";
+    return ToBytes(oss.str());
+}
+
+void CUDAGeneratorImpl::SetState(const std::vector<uint8_t> &state) {
+    const std::string serialized = ToString(state);
+    std::istringstream iss(serialized);
+    iss.imbue(std::locale::classic());
+
+    std::string header;
+    std::getline(iss, header);
+    CHECK_EQ(header, kCUDAStateHeader) << "Invalid CUDA generator state header";
+    const uint64_t seed = ParseStateIntegerLine<uint64_t>(iss, "Invalid CUDA generator seed in state");
+    const uint64_t offset = ParseStateIntegerLine<uint64_t>(iss, "Invalid CUDA generator offset in state");
+    CheckNoTrailingStateData(iss, "CUDA");
+
+    std::lock_guard<std::mutex> lock(mutex_);
+    initial_seed_ = seed;
+    offset_ = offset;
+}
+
+Device CUDAGeneratorImpl::GetDevice() const { return device_; }
+
+std::pair<uint64_t, uint64_t> CUDAGeneratorImpl::ReserveRandomOffset(uint64_t increment) {
+    std::lock_guard<std::mutex> lock(mutex_);
+    CHECK_LE(increment, std::numeric_limits<uint64_t>::max() - offset_) << "CUDA generator offset overflow";
+    const uint64_t offset = offset_;
+    // offset_ addresses individual 32-bit words in one global Philox stream.
+    // It is not cuRAND's per-thread offset, so arbitrary exact increments are safe.
+    offset_ += increment;
+    return {initial_seed_, offset};
+}
+
+std::shared_ptr<Generator> MakeCPUGenerator(uint64_t seed) {
+    return std::shared_ptr<Generator>(new Generator(std::make_shared<CPUGeneratorImpl>(seed)));
+}
+
+std::shared_ptr<Generator> MakeCUDAGenerator(int8_t device_index, uint64_t seed) {
+    return std::shared_ptr<Generator>(new Generator(std::make_shared<CUDAGeneratorImpl>(device_index, seed)));
+}
+
+std::shared_ptr<Generator> GetDefaultCPUGenerator() {
+    static auto generator = MakeCPUGenerator(ProcessDefaultSeed());
+    return generator;
+}
+
+std::shared_ptr<Generator> GetDefaultCUDAGenerator(int8_t device_index) {
+    CHECK_GE(device_index, 0);
+    std::lock_guard<std::mutex> lock(DefaultGeneratorMutex());
+    auto &generators = DefaultCUDAGenerators();
+    auto it = generators.find(device_index);
+    if (it != generators.end()) {
+        return it->second;
+    }
+    auto generator = MakeCUDAGenerator(device_index, DefaultGeneratorSeed());
+    generators.emplace(device_index, generator);
+    return generator;
+}
+
+std::shared_ptr<Generator> GetDefaultGenerator(const Device &device) {
+    if (device.IsCPU()) {
+        return GetDefaultCPUGenerator();
+    }
+    if (device.IsCUDA()) {
+        return GetDefaultCUDAGenerator(device.index());
+    }
+    LOG(FATAL) << "Unsupported default Generator device: " << device;
+    return nullptr;
+}
+
+void ManualSeed(uint64_t seed) { ManualSeedAll(seed); }
+
+void ManualSeedAll(uint64_t seed) {
+    std::lock_guard<std::mutex> lock(DefaultGeneratorMutex());
+    GetDefaultCPUGenerator()->ManualSeed(seed);
+    DefaultGeneratorSeed() = seed;
+    for (auto &[_, generator] : DefaultCUDAGenerators()) { generator->ManualSeed(seed); }
+}
+
+} // namespace infini_train
diff --git a/infini_train/src/generator_internal.h b/infini_train/src/generator_internal.h
new file mode 100644
index 00000000..1eba3ef9
--- /dev/null
+++ b/infini_train/src/generator_internal.h
@@ -0,0 +1,40 @@
+#pragma once
+
+#include <cstddef>
+#include <cstdint>
+#include <memory>
+#include <utility>
+#include <vector>
+
+#include "infini_train/include/device.h"
+
+namespace infini_train {
+
+class Generator;
+
+class GeneratorImpl {
+public:
+    virtual ~GeneratorImpl() = default;
+
+    virtual void ManualSeed(uint64_t seed) = 0;
+    virtual uint64_t Seed() = 0;
+    virtual uint64_t InitialSeed() const = 0;
+    virtual std::vector<uint8_t> GetState() const = 0;
+    virtual void SetState(const std::vector<uint8_t> &state) = 0;
+    virtual Device GetDevice() const = 0;
+};
+
+namespace detail {
+
+class GeneratorAccessor {
+public:
+    static std::pair<uint64_t, uint64_t> ReserveCUDARandomOffset(const std::shared_ptr<Generator> &generator,
+                                                                 uint64_t increment);
+    static void FillCPUUniform(const std::shared_ptr<Generator> &generator, float *data, size_t num_elements,
+                               float from, float to);
+    static void FillCPUNormal(const std::shared_ptr<Generator> &generator, float *data, size_t num_elements, float mean,
+                              float stddev);
+};
+
+} // namespace detail
+} // namespace infini_train
diff --git a/infini_train/src/kernels/cpu/cross_entropy.cc b/infini_train/src/kernels/cpu/cross_entropy.cc
index f520b2e9..5cf95cf2 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/transform.cc b/infini_train/src/kernels/cpu/transform.cc
index 48063c7a..a76102b1 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/elementwise.cu b/infini_train/src/kernels/cuda/elementwise.cu
index fc423b35..96f3bda1 100644
--- a/infini_train/src/kernels/cuda/elementwise.cu
+++ b/infini_train/src/kernels/cuda/elementwise.cu
@@ -1,6 +1,7 @@
 #include <cstddef>
 
 #include <cub/warp/warp_reduce.cuh>
+#include <numeric>
 
 #include "infini_train/include/common/common.h"
 #include "infini_train/include/common/cuda/common_cuda.h"
diff --git a/infini_train/src/kernels/cuda/gather.cu b/infini_train/src/kernels/cuda/gather.cu
index d8b0cffa..a4790035 100644
--- a/infini_train/src/kernels/cuda/gather.cu
+++ b/infini_train/src/kernels/cuda/gather.cu
@@ -1,4 +1,5 @@
 #include "glog/logging.h"
+#include <numeric>
 
 #include "infini_train/include/common/common.h"
 #include "infini_train/include/common/cuda/common_cuda.h"
diff --git a/infini_train/src/kernels/cuda/no_op.cu b/infini_train/src/kernels/cuda/no_op.cu
index ef2c9566..e12fd613 100644
--- a/infini_train/src/kernels/cuda/no_op.cu
+++ b/infini_train/src/kernels/cuda/no_op.cu
@@ -1,4 +1,5 @@
 #include "glog/logging.h"
+#include <numeric>
 
 #include "infini_train/include/dispatcher.h"
 #include "infini_train/include/tensor.h"
diff --git a/infini_train/src/kernels/cuda/random.cu b/infini_train/src/kernels/cuda/random.cu
new file mode 100644
index 00000000..55dcf0c1
--- /dev/null
+++ b/infini_train/src/kernels/cuda/random.cu
@@ -0,0 +1,181 @@
+#include <cmath>
+#include <cstdint>
+#include <limits>
+
+#include "infini_train/include/common/cuda/common_cuda.h"
+#include "infini_train/include/core/runtime/device_guard.h"
+#include "infini_train/include/device.h"
+
+#include "infini_train/src/core/runtime/cuda/cuda_runtime_common.h"
+#include "infini_train/src/random_utils.h"
+
+namespace infini_train::kernels::cuda {
+namespace {
+constexpr int kThreadsPerBlock = 256;
+constexpr float kTwoPi = 6.283185307179586476925286766559f;
+constexpr int kUniformValuesPerThread = 4;
+constexpr int kNormalValuesPerThread = 2;
+
+struct Philox4x32State {
+    uint32_t c0;
+    uint32_t c1;
+    uint32_t c2;
+    uint32_t c3;
+};
+
+__device__ uint32_t MulHi(uint32_t a, uint32_t b) {
+    return static_cast<uint32_t>((static_cast<uint64_t>(a) * b) >> 32);
+}
+
+__device__ Philox4x32State PhiloxRound(Philox4x32State counter, uint32_t key0, uint32_t key1) {
+    constexpr uint32_t kPhiloxM0 = 0xD2511F53;
+    constexpr uint32_t kPhiloxM1 = 0xCD9E8D57;
+
+    const uint32_t lo0 = counter.c0 * kPhiloxM0;
+    const uint32_t hi0 = MulHi(counter.c0, kPhiloxM0);
+    const uint32_t lo1 = counter.c2 * kPhiloxM1;
+    const uint32_t hi1 = MulHi(counter.c2, kPhiloxM1);
+
+    return {hi1 ^ counter.c1 ^ key0, lo1, hi0 ^ counter.c3 ^ key1, lo0};
+}
+
+__device__ Philox4x32State Philox(uint64_t seed, uint64_t counter_index) {
+    constexpr uint32_t kPhiloxW0 = 0x9E3779B9;
+    constexpr uint32_t kPhiloxW1 = 0xBB67AE85;
+
+    Philox4x32State counter{static_cast<uint32_t>(counter_index), static_cast<uint32_t>(counter_index >> 32), 0, 0};
+    uint32_t key0 = static_cast<uint32_t>(seed);
+    uint32_t key1 = static_cast<uint32_t>(seed >> 32);
+
+    for (int round = 0; round < 10; ++round) {
+        counter = PhiloxRound(counter, key0, key1);
+        key0 += kPhiloxW0;
+        key1 += kPhiloxW1;
+    }
+    return counter;
+}
+
+__device__ uint32_t PhiloxStateWord(const Philox4x32State &values, uint32_t lane) {
+    switch (lane) {
+    case 0:
+        return values.c0;
+    case 1:
+        return values.c1;
+    case 2:
+        return values.c2;
+    default:
+        return values.c3;
+    }
+}
+
+__device__ void PhiloxRandomWords(uint64_t seed, uint64_t offset, uint32_t *output, int count) {
+    uint64_t counter_index = offset / 4;
+    uint32_t lane = static_cast<uint32_t>(offset % 4);
+    Philox4x32State values = Philox(seed, counter_index);
+
+#pragma unroll
+    for (int i = 0; i < 4; ++i) {
+        if (i >= count) {
+            return;
+        }
+        output[i] = PhiloxStateWord(values, lane++);
+        if (lane == 4 && i + 1 < count) {
+            values = Philox(seed, ++counter_index);
+            lane = 0;
+        }
+    }
+}
+
+__device__ float UintToUnitFloat(uint32_t value) { return static_cast<float>(value >> 8) * 0x1.0p-24f; }
+
+__global__ void RandomUniformFloat32Kernel(float *data, int64_t num_elements, float from, float to, uint64_t seed,
+                                           uint64_t offset) {
+    const int64_t thread_index = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    const int64_t first_element = thread_index * kUniformValuesPerThread;
+    if (first_element >= num_elements) {
+        return;
+    }
+
+    const int count
+        = static_cast<int>(num_elements - first_element < kUniformValuesPerThread ? num_elements - first_element
+                                                                                  : kUniformValuesPerThread);
+    uint32_t random_words[kUniformValuesPerThread];
+    PhiloxRandomWords(seed, offset + static_cast<uint64_t>(first_element), random_words, count);
+
+#pragma unroll
+    for (int i = 0; i < kUniformValuesPerThread; ++i) {
+        if (i >= count) {
+            return;
+        }
+        const float u = UintToUnitFloat(random_words[i]);
+        const float value = from + (to - from) * u;
+        // Preserve the half-open interval if the final float rounding reaches the upper bound.
+        data[first_element + i] = value == to ? from : value;
+    }
+}
+
+__global__ void RandomNormalFloat32Kernel(float *data, int64_t num_elements, float mean, float stddev, uint64_t seed,
+                                          uint64_t offset) {
+    const int64_t thread_index = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    const int64_t first_element = thread_index * kNormalValuesPerThread;
+    if (first_element >= num_elements) {
+        return;
+    }
+
+    const int count = static_cast<int>(
+        num_elements - first_element < kNormalValuesPerThread ? num_elements - first_element : kNormalValuesPerThread);
+    uint32_t random_words[kNormalValuesPerThread * 2];
+    PhiloxRandomWords(seed, offset + static_cast<uint64_t>(first_element) * 2, random_words, count * 2);
+
+#pragma unroll
+    for (int i = 0; i < kNormalValuesPerThread; ++i) {
+        if (i >= count) {
+            return;
+        }
+        const float u1 = 1.0f - UintToUnitFloat(random_words[i * 2]);
+        const float u2 = UintToUnitFloat(random_words[i * 2 + 1]);
+        const float z = sqrtf(-2.0f * logf(u1)) * cosf(kTwoPi * u2);
+        data[first_element + i] = mean + stddev * z;
+    }
+}
+
+cudaStream_t GetCudaStream(Device device) {
+    auto *stream = dynamic_cast<infini_train::core::cuda::CudaStream *>(
+        infini_train::core::GetDeviceGuardImpl(device.type())->GetStream(device));
+    CHECK_NOTNULL(stream);
+    return stream->cuda_stream();
+}
+} // namespace
+
+void RandomUniformFloat32(void *data, int64_t num_elements, float from, float to, uint64_t seed, uint64_t offset,
+                          Device device) {
+    CHECK_GE(num_elements, 0);
+    infini_train::detail::CheckUniformBounds(from, to);
+    if (num_elements == 0) {
+        return;
+    }
+    const int64_t work_items = num_elements / kUniformValuesPerThread + (num_elements % kUniformValuesPerThread != 0);
+    CHECK_LE(work_items, static_cast<int64_t>(std::numeric_limits<int>::max()) * kThreadsPerBlock)
+        << "Random uniform tensor is too large";
+    const int blocks = static_cast<int>((work_items + kThreadsPerBlock - 1) / kThreadsPerBlock);
+    RandomUniformFloat32Kernel<<<blocks, kThreadsPerBlock, 0, GetCudaStream(device)>>>(
+        static_cast<float *>(data), num_elements, from, to, seed, offset);
+    CUDA_CHECK(cudaGetLastError());
+}
+
+void RandomNormalFloat32(void *data, int64_t num_elements, float mean, float stddev, uint64_t seed, uint64_t offset,
+                         Device device) {
+    CHECK_GE(num_elements, 0);
+    CHECK_GE(stddev, 0.0f);
+    if (num_elements == 0) {
+        return;
+    }
+    const int64_t work_items = num_elements / kNormalValuesPerThread + (num_elements % kNormalValuesPerThread != 0);
+    CHECK_LE(work_items, static_cast<int64_t>(std::numeric_limits<int>::max()) * kThreadsPerBlock)
+        << "Random normal tensor is too large";
+    const int blocks = static_cast<int>((work_items + kThreadsPerBlock - 1) / kThreadsPerBlock);
+    RandomNormalFloat32Kernel<<<blocks, kThreadsPerBlock, 0, GetCudaStream(device)>>>(
+        static_cast<float *>(data), num_elements, mean, stddev, seed, offset);
+    CUDA_CHECK(cudaGetLastError());
+}
+} // namespace infini_train::kernels::cuda
diff --git a/infini_train/src/kernels/cuda/reduction.cu b/infini_train/src/kernels/cuda/reduction.cu
index c56470e3..793e001b 100644
--- a/infini_train/src/kernels/cuda/reduction.cu
+++ b/infini_train/src/kernels/cuda/reduction.cu
@@ -1,4 +1,5 @@
 #include <cub/cub.cuh>
+#include <numeric>
 
 #include "infini_train/include/common/cuda/common_cuda.h"
 #include "infini_train/include/common/cuda/cub_compat.cuh"
diff --git a/infini_train/src/nn/functional.cc b/infini_train/src/nn/functional.cc
index c33e2368..2b912a1a 100644
--- a/infini_train/src/nn/functional.cc
+++ b/infini_train/src/nn/functional.cc
@@ -9,6 +9,7 @@
 #include "infini_train/include/autograd/reduction.h"
 #include "infini_train/include/autograd/softmax.h"
 #include "infini_train/include/autograd/transform.h"
+#include "infini_train/include/generator.h"
 #include "infini_train/include/nn/init.h"
 #include "infini_train/include/tensor.h"
 
@@ -26,6 +27,16 @@ 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, Device device, std::shared_ptr<Generator> generator) {
+    auto tensor = std::make_shared<Tensor>(size, DataType::kFLOAT32, device);
+    return init::Uniform(tensor, 0.0f, 1.0f, generator);
+}
+
+std::shared_ptr<Tensor> Randn(const std::vector<int64_t> &size, Device device, std::shared_ptr<Generator> generator) {
+    auto tensor = std::make_shared<Tensor>(size, DataType::kFLOAT32, device);
+    return init::Normal(tensor, 0.0f, 1.0f, generator);
+}
+
 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 79b4b48b..c3c92813 100644
--- a/infini_train/src/nn/init.cc
+++ b/infini_train/src/nn/init.cc
@@ -1,67 +1,75 @@
 #include "infini_train/include/nn/init.h"
 
-#include <cstring>
+#include <cmath>
 #include <functional>
+#include <limits>
 #include <memory>
+#include <numeric>
 #include <optional>
-#include <random>
 #include <unordered_set>
 
-#ifdef USE_OMP
-#include <omp.h>
-#endif
-
 #include "glog/logging.h"
 
 #include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/device.h"
+#include "infini_train/include/generator.h"
 #include "infini_train/include/tensor.h"
+#include "infini_train/src/generator_internal.h"
+#include "infini_train/src/random_utils.h"
+
+#ifdef USE_CUDA
+namespace infini_train::kernels::cuda {
+void RandomUniformFloat32(void *data, int64_t num_elements, float from, float to, uint64_t seed, uint64_t offset,
+                          Device device);
+void RandomNormalFloat32(void *data, int64_t num_elements, float mean, float stddev, uint64_t seed, uint64_t offset,
+                         Device device);
+} // namespace infini_train::kernels::cuda
+#endif
 
 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);
+int64_t CheckedMultiplyFanFactors(int64_t lhs, int64_t rhs) {
+    CHECK_GE(lhs, 0) << "Fan calculation requires non-negative tensor dimensions";
+    CHECK_GE(rhs, 0) << "Fan calculation requires non-negative tensor dimensions";
+    if (lhs == 0 || rhs == 0) {
+        return 0;
+    }
+    CHECK_LE(lhs, std::numeric_limits<int64_t>::max() / rhs) << "Fan calculation overflow";
+    return lhs * rhs;
+}
 } // namespace
 
-std::shared_ptr<Tensor> Normal(const std::shared_ptr<Tensor> &tensor, float mean, float std,
-                               std::optional<std::mt19937> generator) {
+std::shared_ptr<Tensor> Normal(const std::shared_ptr<Tensor> &tensor, float mean, float stddev,
+                               std::shared_ptr<Generator> generator) {
+    CHECK_GE(stddev, 0.0f);
+    CHECK(tensor->Dtype() == DataType::kFLOAT32) << "Random normal currently supports float32 tensors";
     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);
-        }
-    }
-#else
-    std::normal_distribution<float> dis(mean, std);
-    std::generate(buffer.begin(), buffer.end(), [&]() { return generator ? dis(generator.value()) : dis(gen); });
-#endif
 
     auto device = tensor->GetDevice();
-    core::DeviceGuard guard(device);
-    auto impl = core::GetDeviceGuardImpl(device.type());
+    auto resolved_generator = generator ? generator : GetDefaultGenerator(device);
+    CHECK(resolved_generator->GetDevice().type() == device.type())
+        << "Generator backend must match tensor device backend: generator=" << resolved_generator->GetDevice()
+        << " tensor=" << device;
+    if (num_elements == 0) {
+        return tensor;
+    }
 
-    impl->MemcpyAsync(tensor->DataPtr(), buffer.data(), num_elements * sizeof(float),
-                      device.type() == Device::DeviceType::kCPU ? core::MemcpyKind::kD2D : core::MemcpyKind::kH2D,
-                      impl->GetStream(device));
+#ifdef USE_CUDA
+    if (device.IsCUDA()) {
+        CHECK_LE(num_elements, static_cast<int64_t>(std::numeric_limits<uint64_t>::max() / 2))
+            << "Random normal tensor is too large";
+        const auto [seed, offset] = detail::GeneratorAccessor::ReserveCUDARandomOffset(
+            resolved_generator, static_cast<uint64_t>(num_elements) * 2);
+        core::DeviceGuard guard(device);
+        kernels::cuda::RandomNormalFloat32(tensor->DataPtr(), num_elements, mean, stddev, seed, offset, device);
+        return tensor;
+    }
+#endif
+
+    CHECK(device.IsCPU()) << "Random normal backend is not available for device " << device;
+    detail::GeneratorAccessor::FillCPUNormal(resolved_generator, static_cast<float *>(tensor->DataPtr()),
+                                             tensor->NumElements(), mean, stddev);
     return tensor;
 }
 
@@ -71,13 +79,10 @@ std::pair<int64_t, int64_t> CalculateFanInAndFanOut(const std::shared_ptr<Tensor
     }
     const auto num_input_fmaps = tensor->Dims()[1];
     const auto num_output_fmaps = tensor->Dims()[0];
-    int64_t receptive_field_size = 1;
-    if (tensor->Dims().size() > 2) {
-        receptive_field_size
-            *= std::accumulate(tensor->Dims().begin() + 2, tensor->Dims().end(), 1, std::multiplies<int64_t>());
-    }
-    const auto fan_in = num_input_fmaps * receptive_field_size;
-    const auto fan_out = num_output_fmaps * receptive_field_size;
+    const int64_t receptive_field_size
+        = std::accumulate(tensor->Dims().begin() + 2, tensor->Dims().end(), int64_t{1}, CheckedMultiplyFanFactors);
+    const int64_t fan_in = CheckedMultiplyFanFactors(num_input_fmaps, receptive_field_size);
+    const int64_t fan_out = CheckedMultiplyFanFactors(num_output_fmaps, receptive_field_size);
     return {fan_in, fan_out};
 }
 
@@ -113,7 +118,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::shared_ptr<Generator> generator) {
     for (const auto dim : tensor->Dims()) {
         if (dim == 0) {
             LOG(WARNING) << "Initializing zero-element tensors is a no-op";
@@ -122,39 +127,39 @@ std::shared_ptr<Tensor> KaimingUniform(const std::shared_ptr<Tensor> &tensor, fl
     }
     const auto fan = CalculateCorrectFan(tensor, mode);
     const auto gain = CalculateGain(nonlinearity, a);
-    const float std = gain / sqrt(fan);
-    const float bound = sqrt(3.0f) * std; // Calculate uniform bounds from standard deviation
+    const float stddev = gain / sqrt(fan);
+    const float bound = sqrt(3.0f) * stddev; // Calculate uniform bounds from standard deviation
     return tensor->Uniform(-bound, bound, generator);
 }
 
 std::shared_ptr<Tensor> Uniform(const std::shared_ptr<Tensor> &tensor, float a, float b,
-                                std::optional<std::mt19937> generator) {
+                                std::shared_ptr<Generator> generator) {
+    detail::CheckUniformBounds(a, b);
+    CHECK(tensor->Dtype() == DataType::kFLOAT32) << "Random uniform currently supports float32 tensors";
     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
 
     auto device = tensor->GetDevice();
+    auto resolved_generator = generator ? generator : GetDefaultGenerator(device);
+    CHECK(resolved_generator->GetDevice().type() == device.type())
+        << "Generator backend must match tensor device backend: generator=" << resolved_generator->GetDevice()
+        << " tensor=" << device;
+    if (num_elements == 0) {
+        return tensor;
+    }
 
-    core::DeviceGuard guard(device);
-    auto impl = core::GetDeviceGuardImpl(device.type());
+#ifdef USE_CUDA
+    if (device.IsCUDA()) {
+        const auto [seed, offset] = detail::GeneratorAccessor::ReserveCUDARandomOffset(
+            resolved_generator, static_cast<uint64_t>(num_elements));
+        core::DeviceGuard guard(device);
+        kernels::cuda::RandomUniformFloat32(tensor->DataPtr(), num_elements, a, b, seed, offset, device);
+        return tensor;
+    }
+#endif
 
-    impl->MemcpyAsync(tensor->DataPtr(), buffer.data(), num_elements * sizeof(float),
-                      device.type() == Device::DeviceType::kCPU ? core::MemcpyKind::kD2D : core::MemcpyKind::kH2D,
-                      impl->GetStream(device));
+    CHECK(device.IsCPU()) << "Random uniform backend is not available for device " << device;
+    detail::GeneratorAccessor::FillCPUUniform(resolved_generator, static_cast<float *>(tensor->DataPtr()),
+                                              tensor->NumElements(), a, b);
 
     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 b08b64fb..053d56d4 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>
 
@@ -22,17 +24,19 @@ DistributedDataParallel::DistributedDataParallel(std::shared_ptr<nn::Module> mod
                                                  const DistributedDataParallelConfig ddp_config)
     : ddp_config_(ddp_config),
       ddp_pg_(ProcessGroupFactory::Instance()->Get(GetDataParallelProcessGroupName(rank.GlobalRank()))) {
+    CHECK(module != nullptr) << "DistributedDataParallel: module must not be null";
     CHECK(ddp_config_.zero_stage >= 0 && ddp_config_.zero_stage <= 3)
         << "DistributedDataParallel: zero_stage must be in 0/1/2/3.";
     if (ddp_config_.zero_stage == 3) {
         LOG(FATAL) << "DistributedDataParallel: ZeRO-3 is not implemented yet.";
     }
     for (auto &param : module->Parameters()) {
+        CHECK(param != nullptr) << "DistributedDataParallel: module contains a null parameter";
+        auto device = param->GetDevice();
+        CHECK_EQ(device.index(), rank.thread_rank()) << "All parameters must be on the same device as the module";
         if (!param->requires_grad()) {
             continue;
         }
-        auto device = param->GetDevice();
-        CHECK_EQ(device.index(), rank.thread_rank()) << "All parameters must be on the same device as the module";
         if (!ddp_config.gradient_bucketing_enabled && ddp_config.zero_stage < 1) {
             auto hook = std::make_unique<infini_train::autograd::AllReducePostAccumulateHook>(
                 function::ReduceOpType::kAvg, ddp_pg_);
@@ -40,10 +44,12 @@ DistributedDataParallel::DistributedDataParallel(std::shared_ptr<nn::Module> mod
         }
     }
     for (auto &buffer : module->Buffers()) {
+        CHECK(buffer != nullptr) << "DistributedDataParallel: module contains a null buffer";
         CHECK_EQ(buffer->GetDevice().index(), rank.thread_rank())
             << "All buffers must be on the same device as the module";
     }
     modules_[kModuleName] = std::move(module);
+    SynchronizeModuleState();
 
     if (ddp_config.zero_stage >= 1) {
         BuildParamAndGradBuffers();
@@ -61,6 +67,32 @@ DistributedDataParallel::DistributedDataParallel(std::shared_ptr<nn::Module> mod
     }
 }
 
+void DistributedDataParallel::SynchronizeModuleState() {
+    auto state = modules_.at(kModuleName)->StateDict();
+    std::vector<std::pair<std::string, std::shared_ptr<Tensor>>> named_tensors(state.begin(), state.end());
+    std::sort(named_tensors.begin(), named_tensors.end(),
+              [](const auto &lhs, const auto &rhs) { return lhs.first < rhs.first; });
+
+    std::vector<std::shared_ptr<Tensor>> tensors;
+    tensors.reserve(named_tensors.size());
+    for (const auto &[_, tensor] : named_tensors) { tensors.push_back(tensor); }
+
+    if (tensors.empty()) {
+        return;
+    }
+    CHECK(tensors.front() != nullptr) << "DDP module state tensor '" << named_tensors.front().first << "' is null";
+    const auto expected_device = tensors.front()->GetDevice();
+    for (size_t i = 0; i < tensors.size(); ++i) {
+        CHECK(tensors[i] != nullptr) << "DDP module state tensor '" << named_tensors[i].first << "' is null";
+        CHECK_EQ(tensors[i]->GetDevice(), expected_device)
+            << "DDP module state tensors must be on one device; tensor '" << named_tensors[i].first << "' is on "
+            << tensors[i]->GetDevice() << ", expected " << expected_device;
+    }
+
+    constexpr int kDataParallelRootRank = 0;
+    ddp_pg_->Broadcast(tensors, kDataParallelRootRank);
+}
+
 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 ab3a8002..b5afd48b 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/nn/parallel/process_group.cc b/infini_train/src/nn/parallel/process_group.cc
index 174aa645..9d9ffb4a 100644
--- a/infini_train/src/nn/parallel/process_group.cc
+++ b/infini_train/src/nn/parallel/process_group.cc
@@ -36,8 +36,8 @@ ProcessGroup::ProcessGroup(int world_size, const std::string &name) : world_size
 
 ProcessGroup::ProcessGroup(Device::DeviceType backend, const std::string &process_group_name,
                            const std::vector<int> &ranks)
-    : backend_(backend), runtime_impl_(core::GetDeviceGuardImpl(backend)), ccl_impl_(core::GetCclImpl(backend)),
-      world_size_(ranks.size()), name_(process_group_name) {
+    : world_size_(ranks.size()), name_(process_group_name), backend_(backend),
+      runtime_impl_(core::GetDeviceGuardImpl(backend)), ccl_impl_(core::GetCclImpl(backend)) {
     CHECK_GT(world_size_, 0);
     if (global::GetNnodes() == 1 && global::GetNprocPerNode() == 1) {
         InitSingleProcess(ranks);
@@ -63,8 +63,8 @@ void ProcessGroup::InitSingleProcess(const std::vector<int> &ranks) {
     std::vector<core::CclComm *> comm_ptrs(static_cast<size_t>(world_size_), nullptr);
     ccl_impl_->CommInitAll(comm_ptrs.data(), world_size_, ranks.data());
 
-    for (int i = 0; i < ranks.size(); ++i) {
-        auto *comm_raw = comm_ptrs[static_cast<size_t>(i)];
+    for (size_t i = 0; i < ranks.size(); ++i) {
+        auto *comm_raw = comm_ptrs[i];
         CHECK_NOTNULL(comm_raw);
         comms_.emplace_back(comm_raw);
 
@@ -72,7 +72,7 @@ void ProcessGroup::InitSingleProcess(const std::vector<int> &ranks) {
         devices_.push_back(device);
 
         device_comm_map_[device.index()] = comm_raw;
-        global_group_rank_map_[device.Rank().GlobalRank()] = i;
+        global_group_rank_map_[device.Rank().GlobalRank()] = static_cast<int>(i);
     }
 }
 
@@ -198,10 +198,15 @@ std::shared_ptr<Work> ProcessGroup::Broadcast(const std::vector<std::shared_ptr<
                                               int root_rank_in_group, bool async_op) const {
     CHECK_GE(root_rank_in_group, 0);
     CHECK_LT(root_rank_in_group, world_size_);
-    CHECK_GT(tensors.size(), 0);
+    CHECK(!tensors.empty());
     CHECK_NOTNULL(tensors[0]);
 
-    auto device = tensors[0]->GetDevice();
+    const auto device = tensors[0]->GetDevice();
+    CHECK(device.type() == backend_) << "Broadcast tensor backend must match ProcessGroup backend";
+    for (const auto &tensor : tensors) {
+        CHECK_NOTNULL(tensor);
+        CHECK_EQ(device, tensor->GetDevice());
+    }
     auto group_rank = GetGroupRank(device.Rank().GlobalRank());
     core::DeviceGuard guard(device);
     auto *compute_stream = runtime_impl_->GetStream(device);
@@ -212,8 +217,6 @@ std::shared_ptr<Work> ProcessGroup::Broadcast(const std::vector<std::shared_ptr<
     runtime_impl_->EventRecord(work->ready_event(), compute_stream);
     runtime_impl_->StreamWaitEvent(comm_stream, work->ready_event(), 0);
     for (const auto &tensor : tensors) {
-        CHECK_NOTNULL(tensor);
-        CHECK_EQ(device, tensor->GetDevice());
         const void *send_buffer = (group_rank == root_rank_in_group) ? tensor->DataPtr() : nullptr;
         ccl_impl_->Broadcast(send_buffer, tensor->DataPtr(), tensor->NumElements(), tensor->Dtype(), root_rank_in_group,
                              comm, comm_stream);
@@ -512,27 +515,20 @@ std::shared_ptr<Tensor> ProcessGroup::Gather(const std::vector<std::shared_ptr<T
 
 ProcessGroupFactory *ProcessGroupFactory::Instance() {
     // NOTE(zbl): Instance() with no arguments only gets initialized instance with a certain backend
+    std::lock_guard<std::mutex> lock(g_process_group_factory_mutex);
     auto &instance = g_process_group_factory_instance;
     if (instance == nullptr) {
-        std::lock_guard<std::mutex> lock(g_process_group_factory_mutex);
-        if (instance == nullptr) {
-            LOG(FATAL) << "ProcessGroupFactory is not initialized with backend. "
-                       << "Call ProcessGroupFactory::Instance(backend) first.";
-        }
+        LOG(FATAL) << "ProcessGroupFactory is not initialized with backend. "
+                   << "Call ProcessGroupFactory::Instance(backend) first.";
     }
     return instance.get();
 }
 
 ProcessGroupFactory *ProcessGroupFactory::Instance(Device::DeviceType backend) {
+    std::lock_guard<std::mutex> lock(g_process_group_factory_mutex);
     auto &instance = g_process_group_factory_instance;
     if (instance == nullptr) {
-        std::lock_guard<std::mutex> lock(g_process_group_factory_mutex);
-        if (instance == nullptr) {
-            instance.reset(new ProcessGroupFactory(backend));
-        } else if (instance->backend_ != backend) {
-            LOG(FATAL) << "ProcessGroupFactory backend mismatch. initialized=" << static_cast<int>(instance->backend_)
-                       << ", requested=" << static_cast<int>(backend);
-        }
+        instance.reset(new ProcessGroupFactory(backend));
     } else if (instance->backend_ != backend) {
         LOG(FATAL) << "ProcessGroupFactory backend mismatch. initialized=" << static_cast<int>(instance->backend_)
                    << ", requested=" << static_cast<int>(backend);
@@ -551,13 +547,12 @@ const ProcessGroup *ProcessGroupFactory::GetOrCreate(const std::string &name, co
 }
 
 const ProcessGroup *ProcessGroupFactory::Get(const std::string &name) const {
-    std::lock_guard<std::mutex> lock(mutex_);
+    std::unique_lock<std::mutex> lock(mutex_);
+    cond_.wait(lock, [this, &name]() { return name_to_group_.at(name) != nullptr; });
     return name_to_group_.at(name).get();
 }
 
-const ProcessGroup *ProcessGroupFactory::GetDefaultProcessGroup() const {
-    return name_to_group_.at(kDefaltProcessGroupName).get();
-}
+const ProcessGroup *ProcessGroupFactory::GetDefaultProcessGroup() const { return Get(kDefaltProcessGroupName); }
 
 ProcessGroupFactory::ProcessGroupFactory(Device::DeviceType backend) : backend_(backend) {
     GetOrCreate(kDefaltProcessGroupName, global::GetWorldSize());
diff --git a/infini_train/src/random_utils.h b/infini_train/src/random_utils.h
new file mode 100644
index 00000000..a9153423
--- /dev/null
+++ b/infini_train/src/random_utils.h
@@ -0,0 +1,19 @@
+#pragma once
+
+#include <cmath>
+#include <limits>
+
+#include "glog/logging.h"
+
+namespace infini_train::detail {
+
+inline void CheckUniformBounds(float from, float to) {
+    CHECK_LE(from, to);
+    CHECK(std::isfinite(from)) << "Uniform lower bound must be finite";
+    CHECK(std::isfinite(to)) << "Uniform upper bound must be finite";
+    const double range = static_cast<double>(to) - static_cast<double>(from);
+    CHECK_LE(range, static_cast<double>(std::numeric_limits<float>::max()))
+        << "Uniform bounds range exceeds float maximum";
+}
+
+} // namespace infini_train::detail
diff --git a/infini_train/src/tensor.cc b/infini_train/src/tensor.cc
index 18ca3d22..b454a0e5 100644
--- a/infini_train/src/tensor.cc
+++ b/infini_train/src/tensor.cc
@@ -471,7 +471,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::shared_ptr<Generator> generator) {
     return nn::init::Uniform(shared_from_this(), from, to, generator);
 }
 
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 96776585..d571cf08 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -10,6 +10,9 @@ add_subdirectory(common)
 # Tensor tests
 add_subdirectory(tensor)
 
+# Generator tests
+add_subdirectory(generator)
+
 # Optimizer tests
 add_subdirectory(optimizer)
 
diff --git a/tests/generator/CMakeLists.txt b/tests/generator/CMakeLists.txt
new file mode 100644
index 00000000..8aceceb4
--- /dev/null
+++ b/tests/generator/CMakeLists.txt
@@ -0,0 +1,28 @@
+# Generator tests are kept separate from Tensor tests. Shared behavior is
+# compiled for both backends, while CUDA-specific coverage stays isolated.
+file(GLOB GENERATOR_TEST_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/test_*.cc)
+list(REMOVE_ITEM GENERATOR_TEST_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/test_generator_ddp.cc)
+infini_train_add_test(test_generator_cpu
+  SOURCES ${GENERATOR_TEST_SOURCES}
+  LABELS cpu
+  TEST_FILTER "CPU/*"
+)
+
+if(USE_CUDA)
+  file(GLOB GENERATOR_CUDA_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/cuda_only/test_*.cc)
+  infini_train_add_test(test_generator_cuda
+    SOURCES ${GENERATOR_TEST_SOURCES} ${GENERATOR_CUDA_SOURCES}
+    LABELS cuda
+    TEST_FILTER "CUDA/*"
+  )
+
+  if(USE_NCCL)
+    add_executable(test_generator_ddp test_generator_ddp.cc)
+    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 30 RUN_SERIAL TRUE)
+  endif()
+endif()
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 00000000..0ee38f3b
--- /dev/null
+++ b/tests/generator/cuda_only/test_generator_cuda.cc
@@ -0,0 +1,98 @@
+#include <cstdint>
+#include <sstream>
+#include <string>
+#include <thread>
+#include <vector>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/generator.h"
+#include "infini_train/include/nn/functional.h"
+#include "tests/generator/generator_test_utils.h"
+
+namespace infini_train::test {
+namespace {
+
+uint64_t CUDAOffset(const std::vector<uint8_t> &state) {
+    std::istringstream stream(std::string(state.begin(), state.end()));
+    std::string line;
+    for (int i = 0; i < 3; ++i) { CHECK(std::getline(stream, line)); }
+    return std::stoull(line);
+}
+
+} // namespace
+
+TEST_P(GeneratorDeviceTest, CUDAUsesStablePhiloxStreamAndTracksConsumption) {
+    ONLY_CUDA();
+    // Philox4x32-10 known-answer vector for an all-zero counter and key.
+    const auto unit_float = [](uint32_t word) { return static_cast<float>(word >> 8) * 0x1.0p-24f; };
+    const std::vector<float> expected_first_counter = {
+        unit_float(0x6627e8d5U),
+        unit_float(0xe169c58dU),
+        unit_float(0xbc57ac4cU),
+        unit_float(0x9b00dbd8U),
+    };
+    EXPECT_EQ(CopyToCPUData(nn::function::Rand({4}, GetDevice(), MakeCUDAGenerator(0, 0))), expected_first_counter);
+
+    auto generator = MakeCUDAGenerator(0, 123);
+    EXPECT_EQ(CUDAOffset(generator->GetState()), 0U);
+    (void)nn::function::Rand({13}, GetDevice(), generator);
+    EXPECT_EQ(CUDAOffset(generator->GetState()), 13U);
+    (void)nn::function::Randn({7}, GetDevice(), generator);
+    EXPECT_EQ(CUDAOffset(generator->GetState()), 27U);
+}
+
+TEST_P(GeneratorDeviceTest, CUDAExplicitGeneratorMayTargetAnotherDeviceIndex) {
+    ONLY_CUDA();
+    REQUIRE_MIN_DEVICES(2);
+    const Device cuda1(Device::DeviceType::kCUDA, 1);
+    auto expected_generator = MakeCUDAGenerator(1, 123);
+    auto cross_index_generator = MakeCUDAGenerator(0, 123);
+
+    EXPECT_EQ(CopyToCPUData(nn::function::Rand({32}, cuda1, expected_generator)),
+              CopyToCPUData(nn::function::Rand({32}, cuda1, cross_index_generator)));
+
+    const auto state = cross_index_generator->GetState();
+    auto restored_on_cuda1 = MakeCUDAGenerator(1, 999);
+    restored_on_cuda1->SetState(state);
+    EXPECT_EQ(CopyToCPUData(nn::function::Rand({32}, cuda1, cross_index_generator)),
+              CopyToCPUData(nn::function::Rand({32}, cuda1, restored_on_cuda1)));
+}
+
+TEST_P(GeneratorDeviceTest, CUDADefaultGeneratorsUseGlobalSeedAndRemainIndependent) {
+    ONLY_CUDA();
+    REQUIRE_MIN_DEVICES(2);
+    const Device cuda0(Device::DeviceType::kCUDA, 0);
+    const Device cuda1(Device::DeviceType::kCUDA, 1);
+
+    ManualSeedAll(4242);
+    const auto cuda1_expected = CopyToCPUData(nn::function::Rand({32}, cuda1));
+    EXPECT_EQ(cuda1_expected, CopyToCPUData(nn::function::Rand({32}, cuda1, MakeCUDAGenerator(cuda1.index(), 4242))));
+    ManualSeedAll(4242);
+    (void)nn::function::Rand({32}, cuda0);
+    EXPECT_EQ(cuda1_expected, CopyToCPUData(nn::function::Rand({32}, cuda1)));
+}
+
+TEST_P(GeneratorDeviceTest, ConcurrentManualSeedAllKeepsDefaultBackendsConsistent) {
+    ONLY_CUDA();
+    constexpr int kNumThreads = 8;
+    constexpr int kSeedsPerThread = 128;
+    auto cpu_generator = GetDefaultCPUGenerator();
+    auto cuda_generator = GetDefaultCUDAGenerator(GetDevice().index());
+    std::vector<std::thread> workers;
+
+    for (int thread_rank = 0; thread_rank < kNumThreads; ++thread_rank) {
+        workers.emplace_back([thread_rank]() {
+            for (int i = 0; i < kSeedsPerThread; ++i) {
+                const uint64_t seed = (static_cast<uint64_t>(thread_rank + 1) << 32) | static_cast<uint64_t>(i);
+                ManualSeedAll(seed);
+            }
+        });
+    }
+    for (auto &worker : workers) { worker.join(); }
+
+    EXPECT_EQ(cpu_generator->InitialSeed(), cuda_generator->InitialSeed());
+    ManualSeedAll(2026);
+}
+
+} // namespace infini_train::test
diff --git a/tests/generator/generator_test_utils.h b/tests/generator/generator_test_utils.h
new file mode 100644
index 00000000..4f59e6e7
--- /dev/null
+++ b/tests/generator/generator_test_utils.h
@@ -0,0 +1,31 @@
+#pragma once
+
+#include <cstdint>
+#include <memory>
+#include <vector>
+
+#include "infini_train/include/core/runtime/device_guard.h"
+#include "infini_train/include/generator.h"
+#include "infini_train/include/tensor.h"
+#include "tests/common/test_utils.h"
+
+namespace infini_train::test {
+
+class GeneratorDeviceTest : public InfiniTrainTest {};
+
+inline std::shared_ptr<Generator> MakeGenerator(Device device, uint64_t seed) {
+    return device.IsCPU() ? MakeCPUGenerator(seed) : MakeCUDAGenerator(device.index(), seed);
+}
+
+inline std::vector<float> CopyToCPUData(const std::shared_ptr<Tensor> &tensor) {
+    auto cpu_tensor = tensor->To(Device());
+    auto *impl = core::GetDeviceGuardImpl(tensor->GetDevice().type());
+    impl->SynchronizeDevice(tensor->GetDevice());
+    if (cpu_tensor.NumElements() == 0) {
+        return {};
+    }
+    const auto *data = static_cast<const float *>(cpu_tensor.DataPtr());
+    return std::vector<float>(data, data + cpu_tensor.NumElements());
+}
+
+} // namespace infini_train::test
diff --git a/tests/generator/test_generator_api.cc b/tests/generator/test_generator_api.cc
new file mode 100644
index 00000000..1e91da4f
--- /dev/null
+++ b/tests/generator/test_generator_api.cc
@@ -0,0 +1,76 @@
+#include <memory>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/generator.h"
+#include "infini_train/include/nn/functional.h"
+#include "tests/generator/generator_test_utils.h"
+
+namespace infini_train::test {
+
+TEST_P(GeneratorDeviceTest, ExposesSeedStateAndDevice) {
+    const auto device = GetDevice();
+    auto generator = MakeGenerator(device, 123);
+
+    EXPECT_EQ(generator->InitialSeed(), 123U);
+    EXPECT_EQ(generator->GetDevice(), device);
+    EXPECT_FALSE(generator->GetState().empty());
+
+    (void)nn::function::Randn({1}, device, generator);
+    generator->ManualSeed(456);
+    EXPECT_EQ(generator->InitialSeed(), 456U);
+    EXPECT_EQ(CopyToCPUData(nn::function::Randn({7}, device, generator)),
+              CopyToCPUData(nn::function::Randn({7}, device, MakeGenerator(device, 456))));
+    const auto random_seed = generator->Seed();
+    EXPECT_EQ(random_seed, generator->InitialSeed());
+}
+
+TEST_P(GeneratorDeviceTest, SameSeedReplaysRandAndRandn) {
+    const auto device = GetDevice();
+    auto first = MakeGenerator(device, 123);
+    auto second = MakeGenerator(device, 123);
+    auto different = MakeGenerator(device, 456);
+
+    EXPECT_EQ(CopyToCPUData(nn::function::Rand({3, 5}, device, first)),
+              CopyToCPUData(nn::function::Rand({3, 5}, device, second)));
+    const auto first_normal = CopyToCPUData(nn::function::Randn({3, 5}, device, first));
+    EXPECT_EQ(first_normal, CopyToCPUData(nn::function::Randn({3, 5}, device, second)));
+    EXPECT_NE(first_normal, CopyToCPUData(nn::function::Randn({3, 5}, device, different)));
+}
+
+TEST_P(GeneratorDeviceTest, DefaultGeneratorAdvancesAndManualSeedReplays) {
+    const auto device = GetDevice();
+    ManualSeedAll(2026);
+    const auto first = CopyToCPUData(nn::function::Rand({8}, device));
+    const auto first_normal = CopyToCPUData(nn::function::Randn({8}, device));
+    const auto second = CopyToCPUData(nn::function::Rand({8}, device));
+    EXPECT_NE(first, second);
+
+    ManualSeedAll(2026);
+    EXPECT_EQ(first, CopyToCPUData(nn::function::Rand({8}, device)));
+    EXPECT_EQ(first_normal, CopyToCPUData(nn::function::Randn({8}, device)));
+    EXPECT_EQ(second, CopyToCPUData(nn::function::Rand({8}, device)));
+}
+
+TEST_P(GeneratorDeviceTest, ExplicitGeneratorDoesNotAdvanceDefaultGenerator) {
+    const auto device = GetDevice();
+    ManualSeedAll(777);
+    const auto expected = CopyToCPUData(nn::function::Rand({8}, device));
+
+    ManualSeedAll(777);
+    (void)nn::function::Rand({8}, device, MakeGenerator(device, 123));
+    EXPECT_EQ(expected, CopyToCPUData(nn::function::Rand({8}, device)));
+}
+
+TEST_P(GeneratorDeviceTest, DefaultGeneratorDispatchesByDevice) {
+    const auto device = GetDevice();
+    const auto generator = GetDefaultGenerator(device);
+    EXPECT_EQ(generator->GetDevice(), device);
+    EXPECT_EQ(generator, GetDefaultGenerator(device));
+    const auto backend_generator = device.IsCPU() ? GetDefaultCPUGenerator() : GetDefaultCUDAGenerator(device.index());
+    EXPECT_EQ(generator, backend_generator);
+}
+
+INFINI_TRAIN_REGISTER_TEST(GeneratorDeviceTest);
+
+} // namespace infini_train::test
diff --git a/tests/generator/test_generator_ddp.cc b/tests/generator/test_generator_ddp.cc
new file mode 100644
index 00000000..71006d07
--- /dev/null
+++ b/tests/generator/test_generator_ddp.cc
@@ -0,0 +1,170 @@
+#include <algorithm>
+#include <barrier>
+#include <memory>
+#include <string>
+#include <thread>
+#include <utility>
+#include <vector>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/device.h"
+#include "infini_train/include/generator.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"
+#include "tests/common/test_utils.h"
+#include "tests/generator/generator_test_utils.h"
+
+namespace infini_train::test {
+namespace {
+
+constexpr int kNumRanks = 2;
+constexpr char kRankBufferName[] = "rank_buffer";
+
+class StatefulLinear : public nn::Linear {
+public:
+    explicit StatefulLinear(float rank_value) : nn::Linear(4, 3, true, Device()) {
+        parameters_[nn::Linear::kParamBiasName]->set_requires_grad(false);
+        auto rank_buffer = std::make_shared<Tensor>(std::vector<int64_t>{5}, DataType::kFLOAT32, Device());
+        rank_buffer->Fill(rank_value);
+        buffers_[kRankBufferName] = std::move(rank_buffer);
+    }
+};
+
+std::vector<float> FlattenState(const std::shared_ptr<nn::Module> &module) {
+    auto state = module->StateDict();
+    std::vector<std::pair<std::string, std::shared_ptr<Tensor>>> named_tensors(state.begin(), state.end());
+    std::sort(named_tensors.begin(), named_tensors.end(),
+              [](const auto &lhs, const auto &rhs) { return lhs.first < rhs.first; });
+
+    std::vector<float> values;
+    for (const auto &[_, tensor] : named_tensors) {
+        auto tensor_values = CopyToCPUData(tensor);
+        values.insert(values.end(), tensor_values.begin(), tensor_values.end());
+    }
+    return values;
+}
+
+void EnsureDataParallelProcessGroup(const nn::parallel::Rank &rank, Device device) {
+    auto *factory = nn::parallel::ProcessGroupFactory::Instance(device.type());
+    factory->GetOrCreate(nn::parallel::GetDataParallelProcessGroupName(rank.GlobalRank()),
+                         nn::parallel::GetDataParallelGroupRanks(rank.GlobalRank()));
+}
+
+} // namespace
+
+TEST(GeneratorDDPTest, SynchronizesCPUInitializedStateBeforeFirstForward) {
+    REQUIRE_MIN_DEVICES(2);
+    ManualSeedAll(2026);
+    std::vector<std::vector<float>> pre_sync_state(kNumRanks);
+    std::vector<std::vector<float>> pre_sync_frozen_bias(kNumRanks);
+    std::vector<std::vector<float>> synchronized_state(kNumRanks);
+    std::vector<std::vector<float>> synchronized_buffers(kNumRanks);
+    std::vector<std::vector<float>> forward_outputs(kNumRanks);
+    std::vector<Tensor *> pre_sync_weights(kNumRanks);
+    std::vector<Tensor *> synchronized_weights(kNumRanks);
+    std::vector<std::shared_ptr<nn::Module>> synchronized_modules(kNumRanks);
+    std::barrier models_ready(kNumRanks);
+    std::vector<std::thread> threads;
+
+    for (int thread_rank = 0; thread_rank < kNumRanks; ++thread_rank) {
+        threads.emplace_back([&, thread_rank]() {
+            nn::parallel::Rank rank(0, thread_rank, 1, kNumRanks);
+            nn::parallel::global::thread_global_rank = rank.GlobalRank();
+            const Device device(Device::DeviceType::kCUDA, thread_rank);
+
+            auto module = std::make_shared<StatefulLinear>(static_cast<float>(thread_rank + 1));
+            module->To(device);
+            auto state = module->StateDict();
+            pre_sync_state[thread_rank] = FlattenState(module);
+            pre_sync_frozen_bias[thread_rank] = CopyToCPUData(state.at(nn::Linear::kParamBiasName));
+            pre_sync_weights[thread_rank] = state.at(nn::Linear::kParamWeightName).get();
+            models_ready.arrive_and_wait();
+
+            EnsureDataParallelProcessGroup(rank, device);
+
+            nn::parallel::DistributedDataParallelConfig config;
+            auto ddp = std::make_shared<nn::parallel::DistributedDataParallel>(module, rank, config);
+            state = module->StateDict();
+            synchronized_state[thread_rank] = FlattenState(module);
+            synchronized_buffers[thread_rank] = CopyToCPUData(state.at(kRankBufferName));
+            synchronized_weights[thread_rank] = state.at(nn::Linear::kParamWeightName).get();
+            synchronized_modules[thread_rank] = module;
+
+            auto input = std::make_shared<Tensor>(std::vector<int64_t>{2, 4}, DataType::kFLOAT32, device);
+            input->Fill(0.25f);
+            auto output = (*ddp)({input})[0];
+            forward_outputs[thread_rank] = CopyToCPUData(output);
+        });
+    }
+    for (auto &thread : threads) { thread.join(); }
+
+    ASSERT_FALSE(pre_sync_state[0].empty());
+    ASSERT_FALSE(pre_sync_state[1].empty());
+    EXPECT_NE(pre_sync_state[0], pre_sync_state[1]);
+    EXPECT_NE(pre_sync_frozen_bias[0], pre_sync_frozen_bias[1]);
+    for (int rank = 0; rank < kNumRanks; ++rank) {
+        EXPECT_EQ(synchronized_state[rank], pre_sync_state[0]);
+        EXPECT_EQ(synchronized_buffers[rank], std::vector<float>(5, 1.0f));
+        EXPECT_EQ(synchronized_weights[rank], pre_sync_weights[rank]);
+        EXPECT_FALSE(synchronized_modules[rank]->StateDict().at(nn::Linear::kParamBiasName)->requires_grad());
+    }
+    EXPECT_EQ(forward_outputs[0], forward_outputs[1]);
+}
+
+TEST(GeneratorDDPTest, DoesNotRebroadcastParametersAfterInitialSynchronization) {
+    REQUIRE_MIN_DEVICES(2);
+    ManualSeedAll(2027);
+    std::vector<std::vector<float>> post_forward_weights(kNumRanks);
+    std::vector<std::vector<float>> forward_outputs(kNumRanks);
+    std::barrier models_ready(kNumRanks);
+    std::barrier ddp_ready(kNumRanks);
+    std::barrier parameters_mutated(kNumRanks);
+    std::vector<std::thread> threads;
+
+    for (int thread_rank = 0; thread_rank < kNumRanks; ++thread_rank) {
+        threads.emplace_back([&, thread_rank]() {
+            nn::parallel::Rank rank(0, thread_rank, 1, kNumRanks);
+            nn::parallel::global::thread_global_rank = rank.GlobalRank();
+            const Device device(Device::DeviceType::kCUDA, thread_rank);
+
+            auto module = std::make_shared<StatefulLinear>(static_cast<float>(thread_rank + 1));
+            module->To(device);
+            models_ready.arrive_and_wait();
+            EnsureDataParallelProcessGroup(rank, device);
+
+            nn::parallel::DistributedDataParallelConfig config;
+            auto ddp = std::make_shared<nn::parallel::DistributedDataParallel>(module, rank, config);
+            ddp_ready.arrive_and_wait();
+
+            auto weight = module->StateDict().at(nn::Linear::kParamWeightName);
+            weight->Fill(static_cast<float>(thread_rank + 3));
+            parameters_mutated.arrive_and_wait();
+
+            auto input = std::make_shared<Tensor>(std::vector<int64_t>{2, 4}, DataType::kFLOAT32, device);
+            input->Fill(0.25f);
+            auto output = (*ddp)({input})[0];
+            forward_outputs[thread_rank] = CopyToCPUData(output);
+            post_forward_weights[thread_rank] = CopyToCPUData(weight);
+        });
+    }
+    for (auto &thread : threads) { thread.join(); }
+
+    EXPECT_EQ(post_forward_weights[0], std::vector<float>(12, 3.0f));
+    EXPECT_EQ(post_forward_weights[1], std::vector<float>(12, 4.0f));
+    EXPECT_NE(forward_outputs[0], forward_outputs[1]);
+}
+
+} // namespace infini_train::test
+
+int main(int argc, char **argv) {
+    ::testing::InitGoogleTest(&argc, argv);
+    infini_train::nn::parallel::global::GlobalEnv::Instance().Init(2, 1, false, 1, 1);
+    return RUN_ALL_TESTS();
+}
diff --git a/tests/generator/test_generator_random.cc b/tests/generator/test_generator_random.cc
new file mode 100644
index 00000000..a326fc93
--- /dev/null
+++ b/tests/generator/test_generator_random.cc
@@ -0,0 +1,160 @@
+#include <algorithm>
+#include <cmath>
+#include <limits>
+#include <numeric>
+#include <thread>
+#include <vector>
+
+#include "gtest/gtest.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/generator/generator_test_utils.h"
+
+namespace infini_train::test {
+
+TEST_P(GeneratorDeviceTest, FanCalculationUses64BitShapeProducts) {
+    constexpr int64_t kSpatialExtent = 65536;
+    auto tensor = std::make_shared<Tensor>(
+        std::vector<int64_t>{0, 3, kSpatialExtent, kSpatialExtent}, DataType::kFLOAT32, GetDevice());
+
+    const auto [fan_in, fan_out] = nn::init::CalculateFanInAndFanOut(tensor);
+    EXPECT_EQ(fan_in, int64_t{3} * kSpatialExtent * kSpatialExtent);
+    EXPECT_EQ(fan_out, 0);
+}
+
+TEST_P(GeneratorDeviceTest, UniformAndNormalRespectExplicitGenerator) {
+    const auto device = GetDevice();
+    auto first_uniform = std::make_shared<Tensor>(std::vector<int64_t>{32}, DataType::kFLOAT32, device);
+    auto second_uniform = std::make_shared<Tensor>(std::vector<int64_t>{32}, DataType::kFLOAT32, device);
+    nn::init::Uniform(first_uniform, -2.0f, 3.0f, MakeGenerator(device, 123));
+    nn::init::Uniform(second_uniform, -2.0f, 3.0f, MakeGenerator(device, 123));
+    EXPECT_EQ(CopyToCPUData(first_uniform), CopyToCPUData(second_uniform));
+
+    auto first_normal = std::make_shared<Tensor>(std::vector<int64_t>{32}, DataType::kFLOAT32, device);
+    auto second_normal = std::make_shared<Tensor>(std::vector<int64_t>{32}, DataType::kFLOAT32, device);
+    nn::init::Normal(first_normal, 1.5f, 0.5f, MakeGenerator(device, 456));
+    nn::init::Normal(second_normal, 1.5f, 0.5f, MakeGenerator(device, 456));
+    EXPECT_EQ(CopyToCPUData(first_normal), CopyToCPUData(second_normal));
+}
+
+TEST_P(GeneratorDeviceTest, GlobalSeedMakesKaimingUniformReproducible) {
+    const auto device = GetDevice();
+    auto first = std::make_shared<Tensor>(std::vector<int64_t>{16, 32}, DataType::kFLOAT32, device);
+    auto second = std::make_shared<Tensor>(std::vector<int64_t>{16, 32}, DataType::kFLOAT32, device);
+
+    ManualSeedAll(314159);
+    nn::init::KaimingUniform(first, std::sqrt(5.0f), nn::init::KaimingMode::kFanIn,
+                             nn::init::NonLinearityType::kLeakyReLU);
+    infini_train::ManualSeed(314159);
+    nn::init::KaimingUniform(second, std::sqrt(5.0f), nn::init::KaimingMode::kFanIn,
+                             nn::init::NonLinearityType::kLeakyReLU);
+    EXPECT_EQ(CopyToCPUData(first), CopyToCPUData(second));
+}
+
+TEST_P(GeneratorDeviceTest, ZeroElementCallsDoNotAdvanceState) {
+    const auto device = GetDevice();
+    auto generator = MakeGenerator(device, 123);
+    const auto before = generator->GetState();
+    (void)nn::function::Rand({0}, device, generator);
+    (void)nn::function::Randn({0}, device, generator);
+    EXPECT_EQ(before, generator->GetState());
+}
+
+TEST_P(GeneratorDeviceTest, ValidatesDistributionParameters) {
+    const auto device = GetDevice();
+    auto tensor = std::make_shared<Tensor>(std::vector<int64_t>{4}, DataType::kFLOAT32, device);
+    auto generator = MakeGenerator(device, 123);
+
+    EXPECT_DEATH(nn::init::Uniform(tensor, 2.0f, 1.0f, generator), "Check failed");
+    EXPECT_DEATH(nn::init::Uniform(tensor, 0.0f, std::numeric_limits<float>::infinity(), generator),
+                 "Uniform upper bound must be finite");
+    EXPECT_DEATH(
+        nn::init::Uniform(tensor, std::numeric_limits<float>::lowest(), std::numeric_limits<float>::max(), generator),
+        "Uniform bounds range exceeds float maximum");
+    EXPECT_DEATH(nn::init::Normal(tensor, 0.0f, -1.0f, generator), "Check failed");
+
+    auto wrong_backend = device.IsCPU() ? MakeCUDAGenerator(0, 123) : MakeCPUGenerator(123);
+    EXPECT_DEATH(nn::init::Uniform(tensor, 0.0f, 1.0f, wrong_backend), "Generator backend must match");
+}
+
+TEST_P(GeneratorDeviceTest, EqualUniformBoundsAndZeroStdReturnConstants) {
+    const auto device = GetDevice();
+    auto uniform = std::make_shared<Tensor>(std::vector<int64_t>{17}, DataType::kFLOAT32, device);
+    auto normal = std::make_shared<Tensor>(std::vector<int64_t>{17}, DataType::kFLOAT32, device);
+    auto uniform_generator = MakeGenerator(device, 123);
+    auto normal_generator = MakeGenerator(device, 123);
+    const auto uniform_state = uniform_generator->GetState();
+    const auto normal_state = normal_generator->GetState();
+
+    nn::init::Uniform(uniform, 2.5f, 2.5f, uniform_generator);
+    nn::init::Normal(normal, -3.0f, 0.0f, normal_generator);
+    EXPECT_EQ(CopyToCPUData(uniform), std::vector<float>(17, 2.5f));
+    EXPECT_EQ(CopyToCPUData(normal), std::vector<float>(17, -3.0f));
+    EXPECT_NE(uniform_state, uniform_generator->GetState());
+    EXPECT_NE(normal_state, normal_generator->GetState());
+}
+
+TEST_P(GeneratorDeviceTest, UniformNeverReturnsExclusiveUpperBoundAfterRounding) {
+    const auto device = GetDevice();
+    // These seeds hit the float-rounding edge in the current CPU and CUDA engines.
+    const uint64_t seed = device.IsCPU() ? 179040 : 16390493;
+    const int64_t size = device.IsCPU() ? 28 : 1;
+    auto tensor = std::make_shared<Tensor>(std::vector<int64_t>{size}, DataType::kFLOAT32, device);
+
+    nn::init::Uniform(tensor, 1.0f, 2.0f, MakeGenerator(device, seed));
+    const auto values = CopyToCPUData(tensor);
+    EXPECT_TRUE(std::all_of(values.begin(), values.end(), [](float value) { return value >= 1.0f && value < 2.0f; }));
+}
+
+TEST_P(GeneratorDeviceTest, ConcurrentCallsConsumeDisjointStreamSegments) {
+    constexpr int kNumThreads = 4;
+    constexpr int64_t kChunkSize = 257;
+    const auto device = GetDevice();
+
+    for (const bool normal : {false, true}) {
+        SCOPED_TRACE(normal ? "Randn" : "Rand");
+        auto generator = MakeGenerator(device, 123);
+        const auto random_values = [&](int64_t size, const auto &current_generator) {
+            auto tensor = normal ? nn::function::Randn({size}, device, current_generator)
+                                 : nn::function::Rand({size}, device, current_generator);
+            return CopyToCPUData(tensor);
+        };
+        std::vector<std::vector<float>> chunks(kNumThreads);
+        std::vector<std::thread> workers;
+
+        for (int i = 0; i < kNumThreads; ++i) {
+            workers.emplace_back([&, i]() { chunks[i] = random_values(kChunkSize, generator); });
+        }
+        for (auto &worker : workers) { worker.join(); }
+
+        std::vector<float> concurrent_values;
+        concurrent_values.reserve(kNumThreads * kChunkSize);
+        for (const auto &chunk : chunks) {
+            concurrent_values.insert(concurrent_values.end(), chunk.begin(), chunk.end());
+        }
+        auto serial_generator = MakeGenerator(device, 123);
+        auto serial_values = random_values(kNumThreads * kChunkSize, serial_generator);
+        std::sort(concurrent_values.begin(), concurrent_values.end());
+        std::sort(serial_values.begin(), serial_values.end());
+        EXPECT_EQ(concurrent_values, serial_values);
+        EXPECT_EQ(generator->GetState(), serial_generator->GetState());
+    }
+}
+
+TEST_P(GeneratorDeviceTest, RandomValuesHaveSaneRangeAndMoments) {
+    const auto device = GetDevice();
+    const auto uniform = CopyToCPUData(nn::function::Rand({65537}, device, MakeGenerator(device, 123)));
+    EXPECT_TRUE(std::all_of(uniform.begin(), uniform.end(), [](float value) { return value >= 0.0f && value < 1.0f; }));
+
+    const auto normal = CopyToCPUData(nn::function::Randn({65537}, device, MakeGenerator(device, 456)));
+    const double mean = std::accumulate(normal.begin(), normal.end(), 0.0) / normal.size();
+    const double squared_sum = std::inner_product(normal.begin(), normal.end(), normal.begin(), 0.0);
+    const double variance = squared_sum / normal.size() - mean * mean;
+    EXPECT_NEAR(mean, 0.0, 0.03);
+    EXPECT_NEAR(variance, 1.0, 0.06);
+}
+
+} // namespace infini_train::test
diff --git a/tests/generator/test_generator_state.cc b/tests/generator/test_generator_state.cc
new file mode 100644
index 00000000..9fb0900f
--- /dev/null
+++ b/tests/generator/test_generator_state.cc
@@ -0,0 +1,154 @@
+#include <cstdint>
+#include <locale>
+#include <memory>
+#include <string>
+#include <vector>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/generator.h"
+#include "infini_train/include/nn/functional.h"
+#include "tests/generator/generator_test_utils.h"
+
+namespace infini_train::test {
+namespace {
+
+std::vector<uint8_t> ReplaceStateHeader(const std::vector<uint8_t> &state, const std::string &header) {
+    std::string serialized(state.begin(), state.end());
+    const auto newline = serialized.find('\n');
+    CHECK_NE(newline, std::string::npos);
+    serialized.replace(0, newline, header);
+    return {serialized.begin(), serialized.end()};
+}
+
+class GroupedNumericPunct : public std::numpunct<char> {
+protected:
+    char do_thousands_sep() const override { return ','; }
+    std::string do_grouping() const override { return "\3"; }
+};
+
+class ScopedGroupedNumericLocale {
+public:
+    ScopedGroupedNumericLocale() : original_(std::locale()) {
+        std::locale::global(std::locale(std::locale::classic(), new GroupedNumericPunct));
+    }
+
+    ~ScopedGroupedNumericLocale() { std::locale::global(original_); }
+
+    ScopedGroupedNumericLocale(const ScopedGroupedNumericLocale &) = delete;
+    ScopedGroupedNumericLocale &operator=(const ScopedGroupedNumericLocale &) = delete;
+
+private:
+    std::locale original_;
+};
+
+} // namespace
+
+TEST_P(GeneratorDeviceTest, StateRestoreReplaysSequence) {
+    const auto device = GetDevice();
+    auto generator = MakeGenerator(device, 123);
+    (void)nn::function::Rand({7}, device, generator);
+    const auto state = generator->GetState();
+
+    const auto expected = CopyToCPUData(nn::function::Randn({11}, device, generator));
+    generator->ManualSeed(999);
+    generator->SetState(state);
+    EXPECT_EQ(generator->InitialSeed(), 123U);
+    EXPECT_EQ(expected, CopyToCPUData(nn::function::Randn({11}, device, generator)));
+}
+
+TEST_P(GeneratorDeviceTest, StateRoundTripIgnoresGlobalNumericLocale) {
+    const auto device = GetDevice();
+    auto generator = MakeGenerator(device, 1234567890123ULL);
+    (void)nn::function::Randn({1}, device, generator);
+
+    std::vector<uint8_t> state;
+    {
+        ScopedGroupedNumericLocale locale;
+        state = generator->GetState();
+    }
+    const auto expected = CopyToCPUData(nn::function::Randn({8}, device, generator));
+    {
+        ScopedGroupedNumericLocale locale;
+        generator->SetState(state);
+    }
+
+    EXPECT_EQ(expected, CopyToCPUData(nn::function::Randn({8}, device, generator)));
+}
+
+TEST_P(GeneratorDeviceTest, RandomStreamsSurviveCallPartitioningAndStateRestore) {
+    constexpr int64_t kTotalSize = 31;
+    const auto device = GetDevice();
+
+    for (const bool normal : {false, true}) {
+        SCOPED_TRACE(normal ? "normal" : "uniform");
+        const auto random_values = [&](int64_t size, const std::shared_ptr<Generator> &generator) {
+            auto tensor = normal ? nn::function::Randn({size}, device, generator)
+                                 : nn::function::Rand({size}, device, generator);
+            return CopyToCPUData(tensor);
+        };
+
+        for (const int64_t first_size : {1, 2, 3, 16, 17, 30}) {
+            SCOPED_TRACE(first_size);
+            auto whole_generator = MakeGenerator(device, 123);
+            auto split_generator = MakeGenerator(device, 123);
+
+            const auto whole = random_values(kTotalSize, whole_generator);
+            auto split = random_values(first_size, split_generator);
+            const auto cached_state = split_generator->GetState();
+            const auto tail = random_values(kTotalSize - first_size, split_generator);
+            split.insert(split.end(), tail.begin(), tail.end());
+
+            EXPECT_EQ(whole, split);
+            EXPECT_EQ(whole_generator->GetState(), split_generator->GetState());
+
+            split_generator->SetState(cached_state);
+            EXPECT_EQ(tail, random_values(kTotalSize - first_size, split_generator));
+        }
+    }
+}
+
+TEST_P(GeneratorDeviceTest, StateChangesAfterRandomConsumption) {
+    const auto device = GetDevice();
+    auto generator = MakeGenerator(device, 123);
+    const auto before = generator->GetState();
+    (void)nn::function::Rand({7}, device, generator);
+    EXPECT_NE(before, generator->GetState());
+}
+
+TEST_P(GeneratorDeviceTest, RejectsEmptyTruncatedAndWrongBackendState) {
+    const auto device = GetDevice();
+    auto generator = MakeGenerator(device, 123);
+    const auto state = generator->GetState();
+    ASSERT_GT(state.size(), 1U);
+
+    EXPECT_DEATH(generator->SetState({}), "Invalid .* generator state");
+    EXPECT_DEATH(generator->SetState(std::vector<uint8_t>(state.begin(), state.begin() + state.size() / 2)),
+                 "Invalid .* generator");
+    EXPECT_DEATH(generator->SetState(ReplaceStateHeader(state, "not a generator state")),
+                 "Invalid .* generator state header");
+
+    auto with_trailing_data = state;
+    with_trailing_data.push_back('x');
+    EXPECT_DEATH(generator->SetState(with_trailing_data), "Invalid .* generator");
+}
+
+TEST_P(GeneratorDeviceTest, RejectsStateFromDifferentBackend) {
+    const auto device = GetDevice();
+    auto generator = MakeGenerator(device, 123);
+    auto other = device.IsCPU() ? MakeCUDAGenerator(0, 123) : MakeCPUGenerator(123);
+    EXPECT_DEATH(generator->SetState(other->GetState()), "Invalid .* generator state header");
+}
+
+TEST_P(GeneratorDeviceTest, UsesAllSeedBits) {
+    const auto device = GetDevice();
+    constexpr uint64_t low_seed = 7;
+    constexpr uint64_t high_seed = (uint64_t{1} << 63) | low_seed;
+    auto low = MakeGenerator(device, low_seed);
+    auto high = MakeGenerator(device, high_seed);
+
+    EXPECT_NE(CopyToCPUData(nn::function::Rand({16}, device, low)),
+              CopyToCPUData(nn::function::Rand({16}, device, high)));
+}
+
+} // namespace infini_train::test