cuRANDDx

https://docs.nvidia.com/cuda/curanddx/
Admin
cuRANDDx (cuRAND Device Extensions) is a CUDA C++ library that enables high-performance random...
Tokens:29,586
Snippets:202
Trust Score:10
Update:1 month ago
Show doc for...
Tokens:
Raw
### Generate random numbers with scrambled Sobol64

Source: https://docs.nvidia.com/cuda/curanddx/examples/sobol.html

This example demonstrates the kernel implementation and host-side setup for generating random numbers using the scrambled Sobol64 generator. It includes memory allocation for direction vectors and constants required for the generator.

```cpp
/*
 * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */ 

#include <iostream>
#include <vector>
#include <cassert>

#include <curanddx.hpp>
#include "../common.hpp"

// This example demonstrates how to use quasi-random generator scrambled sobol64 and cuRANDDx thread-level operator
// to generate a sequence of random numbers and compare with the results generated using cuRAND host API

template<class RNG, typename DataType>
__global__ void generate_kernel(DataType*                                 d_out,
                                typename RNG::direction_vector_type*      direction_vectors,
                                const typename RNG::scrambled_const_type* scrambled_consts,
                                const typename RNG::offset_type           offset,
                                const unsigned int                        num_dims,
                                const size_t                              size,
                                DataType                                  input1,
                                DataType                                  input2) {
    CURANDDX_SKIP_IF_NOT_APPLICABLE_SM(RNG);

    int       tid     = blockDim.x * blockIdx.x + threadIdx.x;
    const int threads = blockDim.x * gridDim.x;

    curanddx::uniform<float> my_uniform(input1, input2);

    // When generating n results in m dimensions, the output will consist of n//m results from dimension 0, followed by
    // another n/m results from dimension 1, etc.
    const int size_per_dim = size / num_dims;

    for (auto idx = tid; idx < size; idx += threads) {
        // for each dimension, tid0 generates the first number, tid1 the second number etc.
        const int idex_per_dim = idx % size_per_dim;
        const int dim          = idx / size_per_dim;
        RNG       rng(dim, direction_vectors, offset + idex_per_dim, scrambled_consts);


        d_out[idx] = my_uniform.generate(rng);
    }
}

template<unsigned int Arch>
int sobol_thread_api() {
    using RNG =
        decltype(curanddx::Generator<curanddx::scrambled_sobol64>() + curanddx::SM<Arch>() + curanddx::Thread());

    using DataType = float;

    // Allocate output memory
    DataType*    d_out;
    const size_t size = 5000;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&d_out, size * sizeof(DataType)));

    // size has to be evenly divisible by dimensions
    const unsigned int sobol_dims = 200;
    assert(size % sobol_dims == 0);

    const typename RNG::offset_type offset = 10ULL;
    const DataType                  min_v  = -1;
    const DataType                  max_v  = 1;

    // Step 1: Get pointers to the host direction vectors using cuRAND API
    using direction_vector_type = typename RNG::direction_vector_type;
    direction_vector_type* h_direction_vectors;
    CURAND_CHECK_AND_EXIT(
        curandGetDirectionVectors64(&h_direction_vectors, CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6));

    // Step 2: Allocate memory and copy one direction vector per dimension to the device
    direction_vector_type* direction_vectors_gmem_ptr;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&(direction_vectors_gmem_ptr), sobol_dims * sizeof(direction_vector_type)));

    CUDA_CHECK_AND_EXIT(cudaMemcpy(direction_vectors_gmem_ptr,
                                   h_direction_vectors,
                                   sobol_dims * sizeof(direction_vector_type),
                                   cudaMemcpyHostToDevice));

    // Step 3: Get pointers to the host scrambled constants
    using scrambled_const_type = typename RNG::scrambled_const_type;
    scrambled_const_type* h_scrambled_consts;
    scrambled_const_type* scrambled_consts_gmem_ptr;
    CURAND_CHECK_AND_EXIT(curandGetScrambleConstants64(&h_scrambled_consts));
```

--------------------------------

### Main function to run the example

Source: https://docs.nvidia.com/cuda/curanddx/examples/xorwow_init_and_generate.html

The main function calls the example runner with the functor to execute the xorwow_init_and_generate_thread_api example.

```cpp
int main(int, char**) {
    return example::sm_runner<xorwow_init_and_generate_thread_api_functor>();
}
```

--------------------------------

### Simple PCG example with hash comparison

Source: https://docs.nvidia.com/cuda/curanddx/examples/simple_pcg.html

This C++ snippet implements a simple PCG example using cuRANDDx. It generates random numbers, computes a hash of the output, and compares it to a known hash value (0xaa706742) to verify that the same sequence is generated with NVPL RAND and cuRANDDx using STRICT ordering. The code includes a functor and a main function that runs the example on supported architectures. The output indicates success if the hash matches, otherwise it prints a failure message and returns -1.

