[fbusato]

Description

Closes #12287

The PR provides cuda::experimental::mwc64x32 and cuda::experimental::mwc128x32 to <cuda/random>, implementing Sebastiano Vigna’s scrambled Multiply-With-Carry generators.

The implementation follow the RandomNumberEngine interface, support host/device execution, constant evaluation, state comparisons, and stream serialization.

The PR includes:

• Documentation + performance consideration
• Unit test
• NVBench benchmarks
• Codegen test

[coderabbitai]

No actionable comments were generated in the recent review.

Summary

Summary by CodeRabbit

• New Features
  • Added experimental mwc64x32 and mwc128x32 random-number engines for host and device use. Both produce 32-bit values, support seeding, and can efficiently skip ahead in generated sequences. Host stream serialization is also supported.
• Documentation
  • Added guidance on engine characteristics, seeding, performance, and partitioning random-number streams. The engines are listed as available with CCCL 3.7.0 and CUDA Toolkit 13.7.

The PR adds experimental mwc64x32 and mwc128x32 random engines to <cuda/random>. It adds engine tests, documentation, benchmarks, and SASS code-generation checks, with test and CI configuration for the new checks.

Changes: MWC engines

Layer / File(s)	Summary
Engine implementation and export libcudacxx/include/cuda/__random/mwc_engine.h, libcudacxx/include/cuda/random	The engine implements seeding, output generation, comparison, hosted stream operations, and discard. The header defines mwc64x32 and mwc128x32, and cuda/random includes it.
Engine behavior tests libcudacxx/test/libcudacxx/cuda/random/mwc.pass.cpp, libcudacxx/test/libcudacxx/cuda/random/mwc_io.pass.cpp	Tests cover engine properties, seeding, output bounds, discard equivalence, reference outputs, and hosted stream operations.
Engine documentation and benchmark docs/libcudacxx/extended_api/random.rst, docs/libcudacxx/extended_api/random/mwc.rst, libcudacxx/benchmarks/bench/random/mwc.cu	The Random API index links to the new page. The page describes engine properties, seeding, performance comparisons, and stream partitioning with discard. The benchmark measures both engines across selected block sizes and sample counts.
SASS code-generation checks libcudacxx/test/random_codegen/*, libcudacxx/test/CMakeLists.txt, ci/matrix.yaml	The code-generation tests check instructions for one-step calls to both engines. CMake and CI configuration register the random SASS suite.

Assessment against linked issues

Objective	Addressed	Explanation
Add mwc64x32 and mwc128x32 to <cuda/random> with random-engine operations, including seeding, generation, discard, comparison, and stream operators [#12287]	✅

Suggested reviewers: pciolkosz, griwes

Change: Feature

Merge Risk: Low · up to a4cef

The documented kernel example can compute a wrapped thread index for launches over 2^32 threads, so users who copy it would reuse streams and overwrite output. This is a narrow documentation fix and does not affect the engines themselves. It is worth correcting before merge but is not a blocker.

Security architecture risk: Minimal · up to a4cef

The new engines are explicitly noncryptographic and maintain caller-owned state. Review found no introduced security boundary bypass or privilege expansion. Invalid serialized states are rejected without changing the engine.

Retained concerns
No architecture-level concerns identified.

Security review details

Security Blast Radius

• inferred — Within the inspected implementation, caller-controlled seeds and serialized words affect the supplied engine's state and outputs, rather than an identity, authorization decision, credential store, or service-wide state. Any security-sensitive downstream use would require a separate assessment.

Trust Boundaries and Controls

• observed — The routed main and runtime_test functions belong to the behavioral test executable. Main ignores command-line arguments, while runtime_test invokes engine checks with fixed seeds and reference vectors; these functions do not introduce a production input boundary.

Resilience and Maintainability Implications

• observed — Failed parsing or invalid state validation cannot partially replace the engine: assignment occurs only after all words are read and validated. A seed-sequence exception before generate returns similarly precedes engine assignment. These controls contain failure within the caller's operation rather than requiring rollback of shared state.

- [ ] Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

Recent review info

Run configuration

• Source mode: Auto
• Repository: NVIDIA/cccl
• Branch: main
• File: .coderabbit.yaml
• Commit: 22bbd61
• Layers: Repository YAML
• Review profile: CHILL
• Plan: Enterprise
• Run ID: 6497c1e7-7ee1-43a2-b0a8-c1aa17decdd8

Commits

Reviewing files that changed from the base of the PR and between 27eae3f and a4cefab.

Files selected for processing: 1

• libcudacxx/test/libcudacxx/cuda/random/mwc.pass.cpp

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.

[coderabbitai]

Actionable comments posted: 1

Nitpick comments: 1

libcudacxx/test/libcudacxx/cuda/random/mwc.pass.cpp: 1

162-162: Trivial · Maintainability & Code Quality · Low value

suggestion: Remove const from the local arrays on Lines 162 and 189. The test guidelines discourage const-qualification.

Source: Coding guidelines

Review info

Run configuration

• Source mode: Auto
• Repository: NVIDIA/cccl
• Branch: main
• File: .coderabbit.yaml
• Commit: 22bbd61
• Layers: Repository YAML
• Review profile: CHILL
• Plan: Enterprise
• Run ID: 034d9d3c-f5fc-4989-b0ed-810d8cdfab33

Commits

Reviewing files that changed from the base of the PR and between 22bbd61 and fad1a23.

Files selected for processing: 11

• ci/matrix.yaml
• docs/libcudacxx/extended_api/random.rst
• docs/libcudacxx/extended_api/random/mwc.rst
• libcudacxx/benchmarks/bench/random/mwc.cu
• libcudacxx/include/cuda/__random/mwc_engine.h
• libcudacxx/include/cuda/random
• libcudacxx/test/CMakeLists.txt
• libcudacxx/test/libcudacxx/cuda/random/mwc.pass.cpp
• libcudacxx/test/libcudacxx/cuda/random/mwc_io.pass.cpp
• libcudacxx/test/random_codegen/CMakeLists.txt
• libcudacxx/test/random_codegen/mwc_single_step.cu

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.

[davebayer]

Why isn't this in cudax?

[miscco]

Looks mostly good to me I would like to get @RAMitchell to look at it

[miscco]

I believe we should make those [[nodiscard]], because there is an explicit discard method

[fbusato]

I wrote it in this way mostly for consistency. @RAMitchell what do you think?

[RAMitchell]

There are legitimate reasons to discard (e.g. use every k'th engine output). I leave it up to you both to use nodiscard or not - consistency is probably good.

[fbusato]

Why isn't this in cudax?

Mostly for keeping all non-standard utilities close together. Also, the reason to expose it as experiemental is weak.
@miscco what do you think?

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 2h 37m: Pass: 100%/514 | Total: 4d 16h | Max: 1h 50m | Hits: 99%/926104

See results here.

[RAMitchell]

I ran some brute force fuzzing tests (hundreds of millions of outputs) and they all passed.

Tested seeding, saving and restoring state, and skipping ahead, including very large skip counts.

I (read, codex) used ESBMC and Z3 for formal verification. Examples of statements successfully checked include:

• Generating the next number: for every valid starting state, the returned number and every word of the updated state match the paper.
• Modular addition: for every pair of reduced operands, the helper returns the correct sum with wraparound at the engine’s modulus. This includes cases where the intermediate sum exceeds the storage width.
• Seed-sequence initialization: for every possible set of generated seed words, the state digits are preserved and the carry is placed within the intended safe range.
• Short skips: for every valid starting state and every skip count from 0 through 15, discard(n) leaves exactly the same state as generating and ignoring n numbers.
  For larger skips, the reasoning was split smaller proofs. These establish that the multiplication algorithm computes the correct modular product, and that the skip-ahead algorithm applies exactly the requested number of transitions. The mathematical argument covers every 64-bit skip count, rather than just the values tested.

There is a mistake in the docs:

the example calculates blockIdx.x * blockDim.x using 32-bit arithmetic before assigning it to a 64-bit variable. Large launches can therefore wrap around and give different threads the same index. Widen an operand before multiplying:

auto tid = static_cast<unsigned long long>(blockIdx.x) * blockDim.x
         + threadIdx.x;