
.. _program_listing_file_library_include_rocrand_rocrand_normal.h:

Program Listing for File rocrand_normal.h
=========================================

|exhale_lsh| :ref:`Return to documentation for file <file_library_include_rocrand_rocrand_normal.h>` (``library/include/rocrand/rocrand_normal.h``)

.. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS

.. code-block:: cpp

   // 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_