```cpp
#else
    // Compare hash
    unsigned int xor_curand = 0x0;
    for (auto i = 0U; i < size; i++) {
        xor_curand ^= h_out[i];
    }
    if (xor_curand == 0xaa706742) {
        std::cout << "Compared to the hash value: Same sequence is generated with NVPL RAND and cuRANDDx generator "
                     "using STRICT ordering.\n";
        std::cout << "SUCCESS \n";
        return 0;
    } else {
        std::cout
            << "FAILED: different sequence is generated with NVPL RAND and cuRANDDx generator using STRICT ordering.\n";
        return -1;
    }
#endif
}

template<unsigned int Arch>
struct simple_pcg_thread_api_functor {
    int operator()() { return simple_pcg_thread_api<Arch>(); }
};

int main(int, char**) {
    return example::sm_runner<simple_pcg_thread_api_functor>();
}
```

--------------------------------

### Using Description Traits in User Code

Source: https://docs.nvidia.com/cuda/curanddx/api/description_traits.html

Example demonstrating how to use description traits to inspect an RNG configuration.

```cpp
#include <curanddx.hpp>

using RNG = decltype(curanddx::Generator<curanddx::philox4_32>()
                   + curanddx::PhiloxRounds<7>()
                   + curanddx::SM<800>());

if(curanddx::is_complete_rand<RNG>::value) {
  if (curanddx::generator_of_v<RNG> == curanddx::philox4_32) {
    std::cout << "Use Philox generator with " << curanddx::philox_rounds_of_v<RNG> << " rounds \n";
  }
}
```

--------------------------------

### XORWOW init and generate kernels with host setup

Source: https://docs.nvidia.com/cuda/curanddx/examples/xorwow_init_and_generate.html

Shows how to initialize XORWOW generator states in one kernel and generate normally distributed random numbers in a later kernel using the Box-Muller method. The init_kernel uses rng.init(seed, subsequence, offset) or the constructor, and the generate_kernel uses dist.generate2(rng) and updates the states in global memory. The host code allocates memory, launches both kernels, and verifies results against the cuRAND host API with CURAND_ORDERING_PSEUDO_LEGACY ordering.

```cpp
/*
 * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */ 

#include <iostream>
#include <vector>

#include <curanddx.hpp>
#include "../common.hpp"

// This example demonstrates how to launch two kernels, one for initializing the states, one for generating normally distribute 
// random numbers with Box-Muller method, using XORWOW generator. 
// 
// cuRANDDx functions used in the example kernels:
// (1) state initialization using init() function, or the constructor of the RNG object
// (2) generate2() function as Box-Muller method is selected for normal distribution
//
// The generated RNs are the same as the results using cuRAND host API with CURAND_ORDERING_PSEUDO_LEGACY ordering
// The ordering requires that the result at offset n is from position  
// (n mod 4096) * 2^67 + (n/4096) in the original XORWOW sequence
// https://docs.nvidia.com/cuda/curand/host-api-overview.html#generator-types 

constexpr unsigned int subsequences = 4096;

template<class RNG>
__global__ void init_kernel(RNG* states, const unsigned long long seed, const typename RNG::offset_type offset) {
    CURANDDX_SKIP_IF_NOT_APPLICABLE_SM(RNG);

    int tid = blockDim.x * blockIdx.x + threadIdx.x;

    RNG rng;
    rng.init(seed, ((offset + tid) % subsequences), ((offset + tid) / subsequences)); // seed, subsequence, offset

    // Alternatively one can use the constructor directly
    //RNG rng(seed, ((offset + tid) % subsequences), ((offset + tid) / subsequences)); 
    
    states[tid] = rng;
}

template<class RNG, typename DataType>
__global__ void generate_kernel(RNG* states, float2* d_out, const size_t size, const DataType mean, const DataType stddev) {
    CURANDDX_SKIP_IF_NOT_APPLICABLE_SM(RNG);

    int       tid     = blockDim.x * blockIdx.x + threadIdx.x;
    const int threads = blockDim.x * gridDim.x;

    curanddx::normal<DataType, curanddx::box_muller> dist(mean, stddev);

    RNG rng = states[tid];

    for (auto idx = tid; idx < size / 2; idx += threads) {
        d_out[idx] = dist.generate2(rng);
    }

    // Each thread updates "states" in global memory so "states" can be used for later kernels
    states[tid] = rng;
}


template<unsigned int Arch>
int xorwow_init_and_generate_thread_api() {
    using RNG      = decltype(curanddx::Generator<curanddx::xorwow>() + curanddx::SM<Arch>() + curanddx::Thread());
    using DataType = float;

    // Allocate output memory
    DataType*    d_out;
    const size_t size = 50000;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&d_out, size * sizeof(DataType)));

    const unsigned long long seed   = 1234ULL;
    const typename RNG::offset_type offset = 2ULL;

    const unsigned int block_dim = 256;
    const unsigned int grid_size = 16;

    const DataType mean = 0; 
    const DataType stddev = 2;

    // Allocate an array of states
    RNG* states;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&states, block_dim * grid_size * sizeof(RNG)));

    // Invoke the init kernel first to set up the states
    init_kernel<RNG><<<grid_size, block_dim>>>(states, seed, offset);
    CUDA_CHECK_AND_EXIT(cudaPeekAtLastError());

    // Invoke the generate kernel to generate RNs and update the states
    generate_kernel<RNG, DataType><<<grid_size, block_dim>>>(states, (float2*)d_out, size, mean, stddev);
    CUDA_CHECK_AND_EXIT(cudaPeekAtLastError());
    CUDA_CHECK_AND_EXIT(cudaDeviceSynchronize());

    std::vector<DataType> h_out(size);
    CUDA_CHECK_AND_EXIT(cudaMemcpy(h_out.data(), d_out, size * sizeof(DataType), cudaMemcpyDeviceToHost));

    // cuRAND host API
    curandGenerator_t gen_curand;
    DataType*         d_ref;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&d_ref, size * sizeof(DataType)));
}
```

--------------------------------

### Main NVRTC cuRANDDx example: kernel definition and runtime compilation

Source: https://docs.nvidia.com/cuda/curanddx/examples/nvrtc_pcg.html

Includes the NVRTC cuDNN example's kernel source, compilation program, options, and error handling. It creates a cuRANDDx operator-based kernel that generates random numbers, compiles it with NVRTC, and retrieves the cubin. Note: the snippet ends before the kernel execution and validation part; the missing file would use the compiled cubin.

```cpp
/*
 * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

#include <iostream>
#include <vector>
#include <memory>

#include <nvrtc.h>
#include <cuda.h>
#include <cuda_runtime_api.h>

#include "nvrtc_helper.hpp" 
#include "../common.hpp"

const char* generate_kernel = R"kernel(
#include <curanddx.hpp>

// CURANDDX Operators 

using RNG = decltype(curanddx::Generator<curanddx::pcg>() + curanddx::SM<RNG_SM>() + curanddx::Thread());

extern "C" __global__ void generate_kernel(unsigned int*        d_out,
                                const unsigned long long        seed,
                                const typename RNG::offset_type offset,
                                const size_t                    size) {
    const int tid = blockDim.x * blockIdx.x + threadIdx.x;
    if (tid >= size)
        return;

    curanddx::uniform_bits<unsigned int> dist;

    // compare with NVPL RAND PCG strict ordering
    RNG rng(seed, 0, offset + tid);

    d_out[tid] = dist.generate(rng);

}
)kernel";


// This example demonstrates how to use NVRTC to runtime compile a kernel using CURANDDX functions and check correctness

int main(int, char**) {

    // Get current device
    int current_device;
    CUDA_CHECK_AND_EXIT(cudaGetDevice(&current_device));

    // Create NVRTC program
    nvrtcProgram program;
    NVRTC_SAFE_CALL(nvrtcCreateProgram(&program,           // program
                                       generate_kernel,      // buffer
                                       "generate_kernel.cu", // name
                                       0,                  // numHeaders
                                       NULL,               // headers
                                       NULL));             // includeNames

    // Prepare compilation options
    std::vector<const char*> opts = {
        "--std=c++17",
        "--device-as-default-execution-space",
        "--include-path=" CUDA_INCLUDE_DIR, // Add path to CUDA include directory
        "--include-path=" CUDA_CCCL_INCLUDE_DIR // Add path to CCCL include directory for CTK 13.0
    };

    // Parse cuRANDDx include dirs
    std::vector<std::string> curanddx_include_dirs = example::nvrtc::get_curanddx_include_dirs();
    // Add cuBLASDx include dirs to opts
    for (auto& d : curanddx_include_dirs) {
        opts.push_back(d.c_str());
    }

    // Add GPU_ARCHITECTURE definition to opts
    std::string gpu_architecture_definition =
        "-DRNG_SM=" + std::to_string(example::nvrtc::get_device_architecture(current_device) * 10);
    opts.push_back(gpu_architecture_definition.c_str());

    // Add gpu-architecture to opts
    std::string gpu_architecture_option = example::nvrtc::get_device_architecture_option(current_device);
    opts.push_back(gpu_architecture_option.c_str());

    nvrtcResult compileResult = nvrtcCompileProgram(program,                       // program
                                                    static_cast<int>(opts.size()), // numOptions
                                                    opts.data());                  // options

    // Obtain compilation log from the program
    if (compileResult != NVRTC_SUCCESS) {
        for (auto o : opts) {
            std::cout << o << std::endl;
        }
        example::nvrtc::print_program_log(program);
        std::exit(1);
    }

    // Obtain cubin from the program.
    size_t cubin_size;
    NVRTC_SAFE_CALL(nvrtcGetCUBINSize(program, &cubin_size));
    auto cubin = std::make_unique<char[]>(cubin_size);
    NVRTC_SAFE_CALL(nvrtcGetCUBIN(program, cubin.get()));

    // Destroy the program.
    NVRTC_SAFE_CALL(nvrtcDestroyProgram(&program))
```

--------------------------------

### PCG example: hash comparison and main runner

Source: https://docs.nvidia.com/cuda/curanddx/examples/pcg.html

The first block computes a hash of the reference output when NVPL RAND is unavailable, then compares the cuRANDDx output hash to a known value (0xaa706742) to verify STRICT ordering. The second block defines a functor that calls the PCG thread API and runs it via example::sm_runner in main. Use together to test PCG generation with STRICT ordering.

