# Program Listing for File rocrand_normal.h

↰ Return to documentation for file (`library/include/rocrand/rocrand_normal.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_NORMAL_H_
    #define ROCRAND_NORMAL_H_

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

    #include <math.h>

    #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_uniform.h"

    namespace rocrand_device {
    namespace detail {

    FQUALIFIERS
    float2 box_muller(unsigned int x, unsigned int y)
    {
        float2 result;
        float u = ROCRAND_2POW32_INV + (x * ROCRAND_2POW32_INV);
        float v = ROCRAND_2POW32_INV_2PI + (y * ROCRAND_2POW32_INV_2PI);
        float s = sqrtf(-2.0f * logf(u));
        #ifdef __HIP_DEVICE_COMPILE__
            __sincosf(v, &result.x, &result.y);
            result.x *= s;
            result.y *= s;
        #else
            result.x = sinf(v) * s;
            result.y = cosf(v) * s;
        #endif
        return result;
    }

    FQUALIFIERS float2 box_muller(unsigned long long v)
    {
        unsigned int x = static_cast<unsigned int>(v);
        unsigned int y = static_cast<unsigned int>(v >> 32);

        return box_muller(x, y);
    }

    FQUALIFIERS
    double2 box_muller_double(uint4 v)
    {
        double2 result;
        unsigned long long int v1 = (unsigned long long int)v.x ^
            ((unsigned long long int)v.y << (53 - 32));
        double u = ROCRAND_2POW53_INV_DOUBLE + (v1 * ROCRAND_2POW53_INV_DOUBLE);
        unsigned long long int v2 = (unsigned long long int)v.z ^
            ((unsigned long long int)v.w << (53 - 32));
        double w = (ROCRAND_2POW53_INV_DOUBLE * 2.0) +
            (v2 * (ROCRAND_2POW53_INV_DOUBLE * 2.0));
        double s = sqrt(-2.0 * log(u));
        #ifdef __HIP_DEVICE_COMPILE__
            sincospi(w, &result.x, &result.y);
            result.x *= s;
            result.y *= s;
        #else
            result.x = sin(w * ROCRAND_PI_DOUBLE) * s;
            result.y = cos(w * ROCRAND_PI_DOUBLE) * s;
        #endif
        return result;
    }

    FQUALIFIERS double2 box_muller_double(ulonglong2 v)
    {
        unsigned int x = static_cast<unsigned int>(v.x);
        unsigned int y = static_cast<unsigned int>(v.x >> 32);
        unsigned int z = static_cast<unsigned int>(v.y);
        unsigned int w = static_cast<unsigned int>(v.y >> 32);

        return box_muller_double(make_uint4(x, y, z, w));
    }

    FQUALIFIERS
    __half2 box_muller_half(unsigned short x, unsigned short y)
    {
        #if defined(ROCRAND_HALF_MATH_SUPPORTED)
        __half u = __float2half(ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV));
        __half v = __float2half(ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI));
        __half s = hsqrt(__hmul(__float2half(-2.0f), hlog(u)));
        return __half2 {
            __hmul(hsin(v), s),
            __hmul(hcos(v), s)
        };
        #else
        float2 r;
        float u = ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV);
        float v = ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI);
        float s = sqrtf(-2.0f * logf(u));
        #ifdef __HIP_DEVICE_COMPILE__
            __sincosf(v, &r.x, &r.y);
            r.x *= s;
            r.y *= s;
        #else
            r.x = sinf(v) * s;
            r.y = cosf(v) * s;
        #endif
        return __half2 {
            __float2half(r.x),
            __float2half(r.y)
        };
        #endif
    }

    template<typename state_type>
    FQUALIFIERS float2 mrg_box_muller(unsigned int x, unsigned int y)
    {
        float2 result;
        float  u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);
        float  v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * ROCRAND_2PI;
        float s = sqrtf(-2.0f * logf(u));
        #ifdef __HIP_DEVICE_COMPILE__
            __sincosf(v, &result.x, &result.y);
            result.x *= s;
            result.y *= s;
        #else
            result.x = sinf(v) * s;
            result.y = cosf(v) * s;
        #endif
        return result;
    }

    template<typename state_type>
    FQUALIFIERS double2 mrg_box_muller_double(unsigned int x, unsigned int y)
    {
        double2 result;
        double  u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);
        double  v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * 2.0;
        double s = sqrt(-2.0 * log(u));
        #ifdef __HIP_DEVICE_COMPILE__
            sincospi(v, &result.x, &result.y);
            result.x *= s;
            result.y *= s;
        #else
            result.x = sin(v * ROCRAND_PI_DOUBLE) * s;
            result.y = cos(v * ROCRAND_PI_DOUBLE) * s;
        #endif
        return result;
    }

    FQUALIFIERS
    float roc_f_erfinv(float x)
    {
        float tt1, tt2, lnx, sgn;
        sgn = (x < 0.0f) ? -1.0f : 1.0f;

        x = (1.0f - x) * (1.0f + x);
        lnx = logf(x);

        #ifdef __HIP_DEVICE_COMPILE__
        if (isnan(lnx))
        #else
        if (std::isnan(lnx))
        #endif
            return 1.0f;
        #ifdef __HIP_DEVICE_COMPILE__
        else if (isinf(lnx))
        #else
        else if (std::isinf(lnx))
        #endif
            return 0.0f;

        tt1 = 2.0f / (ROCRAND_PI * 0.147f) + 0.5f * lnx;
        tt2 = 1.0f / (0.147f) * lnx;

        return(sgn * sqrtf(-tt1 + sqrtf(tt1 * tt1 - tt2)));
    }

    FQUALIFIERS
    double roc_d_erfinv(double x)
    {
        double tt1, tt2, lnx, sgn;
        sgn = (x < 0.0) ? -1.0 : 1.0;

        x = (1.0 - x) * (1.0 + x);
        lnx = log(x);

        #ifdef __HIP_DEVICE_COMPILE__
        if (isnan(lnx))
        #else
        if (std::isnan(lnx))
        #endif
            return 1.0;
        #ifdef __HIP_DEVICE_COMPILE__
        else if (isinf(lnx))
        #else
        else if (std::isinf(lnx))
        #endif
            return 0.0;

        tt1 = 2.0 / (ROCRAND_PI_DOUBLE * 0.147) + 0.5 * lnx;
        tt2 = 1.0 / (0.147) * lnx;

        return(sgn * sqrt(-tt1 + sqrt(tt1 * tt1 - tt2)));
    }

    FQUALIFIERS
    float normal_distribution(unsigned int x)
    {
        float p = ::rocrand_device::detail::uniform_distribution(x);
        float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);
        return v;
    }

    FQUALIFIERS
    float normal_distribution(unsigned long long int x)
    {
        float p = ::rocrand_device::detail::uniform_distribution(x);
        float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);
        return v;
    }

    FQUALIFIERS
    float2 normal_distribution2(unsigned int v1, unsigned int v2)
    {
        return ::rocrand_device::detail::box_muller(v1, v2);
    }

    FQUALIFIERS float2 normal_distribution2(uint2 v)
    {
        return ::rocrand_device::detail::box_muller(v.x, v.y);
    }

    FQUALIFIERS float2 normal_distribution2(unsigned long long v)
    {
        return ::rocrand_device::detail::box_muller(v);
    }

    FQUALIFIERS
    float4 normal_distribution4(uint4 v)
    {
        float2 r1 = ::rocrand_device::detail::box_muller(v.x, v.y);
        float2 r2 = ::rocrand_device::detail::box_muller(v.z, v.w);
        return float4{
            r1.x,
            r1.y,
            r2.x,
            r2.y
        };
    }

    FQUALIFIERS float4 normal_distribution4(longlong2 v)
    {
        float2 r1 = ::rocrand_device::detail::box_muller(v.x);
        float2 r2 = ::rocrand_device::detail::box_muller(v.y);
        return float4{r1.x, r1.y, r2.x, r2.y};
    }

    FQUALIFIERS float4 normal_distribution4(unsigned long long v1, unsigned long long v2)
    {
        float2 r1 = ::rocrand_device::detail::box_muller(v1);
        float2 r2 = ::rocrand_device::detail::box_muller(v2);
        return float4{r1.x, r1.y, r2.x, r2.y};
    }

    FQUALIFIERS
    double normal_distribution_double(unsigned int x)
    {
        double p = ::rocrand_device::detail::uniform_distribution_double(x);
        double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);
        return v;
    }

    FQUALIFIERS
    double normal_distribution_double(unsigned long long int x)
    {
        double p = ::rocrand_device::detail::uniform_distribution_double(x);
        double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);
        return v;
    }

    FQUALIFIERS
    double2 normal_distribution_double2(uint4 v)
    {
        return ::rocrand_device::detail::box_muller_double(v);
    }

    FQUALIFIERS double2 normal_distribution_double2(ulonglong2 v)
    {
        return ::rocrand_device::detail::box_muller_double(v);
    }

    FQUALIFIERS
    __half2 normal_distribution_half2(unsigned int v)
    {
        return ::rocrand_device::detail::box_muller_half(
            static_cast<unsigned short>(v),
            static_cast<unsigned short>(v >> 16)
        );
    }

    FQUALIFIERS __half2 normal_distribution_half2(unsigned long long v)
    {
        return ::rocrand_device::detail::box_muller_half(static_cast<unsigned short>(v),
                                                         static_cast<unsigned short>(v >> 32));
    }

    template<typename state_type>
    FQUALIFIERS float2 mrg_normal_distribution2(unsigned int v1, unsigned int v2)
    {
        return ::rocrand_device::detail::mrg_box_muller<state_type>(v1, v2);
    }

    template<typename state_type>
    FQUALIFIERS double2 mrg_normal_distribution_double2(unsigned int v1, unsigned int v2)
    {
        return ::rocrand_device::detail::mrg_box_muller_double<state_type>(v1, v2);
    }

    template<typename state_type>
    FQUALIFIERS __half2 mrg_normal_distribution_half2(unsigned int v)
    {
        v = rocrand_device::detail::mrg_uniform_distribution_uint<state_type>(v);
        return ::rocrand_device::detail::box_muller_half(
            static_cast<unsigned short>(v),
            static_cast<unsigned short>(v >> 16)
        );
    }

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

    #ifndef ROCRAND_DETAIL_PHILOX_BM_NOT_IN_STATE
    FQUALIFIERS
    float rocrand_normal(rocrand_state_philox4x32_10 * state)
    {
        typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;

        if(bm_helper::has_float(state))
        {
            return bm_helper::get_float(state);
        }

        auto state1 = rocrand(state);
        auto state2 = rocrand(state);

        float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
        bm_helper::save_float(state, r.y);
        return r.x;
    }
    #endif // ROCRAND_DETAIL_PHILOX_BM_NOT_IN_STATE

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

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

    FQUALIFIERS
    float4 rocrand_normal4(rocrand_state_philox4x32_10 * state)
    {
        return rocrand_device::detail::normal_distribution4(rocrand4(state));
    }

    #ifndef ROCRAND_DETAIL_PHILOX_BM_NOT_IN_STATE
    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_philox4x32_10 * state)
    {
        typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;

        if(bm_helper::has_double(state))
        {
            return bm_helper::get_double(state);
        }
        double2 r = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
        bm_helper::save_double(state, r.y);
        return r.x;
    }
    #endif // ROCRAND_DETAIL_PHILOX_BM_NOT_IN_STATE

    FQUALIFIERS
    double2 rocrand_normal_double2(rocrand_state_philox4x32_10 * state)
    {
        return rocrand_device::detail::normal_distribution_double2(rocrand4(state));
    }

    FQUALIFIERS
    double4 rocrand_normal_double4(rocrand_state_philox4x32_10 * state)
    {
        double2 r1, r2;
        r1 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
        r2 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
        return double4 {
            r1.x, r1.y, r2.x, r2.y
        };
    }

    #ifndef ROCRAND_DETAIL_MRG31K3P_BM_NOT_IN_STATE
    FQUALIFIERS float rocrand_normal(rocrand_state_mrg31k3p* state)
    {
        typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;

        if(bm_helper::has_float(state))
        {
            return bm_helper::get_float(state);
        }

        auto state1 = state->next();
        auto state2 = state->next();

        float2 r
            = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);
        bm_helper::save_float(state, r.y);
        return r.x;
    }
    #endif // ROCRAND_DETAIL_MRG31K3P_BM_NOT_IN_STATE

    FQUALIFIERS float2 rocrand_normal2(rocrand_state_mrg31k3p* state)
    {
        auto state1 = state->next();
        auto state2 = state->next();

        return rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);
    }

    #ifndef ROCRAND_DETAIL_MRG31K3P_BM_NOT_IN_STATE
    FQUALIFIERS double rocrand_normal_double(rocrand_state_mrg31k3p* state)
    {
        typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;

        if(bm_helper::has_double(state))
        {
            return bm_helper::get_double(state);
        }

        auto state1 = state->next();
        auto state2 = state->next();

        double2 r
            = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,
                                                                                              state2);
        bm_helper::save_double(state, r.y);
        return r.x;
    }
    #endif // ROCRAND_DETAIL_MRG31K3P_BM_NOT_IN_STATE

    FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_mrg31k3p* state)
    {
        auto state1 = state->next();
        auto state2 = state->next();

        return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,
                                                                                               state2);
    }

    #ifndef ROCRAND_DETAIL_MRG32K3A_BM_NOT_IN_STATE
    FQUALIFIERS
    float rocrand_normal(rocrand_state_mrg32k3a * state)
    {
        typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;

        if(bm_helper::has_float(state))
        {
            return bm_helper::get_float(state);
        }

        auto state1 = state->next();
        auto state2 = state->next();

        float2 r
            = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
        bm_helper::save_float(state, r.y);
        return r.x;
    }
    #endif // ROCRAND_DETAIL_MRG32K3A_BM_NOT_IN_STATE

    FQUALIFIERS
    float2 rocrand_normal2(rocrand_state_mrg32k3a * state)
    {
        auto state1 = state->next();
        auto state2 = state->next();

        return rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
    }

    #ifndef ROCRAND_DETAIL_MRG32K3A_BM_NOT_IN_STATE
    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_mrg32k3a * state)
    {
        typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;

        if(bm_helper::has_double(state))
        {
            return bm_helper::get_double(state);
        }

        auto state1 = state->next();
        auto state2 = state->next();

        double2 r
            = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
                                                                                              state2);
        bm_helper::save_double(state, r.y);
        return r.x;
    }
    #endif // ROCRAND_DETAIL_MRG32K3A_BM_NOT_IN_STATE

    FQUALIFIERS
    double2 rocrand_normal_double2(rocrand_state_mrg32k3a * state)
    {
        auto state1 = state->next();
        auto state2 = state->next();

        return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
                                                                                               state2);
    }

    #ifndef ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
    FQUALIFIERS
    float rocrand_normal(rocrand_state_xorwow * state)
    {
        typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;

        if(bm_helper::has_float(state))
        {
            return bm_helper::get_float(state);
        }
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);
        float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
        bm_helper::save_float(state, r.y);
        return r.x;
    }
    #endif // ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE

    FQUALIFIERS
    float2 rocrand_normal2(rocrand_state_xorwow * state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);
        return rocrand_device::detail::normal_distribution2(state1, state2);
    }

    #ifndef ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_xorwow * state)
    {
        typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;

        if(bm_helper::has_double(state))
        {
            return bm_helper::get_double(state);
        }

        auto state1 = rocrand(state);
        auto state2 = rocrand(state);
        auto state3 = rocrand(state);
        auto state4 = rocrand(state);

        double2 r = rocrand_device::detail::normal_distribution_double2(
            uint4 { state1, state2, state3, state4 }
        );
        bm_helper::save_double(state, r.y);
        return r.x;
    }
    #endif // ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE

    FQUALIFIERS
    double2 rocrand_normal_double2(rocrand_state_xorwow * state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);
        auto state3 = rocrand(state);
        auto state4 = rocrand(state);

        return rocrand_device::detail::normal_distribution_double2(
            uint4 { state1, state2, state3, state4 }
        );
    }

    FQUALIFIERS
    float rocrand_normal(rocrand_state_mtgp32 * state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_mtgp32 * state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_normal(rocrand_state_sobol32 * state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_sobol32 * state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_normal(rocrand_state_scrambled_sobol32* state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_scrambled_sobol32* state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_normal(rocrand_state_sobol64* state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_sobol64 * state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_normal(rocrand_state_scrambled_sobol64* state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_scrambled_sobol64* state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    float rocrand_normal(rocrand_state_lfsr113* state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS
    float2 rocrand_normal2(rocrand_state_lfsr113* state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);

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

    FQUALIFIERS
    double rocrand_normal_double(rocrand_state_lfsr113* state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS
    double2 rocrand_normal_double2(rocrand_state_lfsr113* state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);
        auto state3 = rocrand(state);
        auto state4 = rocrand(state);

        return rocrand_device::detail::normal_distribution_double2(
            uint4{state1, state2, state3, state4});
    }

    FQUALIFIERS float rocrand_normal(rocrand_state_threefry2x32_20* state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS float2 rocrand_normal2(rocrand_state_threefry2x32_20* state)
    {
        return rocrand_device::detail::normal_distribution2(rocrand2(state));
    }

    FQUALIFIERS double rocrand_normal_double(rocrand_state_threefry2x32_20* state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_threefry2x32_20* state)
    {
        auto state1 = rocrand2(state);
        auto state2 = rocrand2(state);

        return rocrand_device::detail::normal_distribution_double2(
            uint4{state1.x, state1.y, state2.x, state2.y});
    }

    FQUALIFIERS float rocrand_normal(rocrand_state_threefry2x64_20* state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS float2 rocrand_normal2(rocrand_state_threefry2x64_20* state)
    {
        return rocrand_device::detail::normal_distribution2(rocrand(state));
    }

    FQUALIFIERS double rocrand_normal_double(rocrand_state_threefry2x64_20* state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_threefry2x64_20* state)
    {
        return rocrand_device::detail::normal_distribution_double2(rocrand2(state));
    }

    FQUALIFIERS float rocrand_normal(rocrand_state_threefry4x32_20* state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS float2 rocrand_normal2(rocrand_state_threefry4x32_20* state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);

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

    FQUALIFIERS double rocrand_normal_double(rocrand_state_threefry4x32_20* state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_threefry4x32_20* state)
    {
        return rocrand_device::detail::normal_distribution_double2(rocrand4(state));
    }

    FQUALIFIERS float rocrand_normal(rocrand_state_threefry4x64_20* state)
    {
        return rocrand_device::detail::normal_distribution(rocrand(state));
    }

    FQUALIFIERS float2 rocrand_normal2(rocrand_state_threefry4x64_20* state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);

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

    FQUALIFIERS double rocrand_normal_double(rocrand_state_threefry4x64_20* state)
    {
        return rocrand_device::detail::normal_distribution_double(rocrand(state));
    }

    FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_threefry4x64_20* state)
    {
        auto state1 = rocrand(state);
        auto state2 = rocrand(state);

        return rocrand_device::detail::normal_distribution_double2(ulonglong2{state1, state2});
    }

     // end of group rocranddevice
    #endif // ROCRAND_NORMAL_H_