Repository navigation
[libcu++] Experimental Multiply-With-Carry random engines #12590
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Open
fbusato
wants to merge
12
commits into
NVIDIA:main
Choose a base branch
from
fbusato:mwc-random-engines
base: main
Could not load branches
Branch not found: {{ refName }}
Loading
Could not load tags
Nothing to show
Loading
Are you sure you want to change the base?
Some commits from the old base branch may be removed from the timeline,
and old review comments may become outdated.
Open
Changes from all commits
Commits
Show all changes
12 commits
Select commit
Hold shift + click to select a range
5d58f27
implementation
fbusato 80188f1
add documentation
fbusato 2090f39
remove excessive unrolling
fbusato faba9c3
unit test
fbusato bcf1f27
add performance considerations
fbusato 2d0847e
add microbenchmark
fbusato 412efee
add codegen test
fbusato fad1a23
Merge branch 'main' of https://github.com/NVIDIA/cccl into mwc-random…
fbusato 3d42122
Update libcudacxx/include/cuda/__random/mwc_engine.h
fbusato ad015d5
Update libcudacxx/include/cuda/__random/mwc_engine.h
fbusato 27eae3f
address review comments
fbusato a4cefab
fix runtime tests
fbusato File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,73 @@ | ||
| .. _libcudacxx-extended-api-random-mwc: | ||
|
|
||
| ``mwc64x32`` and ``mwc128x32`` | ||
| ============================== | ||
|
|
||
| Defined in the ``<cuda/random>`` header as **experimental** feature in the namespace ``cuda::experimental``. | ||
|
|
||
| These scrambled *Multiply-With-Carry (MWC)* engines implement the `RandomNumberEngine <https://en.cppreference.com/w/cpp/named_req/RandomNumberEngine>`_ interface and produce unsigned 32-bit values over the entire range ``[0, 2^32 - 1]``. | ||
| They use the **MWC64X/32** and **MWC128X/32** algorithms from `Scrambled Multiply-with-Carry Pseudorandom Number Generators <https://arxiv.org/abs/2610.27313>`_. | ||
|
|
||
| .. note:: | ||
|
|
||
| These engines are not cryptographic generators and are not recommended for simulation or scientific computing requiring independently validated streams. | ||
|
|
||
| .. list-table:: | ||
| :header-rows: 1 | ||
|
|
||
| * - Engine | ||
| - State size | ||
| - Multiplier | ||
| - Period | ||
| * - ``cuda::experimental::mwc64x32`` | ||
| - 64 bits | ||
| - ``0xFFF145F8`` | ||
| - ``(0xFFF145F8 * 2^32 - 2) / 2`` (approximately ``2^63``) | ||
| * - ``cuda::experimental::mwc128x32`` | ||
| - 128 bits | ||
| - ``0xFFCCCC66`` | ||
| - ``(0xFFCCCC66 * 2^96 - 2) / 2`` (approximately ``2^127``) | ||
|
|
||
| Both engines support host and device execution, constant evaluation, and logarithmic-time ``discard(unsigned long long)``. | ||
|
|
||
| Performance Considerations | ||
| -------------------------- | ||
|
|
||
| Both engines have similar throughput. The following table shows approximate performance ratio compared to other engines: | ||
|
|
||
| .. list-table:: | ||
| :header-rows: 1 | ||
|
|
||
| * - Compared with | ||
| - ``MWC64X/32``, ``MWC128X/32`` | ||
| * - cuRAND XORWOW | ||
| - 1.4x | ||
| * - ``cuda::pcg64`` | ||
| - 2.6x | ||
| * - cuRAND Philox4x32-10 | ||
| - 3x | ||
| * - ``cuda::std::philox4x32`` | ||
| - 20x | ||
|
|
||
| Seeding | ||
| ------- | ||
|
|
||
| Both engines can be initialized or reseeded with a 64-bit unsigned integer, including zero, or a seed sequence. The default seed is ``1``. | ||
|
|
||
| Parallel streams | ||
| ---------------- | ||
|
|
||
| Randomly seeding many ``mwc64x32`` engines can lead to overlapping streams because its period is relatively short. The user should use ``discard`` instead to partition a sequence into non-overlapping bounded streams. See the example below: | ||
|
|
||
| .. code:: cuda | ||
|
|
||
| #include <cuda/random> | ||
|
|
||
| __global__ void sample(unsigned* output) { | ||
| unsigned long long tid = blockIdx.x * blockDim.x + threadIdx.x; | ||
| cuda::experimental::mwc128x32 engine(42); | ||
| engine.discard(tid * 1024); // Reserve 1024 outputs per thread | ||
| output[tid] = engine(); | ||
| } | ||
|
|
||
| For performance reasons, each stream should be initialized once, then reuse the engine multiple times. Jumping ahead is more expensive than generating a value. | ||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,90 @@ | ||
| //===----------------------------------------------------------------------===// | ||
| // | ||
| // Part of libcu++, the C++ Standard Library for your entire system, | ||
| // under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #include <thrust/device_vector.h> | ||
|
|
||
| #include <cuda/random> | ||
| #include <cuda/std/cstddef> | ||
| #include <cuda/std/cstdint> | ||
| #include <cuda/std/limits> | ||
| #include <cuda/std/utility> | ||
|
|
||
| #include <cuda_runtime_api.h> | ||
| #include <device_side_benchmark.cuh> | ||
|
|
||
| #include <nvbench/nvbench.cuh> | ||
|
|
||
| using cuda::experimental::mwc128x32; | ||
| using cuda::experimental::mwc64x32; | ||
|
|
||
| NVBENCH_DECLARE_TYPE_STRINGS(mwc64x32, "mwc64x32", "cuda::experimental::mwc64x32"); | ||
| NVBENCH_DECLARE_TYPE_STRINGS(mwc128x32, "mwc128x32", "cuda::experimental::mwc128x32"); | ||
|
|
||
| using rng_types = nvbench::type_list<mwc64x32, mwc128x32>; | ||
|
|
||
| template <class Engine> | ||
| __global__ void initialize(Engine* states, const uint64_t seed) | ||
| { | ||
| const auto tid = static_cast<size_t>(blockIdx.x) * blockDim.x + threadIdx.x; | ||
| states[tid] = Engine{seed + tid}; | ||
| } | ||
|
|
||
| template <class Engine> | ||
| __global__ void generate(const Engine* initial_states, const uint32_t samples_per_thread) | ||
| { | ||
| const auto tid = static_cast<size_t>(blockIdx.x) * blockDim.x + threadIdx.x; | ||
| auto engine = initial_states[tid]; | ||
| uint32_t checksum = 0; | ||
| for (uint32_t i = 0; i < samples_per_thread; ++i) | ||
| { | ||
| checksum += engine(); | ||
| } | ||
| // Keep the scrambler and every state transition observable without per-sample stores. | ||
| sink(cuda::std::pair{engine, checksum}); | ||
| } | ||
|
|
||
| template <class Engine> | ||
| void generation(nvbench::state& state, nvbench::type_list<Engine>) | ||
| { | ||
| const auto requested_samples = state.get_int64("SamplesPerThread"); | ||
| const auto samples_per_thread = static_cast<uint32_t>(requested_samples); | ||
| const auto block_size = static_cast<int>(state.get_int64("BlockSize")); | ||
| int blocks_per_sm = 0; | ||
| NVBENCH_CUDA_CALL_NOEXCEPT( | ||
| cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks_per_sm, generate<Engine>, block_size, 0)); | ||
| if (blocks_per_sm == 0) | ||
| { | ||
| state.skip("Insufficient resources for the requested block size"); | ||
| return; | ||
| } | ||
| const int num_sms = state.get_device().value().get_number_of_sms(); // NOLINT(bugprone-unchecked-optional-access) | ||
| // Launch one resident wave for this engine and block size. | ||
| const int grid_size = blocks_per_sm * num_sms; | ||
| const auto threads = static_cast<size_t>(grid_size) * block_size; | ||
| thrust::device_vector<Engine> initial_states(threads); | ||
| auto* const initial = thrust::raw_pointer_cast(initial_states.data()); | ||
|
|
||
| // Seed once outside timing; every sample starts from the same per-thread state. | ||
| initialize<<<grid_size, block_size>>>(initial, 42); | ||
| NVBENCH_CUDA_CALL_NOEXCEPT(cudaGetLastError()); | ||
| NVBENCH_CUDA_CALL_NOEXCEPT(cudaDeviceSynchronize()); | ||
|
|
||
| state.add_element_count(static_cast<uint64_t>(threads) * samples_per_thread); | ||
| state.add_global_memory_reads<Engine>(threads); | ||
| state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch& launch) { | ||
| generate<<<grid_size, block_size, 0, launch.get_stream()>>>(initial, samples_per_thread); | ||
| }); | ||
| } | ||
|
|
||
| NVBENCH_BENCH_TYPES(generation, NVBENCH_TYPE_AXES(rng_types)) | ||
| .set_name("generation") | ||
| .set_type_axes_names({"Engine{ct}"}) | ||
| .add_int64_axis("BlockSize", {128, 256, 512, 1024}) | ||
| .add_int64_power_of_two_axis("SamplesPerThread", {16}); |
Oops, something went wrong.
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Uh oh!
There was an error while loading. Please reload this page.