```cpp
    // compute hash to be used if nvpl rand is not available
    unsigned int xor_nvplrand = 0x0;
    for (auto i = 0U; i < size; i++) {
        xor_nvplrand ^= h_ref[i];
    }
    std::cout << "NVPL RAND reference xor output is " << std::hex << xor_nvplrand << std::endl;
    return 0;

#else
    // Compare hash
    unsigned int xor_curand = 0x0;
    for (auto i = 0U; i < size; i++) {
        xor_curand ^= h_out[i];
    }
    if (xor_curand == 0xaa706742) {
        std::cout << "Compared to the hash value: Same sequence is generated with NVPL RAND and cuRANDDx generator "
                     "using STRICT ordering.\n";
        std::cout << "SUCCESS \n";
        return 0;
    } else {
        std::cout
            << "FAILED: different sequence is generated with NVPL RAND and cuRANDDx generator using STRICT ordering.\n";
        return 1;
    }
#endif
}
```

```cpp
template<unsigned int Arch>
struct simple_pcg_thread_api_functor {
    int operator()() { return simple_pcg_thread_api<Arch>(); }
};

int main(int, char**) {
    return example::sm_runner<simple_pcg_thread_api_functor>();
}
```

--------------------------------

### Simple PCG example with cuRANDDx and NVPL RAND comparison

Source: https://docs.nvidia.com/cuda/curanddx/examples/simple_pcg.html

Complete CUDA C++ example that generates a sequence of 32-bit random numbers using the cuRANDDx PCG generator with the Thread-level operator and compares the output with NVPL RAND using strict ordering. The kernel uses a uniform_bits distribution and a per-thread offset. If NVPL RAND is available, it creates a PCG generator with strict ordering, generates a reference sequence, and compares the results, printing a success or failure message. It also computes and prints an XOR hash of the reference output.

```cpp
// Copyright (c) 2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
//
// NVIDIA CORPORATION and its licensors retain all intellectual property
// and proprietary rights in and to this software, related documentation
// and any modifications thereto.  Any use, reproduction, disclosure or
// distribution of this software and related documentation without an express
// license agreement from NVIDIA CORPORATION is strictly prohibited.

#include <iostream>
#include <vector>

#include <curanddx.hpp>
#include "common.hpp"

#ifdef CURANDDX_EXAMPLE_NVPL_RAND_AVAILABLE
#    include <nvpl_rand.h>
#    define NVPL_RAND_CHECK(x)                                  \
        do {                                                    \
            if ((x) != NVPL_RAND_STATUS_SUCCESS) {              \
                printf("Error at %s:%d\n", __FILE__, __LINE__); \
                _Exit(EXIT_FAILURE);                            \
            }                                                   \
        } while (0)
#endif

// This example demonstrates how to use the default PCG generator and cuRANDDx Thread-level operator to generate a sequence of random 32-bit 
// numbers and compare with the results generated using NVPL RAND, if available, with STRICT ordering

template<class RNG, typename data_type>
__global__ void generate_kernel(data_type*                      d_out,
                                const unsigned long long        seed,
                                const typename RNG::offset_type offset,
                                const size_t                    size) {
    const int tid = blockDim.x * blockIdx.x + threadIdx.x;
    if (tid >= size)
        return;

    curanddx::uniform_bits<data_type> dist;

    // compare with NVPL RAND PCG strict ordering
    RNG rng(seed, 0, offset + tid);

    d_out[tid] = dist.generate(rng);
}

template<unsigned int Arch>
int simple_pcg_thread_api() {
    using RNG = decltype(curanddx::Generator<curanddx::pcg>() + curanddx::SM<Arch>() + curanddx::Thread());

    using DataType = typename RNG::bitgenerator_result_type;

    // Allocate output memory
    DataType*    d_out;
    const size_t size = 5000;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&d_out, size * sizeof(DataType)));

    unsigned long long seed   = 1234ULL;
    unsigned long long offset = 1ULL;

    // Invokes kernel
    const unsigned int block_dim = 256;
    const unsigned int grid_size = (size + block_dim - 1) / block_dim;
    
    generate_kernel<RNG, DataType><<<grid_size, block_dim, 0>>>(d_out, seed, offset, size);
    CUDA_CHECK_AND_EXIT(cudaPeekAtLastError());
    CUDA_CHECK_AND_EXIT(cudaDeviceSynchronize());

    std::vector<DataType> h_out(size);
    CUDA_CHECK_AND_EXIT(cudaMemcpy(h_out.data(), d_out, size * sizeof(DataType), cudaMemcpyDeviceToHost));
    CUDA_CHECK_AND_EXIT(cudaFree(d_out));

#ifdef CURANDDX_EXAMPLE_NVPL_RAND_AVAILABLE
    // nvpl RAND
    std::vector<DataType> h_ref(size);
    nvplRandGenerator_t   gen;
    const auto            generator_type = NVPL_RAND_RNG_PSEUDO_PCG;

    nvplRandMTCreateGeneratorDefault(&gen, generator_type);
    NVPL_RAND_CHECK(nvplRandSetPseudoRandomGeneratorSeed(gen, seed));
    NVPL_RAND_CHECK(nvplRandSetGeneratorOffset(gen, offset));
    NVPL_RAND_CHECK(nvplRandMTSetGeneratorOrdering(gen, NVPL_RAND_ORDERING_STRICT));

    // Generate
    NVPL_RAND_CHECK(nvplRandGenerate(gen, h_ref.data(), size));
    NVPL_RAND_CHECK(nvplRandDestroyGenerator(gen));

    // Compare Results
    if (h_out == h_ref) {
        std::cout << "SUCCESS: Same sequence is generated with cuRANDDx and nvpl RAND API using STRICT ordering.\n";
    } else {
        int count {0};
        for (auto i = 0U; i < size; i++) {
            if (h_out[i] != h_ref[i] && count < 10) {
                printf("array_curanddx[%u] = %u, array_nvplrand[%u] = %u \n", i, h_out[i], i, h_ref[i]);
                count++;
            }
        }
        std::cout << "FAILED: Different sequence is generated with cuRANDDx and NVPL RAND Host API using LEGACY "
                     "ordering.\n";
        return 1;
    }

    // compute hash to be used if nvpl rand is not available
    unsigned int xor_nvplrand = 0x0;
    for (auto i = 0U; i < size; i++) {
        xor_nvplrand ^= h_ref[i];
    }
    std::cout << "NVPL RAND reference xor output is " << std::hex << xor_nvplrand << std::endl;
    return 0;

```

