cuRANDDx#

API reference: cuRANDDx C++ documentation.

Operators#

Operators are set with curanddxSetOperatorInt64():

CURANDDX_OPERATOR_GENERATOR — CURANDDX_GENERATOR_PCG, CURANDDX_GENERATOR_PHILOX4_32, CURANDDX_GENERATOR_XORWOW, CURANDDX_GENERATOR_MRG32K3A, CURANDDX_GENERATOR_SOBOL32, CURANDDX_GENERATOR_SCRAMBLED_SOBOL32, CURANDDX_GENERATOR_SOBOL64, or CURANDDX_GENERATOR_SCRAMBLED_SOBOL64.
CURANDDX_OPERATOR_PHILOX_ROUNDS — number of Philox rounds when the generator is CURANDDX_GENERATOR_PHILOX4_32 (optional, default 10).
CURANDDX_OPERATOR_SM — target SM.
CURANDDX_OPERATOR_EXECUTION — COMMONDX_EXECUTION_THREAD only.
CURANDDX_OPERATOR_DISTRIBUTION — CURANDDX_DISTRIBUTION_UNIFORM, CURANDDX_DISTRIBUTION_NORMAL, CURANDDX_DISTRIBUTION_LOG_NORMAL, CURANDDX_DISTRIBUTION_POISSON, or CURANDDX_DISTRIBUTION_UNIFORM_BITS.
CURANDDX_OPERATOR_OUTPUT_TYPE — e.g. COMMONDX_R_32F.
CURANDDX_OPERATOR_GENERATE_METHOD — CURANDDX_GENERATE_METHOD_SINGLE, CURANDDX_GENERATE_METHOD_PAIR, or CURANDDX_GENERATE_METHOD_QUAD.
CURANDDX_OPERATOR_NORMAL_METHOD — CURANDDX_NORMAL_METHOD_ICDF or CURANDDX_NORMAL_METHOD_BOX_MULLER (optional, for CURANDDX_DISTRIBUTION_NORMAL and CURANDDX_DISTRIBUTION_LOG_NORMAL distributions).
CURANDDX_OPERATOR_DEVICE_FUNCTIONS — bitmask of CURANDDX_DEVICE_FUNCTION_GENERATE, CURANDDX_DEVICE_FUNCTION_INIT_STATE, CURANDDX_DEVICE_FUNCTION_DESTROY_STATE, CURANDDX_DEVICE_FUNCTION_SKIP_OFFSET, CURANDDX_DEVICE_FUNCTION_SKIP_SUBSEQUENCE, and CURANDDX_DEVICE_FUNCTION_SKIP_SEQUENCE.

Distribution parameters (e.g. min/max for uniform, mean/stddev for normal) are passed at call time as an argument to the generated generate device function, not as an operator.

Traits#

curanddxGetTraitStrSize() + curanddxGetTraitStr() — CURANDDX_TRAIT_SYMBOL_GENERATE_NAME, CURANDDX_TRAIT_SYMBOL_INIT_STATE_NAME, CURANDDX_TRAIT_SYMBOL_DESTROY_STATE_NAME, CURANDDX_TRAIT_SYMBOL_SKIP_OFFSET_NAME, CURANDDX_TRAIT_SYMBOL_SKIP_SUBSEQUENCE_NAME, CURANDDX_TRAIT_SYMBOL_SKIP_SEQUENCE_NAME.
curanddxGetTraitInt64() — CURANDDX_TRAIT_STATE_SIZE, CURANDDX_TRAIT_STATE_ALIGNMENT.

Device usage#

Allocate state with CURANDDX_TRAIT_STATE_SIZE and CURANDDX_TRAIT_STATE_ALIGNMENT.
Call init_state to initialize the generator state.
- Pseudorandom: void init_state(unsigned long long* seed, unsigned long long* subsequence, offset_type* offset, void* state)
- Quasirandom: void init_state(unsigned int* dim, direction_vector_type* direction_vector, offset_type* offset, scrambled_const_type* scrambled_consts, void* state)
Call generate to produce random values:

void generate(void* state, output_type* out, param_type* params)

params contains the distribution parameters (e.g. two floats for min/max in uniform, mean/stddev in normal); unused for CURANDDX_DISTRIBUTION_UNIFORM_BITS.
Optionally call skip functions to advance the state:
- void skip_offset(void* state, offset_type* n)
- void skip_subsequence(void* state, unsigned long long* n)
- void skip_sequence(void* state, unsigned long long* n)
Call destroy_state to clean up:

void destroy_state(void* state)

Example#

cuRANDDx example#

