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_