--------------------------------

### Philox Random Bits Thread API Example

Source: https://docs.nvidia.com/cuda/curanddx/examples/philox_random_bits.html

This CUDA C++ example uses the cuRANDDx thread-level operator with the Philox generator to generate random bits. It defines a kernel that generates 4 random values per thread, then compares the output with the cuRAND host API using CURAND_ORDERING_PSEUDO_LEGACY ordering. The example uses a seed of 1234 and an offset of 1, and prints SUCCESS or FAILED based on the comparison.

```cpp
/*
 * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */ 

#include <iostream>
#include <vector>

#include <curanddx.hpp>
#include "../common.hpp"

// This example demonstrates how to use cuRANDDx thread-level operator to generate a sequence of random numbers
// using Philox generator with user-defined round number (default 10). 

// The generated randombits values are the same as the results using cuRAND host API 
// with CURAND_ORDERING_PSEUDO_LEGACY ordering. 
// The ordering is that using subsequences different sequences each thread generates 4 32-bit random values, 
// and each four values from one sequence are followed by four values from next sequence.
// https://docs.nvidia.com/cuda/curand/host-api-overview.html#generator-types Philox_4x32_10 pseudorandom generator

constexpr unsigned int subsequences = 65536;

template<class RNG, typename DataType>
__global__ void generate_kernel(uint4*                        d_out,
                                const unsigned long long        seed,
                                const typename RNG::offset_type offset,
                                const size_t                    size) {
    CURANDDX_SKIP_IF_NOT_APPLICABLE_SM(RNG);

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i >= size / 4)
        return;

    RNG rng(seed, ((offset + i) % subsequences), ((offset + i) / subsequences)); // seed, subsequence, offset

    d_out[i] = rng.generate4();
}

template<unsigned int Arch>
int philox_random_bits_thread_api() {
    using RNG      = decltype(curanddx::Generator<curanddx::philox4_32>() + curanddx::PhiloxRounds<10>() +
                         curanddx::SM<Arch>() + curanddx::Thread());
    using DataType = unsigned int;

    // Allocate output memory
    DataType*    d_out;
    const size_t size = 5000;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&d_out, size * sizeof(DataType)));

    const unsigned long long        seed   = 1234ULL;
    const typename RNG::offset_type offset = 1ULL;

    // Invokes kernel
    const unsigned int block_dim = 256;
    const unsigned int grid_size = (size / 4 + block_dim - 1) / block_dim;
    generate_kernel<RNG, DataType><<<grid_size, block_dim, 0>>>((uint4*)d_out, seed, offset, size);
    CUDA_CHECK_AND_EXIT(cudaPeekAtLastError());
    CUDA_CHECK_AND_EXIT(cudaDeviceSynchronize());

    std::vector<DataType> h_out(size);
    CUDA_CHECK_AND_EXIT(cudaMemcpy(h_out.data(), d_out, size * sizeof(DataType), cudaMemcpyDeviceToHost));

    // cuRAND host API
    curandGenerator_t gen_curand;
    DataType*         d_ref;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&d_ref, size * sizeof(DataType)));

    CURAND_CHECK_AND_EXIT(curandCreateGenerator(&gen_curand, CURAND_RNG_PSEUDO_PHILOX4_32_10));
    CURAND_CHECK_AND_EXIT(curandSetPseudoRandomGeneratorSeed(gen_curand, seed));
    CURAND_CHECK_AND_EXIT(curandSetGeneratorOffset(gen_curand, offset * 4));
    CURAND_CHECK_AND_EXIT(curandSetGeneratorOrdering(gen_curand, CURAND_ORDERING_PSEUDO_LEGACY));

    CURAND_CHECK_AND_EXIT(curandGenerate(gen_curand, d_ref, size));

    std::vector<DataType> h_ref(size);
    CUDA_CHECK_AND_EXIT(cudaMemcpy(h_ref.data(), d_ref, size * sizeof(DataType), cudaMemcpyDeviceToHost));

    CURAND_CHECK_AND_EXIT(curandDestroyGenerator(gen_curand));
    CUDA_CHECK_AND_EXIT(cudaFree(d_out));
    CUDA_CHECK_AND_EXIT(cudaFree(d_ref));

    // Compare Results
    if (h_out == h_ref) {
        std::cout << "SUCCESS: Same sequence is generated with cuRANDDx and cuRAND Host API using LEGACY ordering.\n";
        return 0;
    } else {
        int count {0};
        for (auto i = 0U; i < size; i++) {
            if (h_out[i] != h_ref[i] && count++ < 10) {
                std::cout << "array_curanddx[" << i << "] = " << h_out[i] << " array_curand[" << i << "] = " << h_ref[i]
                          << std::endl;
                ;
            }
        }
        std::cout << "FAILED: Different sequence is generated with cuRANDDx and cuRAND Host API using LEGACY "
                     "ordering.\n";
        return 1;
    }
}

template<unsigned int Arch>
struct philox_random_bits_thread_api_functor {
    int operator()() { return philox_random_bits_thread_api<Arch>(); }
};

int main(int, char**) {
    return example::sm_runner<philox_random_bits_thread_api_functor>();
}

```

--------------------------------

### Define functor to run xorwow_init_and_generate_thread_api

Source: https://docs.nvidia.com/cuda/curanddx/examples/xorwow_init_and_generate.html

Defines a functor that calls the xorwow_init_and_generate_thread_api function with a given architecture. This is used to run the example on a specific GPU architecture.

```cpp
template<unsigned int Arch>
struct xorwow_init_and_generate_thread_api_functor {
    int operator()() { return xorwow_init_and_generate_thread_api<Arch>(); }
};
```

--------------------------------

### Generating random numbers in a CUDA kernel

Source: https://docs.nvidia.com/cuda/curanddx/get_started/introduction.html

This example shows how to define an RNG type, initialize the state within a kernel, specify a normal distribution, and generate random numbers.

```cpp
#include <curanddx.hpp>

using RNG = decltype(curanddx::Generator<curanddx::philox4_32>() +
                     curanddx::SM<800>() +
                     curanddx::Thread());

__global__ void generate_kernel(double4*                        d_out,
                                const unsigned long long        seed,
                                const typename RNG::offset_type offset) {
  const int i = blockDim.x * blockIdx.x + threadIdx.x;

  // Initialize state with seed, subsequence, offset
  RNG rng(seed, ((offset + i) % 65536), ((offset + i) / 65536));

  // Specify the distribution and the method used, either box-muller or icdf
  curanddx::normal<double, curanddx::box_muller> dist(0, 1);


  // Generate the random numbers and copy the output to global memory
  d_out[i] = dist.generate4(rng);
}
```

--------------------------------

### Philox Thread-Level Random Number Generation with cuRANDDx

Source: https://docs.nvidia.com/cuda/curanddx/examples/introduction_philox.html

This CUDA C++ example demonstrates using cuRANDDx's thread-level operator to generate log-normally distributed random numbers with the Philox generator. It uses a user-defined round number (default 10) and compares results with the cuRAND host API using CURAND_ORDERING_PSEUDO_LEGACY ordering. The kernel generates 4 values per thread, and the example includes setup, kernel launch, and verification steps.

```cpp
/*
 * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */ 

#include <iostream>
#include <vector>

#include <curanddx.hpp>
#include "../common.hpp"

// This example demonstrates how to use cuRANDDx thread-level operator to generate a sequence of random numbers
// using Philox generator with user-defined round number (default 10). 

// The generated log-normally distributed values are the same as the results using cuRAND host API 
// with CURAND_ORDERING_PSEUDO_LEGACY ordering. 
// The ordering is that using subsequences different sequences each thread generates 4 32-bit random values, 
// and each four values from one sequence are followed by four values from next sequence.
// https://docs.nvidia.com/cuda/curand/host-api-overview.html#generator-types Philox_4x32_10 pseudorandom generator

constexpr unsigned int subsequences = 65536;

template<class RNG, typename DataType>
__global__ void generate_kernel(double4*                        d_out,
                                const unsigned long long        seed,
                                const typename RNG::offset_type offset,
                                const size_t                    size,
                                const DataType                  mean,
                                const DataType                  stddev) {
    CURANDDX_SKIP_IF_NOT_APPLICABLE_SM(RNG);

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i >= size / 4)
        return;

    curanddx::log_normal<DataType, curanddx::box_muller> dist(mean, stddev);

    RNG rng(seed, ((offset + i) % subsequences), ((offset + i) / subsequences)); // seed, subsequence, offset

    d_out[i] = dist.generate4(rng);
}

template<unsigned int Arch>
int philox_thread_api() {
    using RNG      = decltype(curanddx::Generator<curanddx::philox4_32>() + curanddx::PhiloxRounds<10>() +
                         curanddx::SM<Arch>() + curanddx::Thread());
    using DataType = double;

    // Allocate output memory
    DataType*    d_out;
    const size_t size = 5000;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&d_out, size * sizeof(DataType)));

    const unsigned long long        seed   = 1234ULL;
    const typename RNG::offset_type offset = 1ULL;

    const DataType mean   = 1;
    const DataType stddev = 0.5;

    // Invokes kernel
    const unsigned int block_dim = 256;
    const unsigned int grid_size = (size / 4 + block_dim - 1) / block_dim;
    generate_kernel<RNG, DataType><<<grid_size, block_dim, 0>>>((double4*)d_out, seed, offset, size, mean, stddev);
    CUDA_CHECK_AND_EXIT(cudaPeekAtLastError());
    CUDA_CHECK_AND_EXIT(cudaDeviceSynchronize());

    std::vector<DataType> h_out(size);
    CUDA_CHECK_AND_EXIT(cudaMemcpy(h_out.data(), d_out, size * sizeof(DataType), cudaMemcpyDeviceToHost));

    // cuRAND host API
    curandGenerator_t gen_curand;
    DataType*         d_ref;
    CUDA_CHECK_AND_EXIT(cudaMalloc((void**)&d_ref, size * sizeof(DataType)));

    CURAND_CHECK_AND_EXIT(curandCreateGenerator(&gen_curand, CURAND_RNG_PSEUDO_PHILOX4_32_10));
    CURAND_CHECK_AND_EXIT(curandSetPseudoRandomGeneratorSeed(gen_curand, seed));
    CURAND_CHECK_AND_EXIT(curandSetGeneratorOffset(gen_curand, offset * 4));
    CURAND_CHECK_AND_EXIT(curandSetGeneratorOrdering(gen_curand, CURAND_ORDERING_PSEUDO_LEGACY));

    CURAND_CHECK_AND_EXIT(curandGenerateLogNormalDouble(gen_curand, d_ref, size, mean, stddev));

    std::vector<DataType> h_ref(size);
    CUDA_CHECK_AND_EXIT(cudaMemcpy(h_ref.data(), d_ref, size * sizeof(DataType), cudaMemcpyDeviceToHost));

    CURAND_CHECK_AND_EXIT(curandDestroyGenerator(gen_curand));
    CUDA_CHECK_AND_EXIT(cudaFree(d_out));
    CUDA_CHECK_AND_EXIT(cudaFree(d_ref));
}
```

--------------------------------

### Load cubin and launch generate_kernel with CUDA driver API

Source: https://docs.nvidia.com/cuda/curanddx/examples/nvrtc_pcg.html

Loads a cubin, launches the generate_kernel with a grid of 256 threads per block, copies results to host, and compares the XOR hash against 0xaa706742 to verify STRICT ordering matches NVPL RAND. Requires CUDA driver API and a previously compiled cubin. Prints SUCCESS or FAILED and returns 0 or 1 accordingly.

```cpp
    // Load the generated Cubin and get a handle to the generate_kernel
    CUdevice   cuDevice;
    CUcontext  context;
    CUmodule   module;
    CUfunction kernel;
    CU_CHECK_AND_EXIT(cuInit(0));
    CU_CHECK_AND_EXIT(cuDeviceGet(&cuDevice, current_device));
#if CUDA_VERSION >= 13000
    CU_CHECK_AND_EXIT(cuCtxCreate(&context, (CUctxCreateParams*)0, 0, cuDevice));
#else
    CU_CHECK_AND_EXIT(cuCtxCreate(&context, 0, cuDevice));
#endif
    CU_CHECK_AND_EXIT(cuModuleLoadDataEx(&module, cubin.get(), 0, 0, 0));
    CU_CHECK_AND_EXIT(cuModuleGetFunction(&kernel, module, "generate_kernel"));

    // Allocate output memory
    using DataType = unsigned int;

    CUdeviceptr   d_out;
    size_t size = 5000;
    CU_CHECK_AND_EXIT(cuMemAlloc(&d_out, size * sizeof(DataType)));

    unsigned long long seed   = 1234ULL;
    unsigned long long offset = 1ULL;

    // Invokes kernel
    const unsigned int block_dim = 256;
    const unsigned int grid_size = (size + block_dim - 1) / block_dim;

    // Execute generate_kernel
    void* args[] = {&d_out, &seed, &offset, &size};
    CU_CHECK_AND_EXIT(cuLaunchKernel(kernel,
                                     grid_size, // number of blocks
                                     1,
                                     1,
                                     block_dim, // number of threads
                                     1,
                                     1,
                                     0,    // shared memory size
                                     NULL, // NULL stream
                                     args,
                                     0));
    CU_CHECK_AND_EXIT(cuCtxSynchronize());

    // Copy to host
    std::vector<DataType> h_out(size);
    CU_CHECK_AND_EXIT(cuMemcpyDtoH(h_out.data(), d_out, size * sizeof(DataType)));
    CU_CHECK_AND_EXIT(cuMemFree(d_out));

    CU_CHECK_AND_EXIT(cuModuleUnload(module));
    CU_CHECK_AND_EXIT(cuCtxDestroy(context));

    // Compare hash
    unsigned int xor_curand = 0x0;
    for (auto i = 0U; i < size; i++) {
        xor_curand ^= h_out[i];
    }
    if (xor_curand == 0xaa706742) {
        std::cout << "Compared to the hash value: Same sequence is generated with NVPL RAND and cuRANDDx generator "
                     "using STRICT ordering.\n";
        std::cout << "SUCCESS \n";
        return 0;
    } else {
        std::cout
            << "FAILED: different sequence is generated with NVPL RAND and cuRANDDx generator using STRICT ordering.\n";
        return 1;
    }
}
```