/*
 * 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 <cuda.h>
#include <cuda_runtime.h>
#include <libcuranddx.h>
#include <nvJitLink.h>
#include <nvrtc.h>

#include <array>
#include <cassert>
#include <cstdio>
#include <cstdlib>
#include <iostream>
#include <string>
#include <vector>

#include "common_examples.hpp"
#include "macros.hpp"

using namespace examples;

int main() {

    constexpr int count = 32;

    arch_t dx_sm = get_dx_sm();
    arch_t target_sm = get_target_sm();

    // Create descriptor. See: https://docs.nvidia.com/cuda/curanddx/
    curanddxDescriptor h { 0 };
    LIBMATHDX_CHECK(curanddxCreateDescriptor(&h));
    LIBMATHDX_CHECK(curanddxSetOperatorInt64(h, CURANDDX_OPERATOR_GENERATOR, CURANDDX_GENERATOR_PCG));
    LIBMATHDX_CHECK(curanddxSetOperatorInt64(h, CURANDDX_OPERATOR_EXECUTION, COMMONDX_EXECUTION_THREAD));
    LIBMATHDX_CHECK(curanddxSetOperatorInt64(h, CURANDDX_OPERATOR_DISTRIBUTION, CURANDDX_DISTRIBUTION_UNIFORM));
    LIBMATHDX_CHECK(curanddxSetOperatorInt64(h, CURANDDX_OPERATOR_OUTPUT_TYPE, COMMONDX_R_32F));
    LIBMATHDX_CHECK(curanddxSetOperatorInt64(h, CURANDDX_OPERATOR_GENERATE_METHOD, CURANDDX_GENERATE_METHOD_SINGLE));
    LIBMATHDX_CHECK(curanddxSetOperatorInt64(h, CURANDDX_OPERATOR_SM, dx_sm.operator_sm()));
    LIBMATHDX_CHECK(curanddxSetOperatorInt64(h,
                                             CURANDDX_OPERATOR_DEVICE_FUNCTIONS,
                                             CURANDDX_DEVICE_FUNCTION_GENERATE | CURANDDX_DEVICE_FUNCTION_INIT_STATE |
                                                 CURANDDX_DEVICE_FUNCTION_DESTROY_STATE));

    // Generate device function (LTOIR)
    commondxCode code;
    LIBMATHDX_CHECK(commondxCreateCode(&code));
    LIBMATHDX_CHECK(commondxSetCodeOptionInt64(code, COMMONDX_OPTION_TARGET_SM, target_sm.operator_sm()));
    LIBMATHDX_CHECK(curanddxFinalizeCode(code, h));

    size_t lto_size = 0;
    LIBMATHDX_CHECK(commondxGetCodeLTOIRSize(code, &lto_size));
    std::vector<char> lto(lto_size);
    LIBMATHDX_CHECK(commondxGetCodeLTOIR(code, lto.size(), lto.data()));
    LIBMATHDX_CHECK(commondxDestroyCode(code));

    // Query symbol names for declaration and call
    size_t size = 0;
    LIBMATHDX_CHECK(curanddxGetTraitStrSize(h, CURANDDX_TRAIT_SYMBOL_GENERATE_NAME, &size));
    std::vector<char> generate_name(size);
    LIBMATHDX_CHECK(curanddxGetTraitStr(h, CURANDDX_TRAIT_SYMBOL_GENERATE_NAME, size, generate_name.data()));
    LIBMATHDX_CHECK(curanddxGetTraitStrSize(h, CURANDDX_TRAIT_SYMBOL_INIT_STATE_NAME, &size));
    std::vector<char> init_name(size);
    LIBMATHDX_CHECK(curanddxGetTraitStr(h, CURANDDX_TRAIT_SYMBOL_INIT_STATE_NAME, size, init_name.data()));
    LIBMATHDX_CHECK(curanddxGetTraitStrSize(h, CURANDDX_TRAIT_SYMBOL_DESTROY_STATE_NAME, &size));
    std::vector<char> destroy_name(size);
    LIBMATHDX_CHECK(curanddxGetTraitStr(h, CURANDDX_TRAIT_SYMBOL_DESTROY_STATE_NAME, size, destroy_name.data()));

    // Query opaque state size & alignment
    long long int state_size = 0;
    long long int state_align = 0;
    LIBMATHDX_CHECK(curanddxGetTraitInt64(h, CURANDDX_TRAIT_STATE_SIZE, &state_size));
    LIBMATHDX_CHECK(curanddxGetTraitInt64(h, CURANDDX_TRAIT_STATE_ALIGNMENT, &state_align));

    printf("Successfully generated LTOIR, %zu bytes for cuRANDDx device functions %s, %s, and %s\n",
           lto_size,
           generate_name.data(),
           init_name.data(),
           destroy_name.data());

    LIBMATHDX_CHECK(curanddxDestroyDescriptor(h));

    const float MIN = -2.0;
    const float MAX = 3.0;
    // Create CUDA C++ kernel that calls the device function
    const char kernel_template[] = R"(
    #define RNG_INIT %s
    #define MIN %f
    #define MAX %f
    #define RNG_GENERATE %s
    #define RNG_FREE %s
    constexpr unsigned state_size = %d;
    constexpr unsigned state_alignment = %d;

    extern "C" __device__ void RNG_INIT(unsigned long long* seed,
                                        unsigned long long* subseq,
                                        unsigned long long* offset,
                                        void* state);

    extern "C" __device__ void RNG_FREE(void* state);

    extern "C" __device__ void RNG_GENERATE(void* state,
                                            float* out, 
                                            float* params);

    extern "C" __global__ void curand_example_kernel(float* out) {
        const unsigned tid = threadIdx.x;
        unsigned long long seed = 1234ull;
        unsigned long long subseq = 0ull;
        unsigned long long offset = (unsigned long long)(1 + tid);
        alignas(state_alignment) char state[state_size];
        RNG_INIT(&seed, &subseq, &offset, state);
        float params[2] = {MIN, MAX};
        RNG_GENERATE(state, out + tid, params);
        RNG_FREE(state);
    }
    )";

    // Compile and link
    std::string final_kernel = strprintf(kernel_template,
                                         init_name.data(),
                                         MIN,
                                         MAX,
                                         generate_name.data(),
                                         destroy_name.data(),
                                         (int)state_size,
                                         (int)state_align);
    std::vector<char> cubin = compile_and_link(final_kernel, lto, target_sm);

    // Load and run
    CUmodule module {};
    CUfunction kernel {};
    CUDA_CHECK(cudaSetDevice(0));
    CU_CHECK(cuModuleLoadDataEx(&module, cubin.data(), 0, 0, 0));
    CU_CHECK(cuModuleGetFunction(&kernel, module, "curand_example_kernel"));

    float* out = nullptr;
    CUDA_CHECK(cudaMallocManaged(&out, count * sizeof(float)));
    std::fill(out, out + count, 0.0f);

    std::array args = { static_cast<void*>(&out) };
    CU_CHECK(cuLaunchKernel(kernel, 1, 1, 1, count, 1, 1, 0, 0, args.data(), nullptr));
    CUDA_CHECK(cudaDeviceSynchronize());

    // Print and basic range check
    std::cout << "Generated " << count << " uniform random values in [" << MIN << "," << MAX << "):\n";
    for (unsigned i = 0; i < count; i++) {
        float v = out[i];
        std::cout << v << (i + 1 == count ? '\n' : ' ');
        if (!(v >= MIN && v < MAX)) {
            std::cerr << "Value out of range at index " << i << ": " << v << "\n";
            std::abort();
        }
    }

    CU_CHECK(cuModuleUnload(module));
    CUDA_CHECK(cudaFree(out));

    return 0;
}

API reference#

typedef long long int curanddxDescriptor#

A cuRANDDx descriptor.

Represents a configured cuRANDDx device function description.

enum curanddxDistribution_t#

Distribution type.

Indicate the random number generation to use. Can be set using CURANDDX_OPERATOR_DISTRIBUTION and further customized using CURANDDX_OPERATOR_NORMAL_METHOD. See cuRANDDx distributions in the documentation: https://docs.nvidia.com/cuda/curanddx/api/methods.html#random-number-generation-with-distributions.

Values:

enumerator CURANDDX_DISTRIBUTION_UNIFORM_BITS#: Uniform bits distribution (raw random bits without transformation)

enumerator CURANDDX_DISTRIBUTION_UNIFORM#: Uniform distribution in [0,1)

enumerator CURANDDX_DISTRIBUTION_NORMAL#: Normal distribution

enumerator CURANDDX_DISTRIBUTION_LOG_NORMAL#: Log-normal distribution

enumerator CURANDDX_DISTRIBUTION_POISSON#: Poisson distribution

enum curanddxNormalMethod_t#

Normal distribution method type.

Specifies the method to use for normal and log-normal distributions. Set using CURANDDX_OPERATOR_NORMAL_METHOD . See cuRANDDx normal distribution documentation: https://docs.nvidia.com/cuda/curanddx/api/methods.html#normal-distribution

Values:

enumerator CURANDDX_NORMAL_METHOD_ICDF#: Inverse cumulative distribution function (ICDF), generating one value at a time

enumerator CURANDDX_NORMAL_METHOD_BOX_MULLER#: Box-Muller method, requiring two input values and generating two output values

enum curanddxGenerateMethod_t#

Generation method type.

Specifies which generation method to use for cuRANDDx distributions. Set using CURANDDX_OPERATOR_GENERATE_METHOD . See cuRANDDx execution methods in the documentation: https://docs.nvidia.com/cuda/curanddx/api/methods.html

Values:

enumerator CURANDDX_GENERATE_METHOD_SINGLE#: Generate single value (invalid for Philox with some distributions)

enumerator CURANDDX_GENERATE_METHOD_PAIR#: Generate two values (for Box-Muller or Philox double precision)

enumerator CURANDDX_GENERATE_METHOD_QUAD#: Generate four values (for Philox generator)

enum curanddxGenerator_t#

Generator type.

Matches cuRANDDx generator kinds. Set using CURANDDX_OPERATOR_GENERATOR . See cuRANDDx generators in the documentation: https://docs.nvidia.com/cuda/curanddx/api/description_ops.html#generator-operator-label

Values:

enumerator CURANDDX_GENERATOR_XORWOW#

enumerator CURANDDX_GENERATOR_MRG32K3A#

enumerator CURANDDX_GENERATOR_PHILOX4_32#

enumerator CURANDDX_GENERATOR_PCG#

enumerator CURANDDX_GENERATOR_SOBOL32#

enumerator CURANDDX_GENERATOR_SCRAMBLED_SOBOL32#

enumerator CURANDDX_GENERATOR_SOBOL64#

enumerator CURANDDX_GENERATOR_SCRAMBLED_SOBOL64#

enum curanddxOperatorType_t#

Operators.

The set of supported curandDx operators.

Values:

enumerator CURANDDX_OPERATOR_GENERATOR#: Operator data type: curanddxGenerator_t. Operator definition: required

enumerator CURANDDX_OPERATOR_PHILOX_ROUNDS#: Operator data type: long long int. Expected content: number of philox rounds when generator is PHILOX. See https://docs.nvidia.com/cuda/curanddx/api/description_ops.html#philoxrounds-operator. Operator definition: optional (default is 10)

enumerator CURANDDX_OPERATOR_SM#: Operator data type: long long int. Expected content: 700 (Volta), 800 (Ampere), 890, 900, 1000, 1200, … Operator definition: required

enumerator CURANDDX_OPERATOR_EXECUTION#: Operator data type: commondxExecution_t. Only support COMMONDX_EXECUTION_THREAD. Operator definition: required

enumerator CURANDDX_OPERATOR_DISTRIBUTION#: Operator data type: curanddxDistribution_t. Indicate which distribution to use. Operator definition: required

enumerator CURANDDX_OPERATOR_OUTPUT_TYPE#: Operator data type: commondxValueType_t. Expected content: scale output element value type (e.g., COMMONDX_R_32F for float). The actual output type is determined by this operator and by CURANDDX_OPERATOR_GENERATE_METHOD. Operator definition: required

enumerator CURANDDX_OPERATOR_GENERATE_METHOD#: Operator data type: curanddxGenerateMethod_t. Expected content: generation method (generate, generate2, generate4). Combined with CURANDDX_OPERATOR_OUTPUT_TYPE this determines the output element types. Operator definition: required

enumerator CURANDDX_OPERATOR_NORMAL_METHOD#: Operator data type: curanddxNormalMethod_t. Expected content: normal distribution method (ICDF or Box-Muller). Operator definition: optional

enumerator CURANDDX_OPERATOR_DEVICE_FUNCTIONS#: Operator data type: bit fields of xor’ed curanddxDeviceFunctionType_t. Expected content: what device functions to generate. Operator definition: required.

enum curanddxTraitType_t#

Traits.

The set of supported types of traits that can be accessed from finalized sources that use curanddx.

Values:

enumerator CURANDDX_TRAIT_SYMBOL_GENERATE_NAME#: Trait data type: C-string. Value: CURANDDX_DEVICE_FUNCTION_GENERATE symbol name.

enumerator CURANDDX_TRAIT_SYMBOL_INIT_STATE_NAME#: Trait data type: C-string. Value: CURANDDX_DEVICE_FUNCTION_INIT_STATE symbol name.

enumerator CURANDDX_TRAIT_SYMBOL_DESTROY_STATE_NAME#: Trait data type: C-string. Value: CURANDDX_DEVICE_FUNCTION_DESTROY_STATE symbol name.

enumerator CURANDDX_TRAIT_SYMBOL_SKIP_OFFSET_NAME#: Trait data type: C-string. Value: CURANDDX_DEVICE_FUNCTION_SKIP_OFFSET symbol name.

enumerator CURANDDX_TRAIT_SYMBOL_SKIP_SUBSEQUENCE_NAME#: Trait data type: C-string. Value: CURANDDX_DEVICE_FUNCTION_SKIP_SUBSEQUENCE symbol name.

enumerator CURANDDX_TRAIT_SYMBOL_SKIP_SEQUENCE_NAME#: Trait data type: C-string. Value: CURANDDX_DEVICE_FUNCTION_SKIP_SEQUENCE symbol name.

enumerator CURANDDX_TRAIT_STATE_SIZE#: Trait data type: long long int. Value: RNG state size in bytes.

enumerator CURANDDX_TRAIT_STATE_ALIGNMENT#: Trait data type: long long int. Value: RNG state alignment in bytes.

enum curanddxDeviceFunctionType_t#

Device functions.

The set of supported device functions. Must be passed using CURANDDX_OPERATOR_DEVICE_FUNCTIONS and can be or’ed to indicate that libmathdx should generate multiple device functions.

Values:

enumerator CURANDDX_DEVICE_FUNCTION_GENERATE#

Generate random numbers using pre-initialized state. The device function has signature void(void* state, vector_output_type* out, param_type* params) where

output_type corresponds to CURANDDX_OPERATOR_OUTPUT_TYPE
vector_output_type is a vector of scalars of type output_type and length determined by CURANDDX_OPERATOR_GENERATE_METHOD
param_type is the same as output_type for non-Poisson distributions, and double for the Poisson distribution params points to distribution parameters: two elements (min, max) for uniform, (mean, std) for normal/log-normal, one element (lambda) for Poisson; unused for uniform_bits. See also https://docs.nvidia.com/cuda/curanddx/api/methods.html#random-number-generation-with-distributions and CURANDDX_TRAIT_SYMBOL_GENERATE_NAME.

enumerator CURANDDX_DEVICE_FUNCTION_INIT_STATE#

State initialization function for pseudorandom generators. The device function has signature void(unsigned long long* seed, unsigned long long* subsequence, offset_type* offset, void* state) for pseudo-random number generators, and

void(unsigned int* dim, direction_vector_type* direction_vector, offset_type* offset,


scrambled_const_type* scrambled_consts, void* state)

for quasi-random number generators. offset_type, direction_vector_type and scrambled_const_type are all defined like in cuRANDDx. See CURANDDX_TRAIT_SYMBOL_INIT_STATE_NAME.

enumerator CURANDDX_DEVICE_FUNCTION_DESTROY_STATE#: State destruction function for pseudorandom generators. The device function has signature void(void* state). See CURANDDX_TRAIT_SYMBOL_DESTROY_STATE_NAME.

enumerator CURANDDX_DEVICE_FUNCTION_SKIP_OFFSET#: Skip offset function (valid for all generators except scrambled SOBOL). The device function has signature void(void* state, offset_type* n). See CURANDDX_TRAIT_SYMBOL_SKIP_OFFSET_NAME.

enumerator CURANDDX_DEVICE_FUNCTION_SKIP_SUBSEQUENCE#: Skip subsequence function (valid for pseudorandom generators). The device function has signature void(void* state, unsigned long long* n). See CURANDDX_TRAIT_SYMBOL_SKIP_SUBSEQUENCE_NAME.

enumerator CURANDDX_DEVICE_FUNCTION_SKIP_SEQUENCE#: Skip sequence function (valid for MRG32K3A generator only). The device function has signature void(void* state, unsigned long long* n). See CURANDDX_TRAIT_SYMBOL_SKIP_SEQUENCE_NAME.

commondxStatusType curanddxGetVersion( int *major, int *minor, int *patch, )#

Returns the major.minor.patch version of cuRANDDx.

Parameters:

major – [out] The major version
minor – [out] The minor version
patch – [out] The patch version

Returns:

COMMONDX_SUCCESS

commondxStatusType curanddxCreateDescriptor( curanddxDescriptor *handle, )#

Creates a cuRANDDx descriptor.

Parameters:: handle – [inout] A pointer to a descriptor handle. As output, an initialized cuRANDDx descriptor.
Returns:: COMMONDX_SUCCESS on success, or an error code.

commondxStatusType curanddxSetOptionStr( curanddxDescriptor handle, commondxOption opt, const char *value, )#

Sets a C-string option on a cuRANDDx descriptor.

Parameters:

handle – [in] A cuRANDDx descriptor, output of curanddxCreateDescriptor .
opt – [in] The option to set.
value – [in] The value for the option.

Returns: