# Program Listing for File rocrand_uniform.h

↰ Return to documentation for file (`library/include/rocrand/rocrand_uniform.h`)

    // Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
    //
    // Permission is hereby granted, free of charge, to any person obtaining a copy
    // of this software and associated documentation files (the "Software"), to deal
    // in the Software without restriction, including without limitation the rights
    // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
    // copies of the Software, and to permit persons to whom the Software is
    // furnished to do so, subject to the following conditions:
    //
    // The above copyright notice and this permission notice shall be included in
    // all copies or substantial portions of the Software.
    //
    // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
    // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
    // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
    // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
    // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
    // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
    // THE SOFTWARE.

    #ifndef ROCRAND_UNIFORM_H_
    #define ROCRAND_UNIFORM_H_

    #ifndef FQUALIFIERS
    #define FQUALIFIERS __forceinline__ __device__
    #endif // FQUALIFIERS

    #include "rocrand/rocrand_lfsr113.h"
    #include "rocrand/rocrand_mrg31k3p.h"
    #include "rocrand/rocrand_mrg32k3a.h"
    #include "rocrand/rocrand_mtgp32.h"
    #include "rocrand/rocrand_philox4x32_10.h"
    #include "rocrand/rocrand_scrambled_sobol32.h"
    #include "rocrand/rocrand_scrambled_sobol64.h"
    #include "rocrand/rocrand_sobol32.h"
    #include "rocrand/rocrand_sobol64.h"
    #include "rocrand/rocrand_threefry2x32_20.h"
    #include "rocrand/rocrand_threefry2x64_20.h"
    #include "rocrand/rocrand_threefry4x32_20.h"
    #include "rocrand/rocrand_threefry4x64_20.h"
    #include "rocrand/rocrand_xorwow.h"

    #include "rocrand/rocrand_common.h"

    namespace rocrand_device {
    namespace detail {

    struct two_uints
    {
        unsigned int x;
        unsigned int y;
    };

    union two_uints_to_ulong
    {
        two_uints uint2_value;
        unsigned long long int ulong_value;
    };

    // For unsigned integer between 0 and UINT_MAX, returns value between
    // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
    FQUALIFIERS
    float uniform_distribution(unsigned int v)
    {
        return ROCRAND_2POW32_INV + (v * ROCRAND_2POW32_INV);
    }

    // For unsigned integer between 0 and ULLONG_MAX, returns value between
    // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
    FQUALIFIERS
    float uniform_distribution(unsigned long long int v)
    {
        return ROCRAND_2POW32_INV + (v >> 32) * ROCRAND_2POW32_INV;
    }

    FQUALIFIERS
    float4 uniform_distribution4(uint4 v)
    {
       return float4 {
           ROCRAND_2POW32_INV + (v.x * ROCRAND_2POW32_INV),
           ROCRAND_2POW32_INV + (v.y * ROCRAND_2POW32_INV),
           ROCRAND_2POW32_INV + (v.z * ROCRAND_2POW32_INV),
           ROCRAND_2POW32_INV + (v.w * ROCRAND_2POW32_INV)
       };
    }

    FQUALIFIERS float4 uniform_distribution4(ulonglong4 v)
    {
        return float4{ROCRAND_2POW64_INV + (v.x * ROCRAND_2POW64_INV),
                      ROCRAND_2POW64_INV + (v.y * ROCRAND_2POW64_INV),
                      ROCRAND_2POW64_INV + (v.z * ROCRAND_2POW64_INV),
                      ROCRAND_2POW64_INV + (v.w * ROCRAND_2POW64_INV)};
    }

    // For unsigned integer between 0 and UINT_MAX, returns value between
    // 0.0 and 1.0, excluding 0.0 and including 1.0.
    FQUALIFIERS
    double uniform_distribution_double(unsigned int v)
    {
        return ROCRAND_2POW32_INV_DOUBLE + (v * ROCRAND_2POW32_INV_DOUBLE);
    }

    FQUALIFIERS
    double uniform_distribution_double(unsigned int v1, unsigned int v2)
    {
        two_uints_to_ulong v;
        v.uint2_value.x = v1;
        v.uint2_value.y = (v2 >> 11);
        return ROCRAND_2POW53_INV_DOUBLE + (v.ulong_value * ROCRAND_2POW53_INV_DOUBLE);
    }

    FQUALIFIERS
    double uniform_distribution_double(unsigned long long int v)
    {
        return ROCRAND_2POW53_INV_DOUBLE + (
            // 2^53 is the biggest int that can be stored in double, such
            // that it and all smaller integers can be stored in double
            (v >> 11) * ROCRAND_2POW53_INV_DOUBLE
        );
    }

    FQUALIFIERS
    double2 uniform_distribution_double2(uint4 v)
    {
        return double2 {
            uniform_distribution_double(v.x, v.y),
            uniform_distribution_double(v.z, v.w)
        };
    }

    FQUALIFIERS
    double4 uniform_distribution_double4(uint4 v1, uint4 v2)
    {
        return double4 {
            uniform_distribution_double(v1.x, v1.y),
            uniform_distribution_double(v1.z, v1.w),
            uniform_distribution_double(v2.x, v2.y),
            uniform_distribution_double(v2.z, v2.w)
        };
    }

    FQUALIFIERS double2 uniform_distribution_double2(ulonglong2 v)
    {
        return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
    }

    FQUALIFIERS double2 uniform_distribution_double2(ulonglong4 v)
    {
        return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
    }

    FQUALIFIERS double4 uniform_distribution_double4(ulonglong4 v)
    {
        return double4{uniform_distribution_double(v.x),
                       uniform_distribution_double(v.z),
                       uniform_distribution_double(v.x),
                       uniform_distribution_double(v.z)};
    }

    FQUALIFIERS
    __half uniform_distribution_half(unsigned short v)
    {
        return __float2half(ROCRAND_2POW16_INV + (v * ROCRAND_2POW16_INV));
    }

    // For an unsigned integer produced by an MRG-based engine, returns a value
    // in range [0, UINT32_MAX].
    template<typename state_type>
    FQUALIFIERS unsigned int mrg_uniform_distribution_uint(unsigned int v) = delete;

    template<>
    FQUALIFIERS unsigned int mrg_uniform_distribution_uint<rocrand_state_mrg31k3p>(unsigned int v)
    {
        return static_cast<unsigned int>((v - 1) * ROCRAND_MRG31K3P_UINT32_NORM);
    }

    template<>
    FQUALIFIERS unsigned int mrg_uniform_distribution_uint<rocrand_state_mrg32k3a>(unsigned int v)
    {
        return static_cast<unsigned int>((v - 1) * ROCRAND_MRG32K3A_UINT_NORM);
    }

    // For an unsigned integer produced by an MRG-based engine, returns value between
    // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
    template<typename state_type>
    FQUALIFIERS float mrg_uniform_distribution(unsigned int v) = delete;

    template<>
    FQUALIFIERS float mrg_uniform_distribution<rocrand_state_mrg31k3p>(unsigned int v)
    {
        double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
        return static_cast<float>(ret);
    }

    template<>
    FQUALIFIERS float mrg_uniform_distribution<rocrand_state_mrg32k3a>(unsigned int v)
    {
        double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
        return static_cast<float>(ret);
    }

    // For an unsigned integer produced by an MRG generator, returns value between
    // 0.0 and 1.0, excluding 0.0 and including 1.0.
    template<typename state_type>
    FQUALIFIERS double mrg_uniform_distribution_double(unsigned int v) = delete;

    template<>
    FQUALIFIERS double mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(unsigned int v)
    {
        double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
        return ret;
    }

    template<>
    FQUALIFIERS double mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(unsigned int v)
    {
        double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
        return ret;
    }

    } // end namespace detail
    } // end namespace rocrand_device

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_philox4x32_10 * state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS
    float2 rocrand_uniform2(rocrand_state_philox4x32_10 * state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);

        return float2 {
            rocrand_device::detail::uniform_distribution(state1),
            rocrand_device::detail::uniform_distribution(state2)
        };
    }

    FQUALIFIERS
    float4 rocrand_uniform4(rocrand_state_philox4x32_10 * state)
    {
        return rocrand_device::detail::uniform_distribution4(rocrand4(state));
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_philox4x32_10 * state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);

        return rocrand_device::detail::uniform_distribution_double(state1, state2);
    }

    FQUALIFIERS
    double2 rocrand_uniform_double2(rocrand_state_philox4x32_10 * state)
    {
        return rocrand_device::detail::uniform_distribution_double2(rocrand4(state));
    }

    FQUALIFIERS
    double4 rocrand_uniform_double4(rocrand_state_philox4x32_10 * state)
    {
        return rocrand_device::detail::uniform_distribution_double4(rocrand4(state), rocrand4(state));
    }

    FQUALIFIERS float rocrand_uniform(rocrand_state_mrg31k3p* state)
    {
        return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg31k3p>(state->next());
    }

    FQUALIFIERS double rocrand_uniform_double(rocrand_state_mrg31k3p* state)
    {
        return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(
            state->next());
    }

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_mrg32k3a * state)
    {
        return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg32k3a>(state->next());
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_mrg32k3a * state)
    {
        return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(
            state->next());
    }

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_xorwow * state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_xorwow * state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);

        return rocrand_device::detail::uniform_distribution_double(state1, state2);
    }

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_mtgp32 * state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_mtgp32 * state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_sobol32 * state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_sobol32 * state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_scrambled_sobol32* state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_scrambled_sobol32* state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_sobol64* state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_sobol64 * state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_scrambled_sobol64* state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_scrambled_sobol64* state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_uniform(rocrand_state_lfsr113* state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_uniform_double(rocrand_state_lfsr113* state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS float rocrand_uniform(rocrand_state_threefry2x32_20* state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS double rocrand_uniform_double(rocrand_state_threefry2x32_20* state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS float rocrand_uniform(rocrand_state_threefry2x64_20* state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS double rocrand_uniform_double(rocrand_state_threefry2x64_20* state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS float rocrand_uniform(rocrand_state_threefry4x32_20* state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS double rocrand_uniform_double(rocrand_state_threefry4x32_20* state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

    FQUALIFIERS float rocrand_uniform(rocrand_state_threefry4x64_20* state)
    {
        return rocrand_device::detail::uniform_distribution(rocrand(state));
    }

    FQUALIFIERS double rocrand_uniform_double(rocrand_state_threefry4x64_20* state)
    {
        return rocrand_device::detail::uniform_distribution_double(rocrand(state));
    }

     // end of group rocranddevice
    #endif // ROCRAND_UNIFORM_H_