--------------------------------

### Compile with nvcc using MathDx include directory

Source: https://docs.nvidia.com/cuda/curanddx/get_started/installation.html

Use this command to compile a CUDA source file with cuRANDDx. Replace <mathdx_include_dir> with the path to the MathDx include directory, <your_source_file>.cu with your source file, and <your_binary> with the output binary. Requires C++17 and an appropriate architecture flag.

```bash
nvcc -std=c++17 -arch sm_XY (...) -I<mathdx_include_dir> <your_source_file>.cu -o <your_binary>

```

--------------------------------

### Using execution traits in user code

Source: https://docs.nvidia.com/cuda/curanddx/api/execution_traits.html

Example of retrieving execution traits from a descriptor built with Thread operator. The descriptor uses philox4_32 generator, PhiloxRounds<7>, and SM<800>. The offset_type trait is accessed via typename RNG::offset_type.

```cpp
#include <curanddx.hpp>

using RNG = decltype(curanddx::Generator<curanddx::philox4_32>()
                   + curanddx::PhiloxRounds<7>()
                   + curanddx::SM<800>());

using offset_type = typename RNG::offset_type;

```

--------------------------------

### Create a thread-level execution object with curanddx::Thread()

Source: https://docs.nvidia.com/cuda/curanddx/api/execution_ops.html

Use this snippet to create a thread-level execution object for a cuRANDDx function descriptor. It sets the random number generation to run in a thread context, where each thread independently creates a random number in its register. The example combines the Thread operator with a Generator and SM description operator to form a complete function descriptor.

```cpp
curanddx::Thread()

```

```cpp
#include <curanddx.hpp>

using RNG = decltype(curanddx::Generator<curanddx::pcg>() +
                     curanddx::SM<750>() +
                     curanddx::Thread());

```

--------------------------------

### State Initialization for Thread Execution (init and RNG constructor)

Source: https://docs.nvidia.com/cuda/curanddx/api/methods.html

Shows the explicit init() and implicit thread execution constructor signatures for initializing RNG state. Pseudorandom generators use seed, subsequence, and offset; quasirandom generators require dimension, direction vectors, offset, and optional scrambled constants.

```cpp
using RNG = decltype(curanddx::Thread() + ...);

// Explicit init state function for pseudorandom generators
void init(const unsigned long long seed,
          const unsigned long long subsequence,
          const typename RNG::offset_type offset);

// Implicit init state function using thread execution constructor for pseudorandom generators
void RNG(const unsigned long long seed,
         const unsigned long long subsequence,
         const typename RNG::offset_type offset);

// Explicit init function for quasirandom generators
void init(const unsigned int dim,
          typename RNG::direction_vector_type*      direction_vectors,
          const typename RNG::offset_type           offset,
          const typename RNG::scrambled_const_type* scrambled_consts = nullptr);

// Implicit init function using thread execution constructor for quasirandom generators
void RNG(const unsigned int dim,
         typename RNG::direction_vector_type*      direction_vectors,
         const typename RNG::offset_type           offset,
         const typename RNG::scrambled_const_type* scrambled_consts = nullptr);
```

--------------------------------

### Generate 32-bit Random Numbers (generate and generate4)

Source: https://docs.nvidia.com/cuda/curanddx/api/methods.html

Declares the generate() and generate4() functions for producing 32-bit random numbers. generate() works for all generators except Philox; generate4() is only for Philox. Using the wrong function for a generator causes compilation errors.

```cpp
using RNG = decltype(curanddx::Thread() + ...);

// Generate a single 32-bit pseudo-random or quasi-random number. Supported for all generators except Philox.
unsigned int generate();

// Generate 4 32-bit pseudo-random numbers. Only supported for Philox generator.
uint4 generate4();
```

--------------------------------

### Compare cuRANDDx and cuRAND Host API results with Philox

Source: https://docs.nvidia.com/cuda/curanddx/examples/introduction_philox.html

Compares the output arrays from cuRANDDx and cuRAND Host API for the Philox generator. Prints SUCCESS if the sequences match, otherwise prints FAILED and shows up to 10 mismatched indices. Uses LEGACY ordering.

```cpp
    // Compare Results
    if (h_out == h_ref) {
        std::cout << "SUCCESS: Same sequence is generated with cuRANDDx and cuRAND Host API using LEGACY ordering.\n";
        return 0;
    } else {
        int count {0};
        for (auto i = 0U; i < size; i++) {
            if (h_out[i] != h_ref[i] && count++ < 10) {
                std::cout << "array_curanddx[" << i << "] = " << h_out[i] << " array_curand[" << i << "] = " << h_ref[i]
                          << std::endl;
                ;
            }
        }
        std::cout << "FAILED: Different sequence is generated with cuRANDDx and cuRAND Host API using LEGACY "
                     "ordering.\n";
        return 1;
    }
}
```

```cpp
template<unsigned int Arch>
struct philox_thread_api_functor {
    int operator()() { return philox_thread_api<Arch>(); }
};
```

```cpp
int main(int, char**) {
    return example::sm_runner<philox_thread_api_functor>();
}
```
Tokens:
Raw