/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/docs-6.4.2/library/include/rocrand/rocrand_normal.h Source File

• library
• include
• rocrand

rocrand_normal.h

1 // Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.

2 //

3 // Permission is hereby granted, free of charge, to any person obtaining a copy

4 // of this software and associated documentation files (the "Software"), to deal

5 // in the Software without restriction, including without limitation the rights

6 // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell

7 // copies of the Software, and to permit persons to whom the Software is

8 // furnished to do so, subject to the following conditions:

9 //

10 // The above copyright notice and this permission notice shall be included in

11 // all copies or substantial portions of the Software.

12 //

13 // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR

14 // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,

15 // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE

16 // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER

17 // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,

18 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN

19 // THE SOFTWARE.

20 

21 #ifndef ROCRAND_NORMAL_H_

22 #define ROCRAND_NORMAL_H_

23 

29 #include <math.h>

30 

31 #include "rocrand/rocrand_lfsr113.h"

32 #include "rocrand/rocrand_mrg31k3p.h"

33 #include "rocrand/rocrand_mrg32k3a.h"

34 #include "rocrand/rocrand_mtgp32.h"

35 #include "rocrand/rocrand_philox4x32_10.h"

36 #include "rocrand/rocrand_scrambled_sobol32.h"

37 #include "rocrand/rocrand_scrambled_sobol64.h"

38 #include "rocrand/rocrand_sobol32.h"

39 #include "rocrand/rocrand_sobol64.h"

40 #include "rocrand/rocrand_threefry2x32_20.h"

41 #include "rocrand/rocrand_threefry2x64_20.h"

42 #include "rocrand/rocrand_threefry4x32_20.h"

43 #include "rocrand/rocrand_threefry4x64_20.h"

44 #include "rocrand/rocrand_xorwow.h"

45 

46 #include "rocrand/rocrand_uniform.h"

47 

48 namespace rocrand_device {

49 namespace detail {

50 

51 __forceinline__ __device__ __host__ float2 box_muller(unsigned int x, unsigned int y)

52 {

53  float2 result;

54  float u = ROCRAND_2POW32_INV + (x * ROCRAND_2POW32_INV);

55  float v = ROCRAND_2POW32_INV_2PI + (y * ROCRAND_2POW32_INV_2PI);

56  float s = sqrtf(-2.0f * logf(u));

57  #ifdef __HIP_DEVICE_COMPILE__

58  __sincosf(v, &result.x, &result.y);

59  result.x *= s;

60  result.y *= s;

61  #else

62  result.x = sinf(v) * s;

63  result.y = cosf(v) * s;

64  #endif

65  return result;

66 }

67 

68 __forceinline__ __device__ __host__ float2 box_muller(unsigned long long v)

69 {

70  unsigned int x = static_cast<unsigned int>(v);

71  unsigned int y = static_cast<unsigned int>(v >> 32);

72 

73  return box_muller(x, y);

74 }

75 

76 __forceinline__ __device__ __host__ double2 box_muller_double(uint4 v)

77 {

78  double2 result;

79  unsigned long long int v1 = (unsigned long long int)v.x ^

80  ((unsigned long long int)v.y << (53 - 32));

81  double u = ROCRAND_2POW53_INV_DOUBLE + (v1 * ROCRAND_2POW53_INV_DOUBLE);

82  unsigned long long int v2 = (unsigned long long int)v.z ^

83  ((unsigned long long int)v.w << (53 - 32));

84  double w = (ROCRAND_2POW53_INV_DOUBLE * 2.0) +

85  (v2 * (ROCRAND_2POW53_INV_DOUBLE * 2.0));

86  double s = sqrt(-2.0 * log(u));

87  #ifdef __HIP_DEVICE_COMPILE__

88  sincospi(w, &result.x, &result.y);

89  result.x *= s;

90  result.y *= s;

91  #else

92  result.x = sin(w * ROCRAND_PI_DOUBLE) * s;

93  result.y = cos(w * ROCRAND_PI_DOUBLE) * s;

94  #endif

95  return result;

96 }

97 

98 __forceinline__ __device__ __host__ double2 box_muller_double(ulonglong2 v)

99 {

100  unsigned int x = static_cast<unsigned int>(v.x);

101  unsigned int y = static_cast<unsigned int>(v.x >> 32);

102  unsigned int z = static_cast<unsigned int>(v.y);

103  unsigned int w = static_cast<unsigned int>(v.y >> 32);

104 

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

106 }

107 

108 __forceinline__ __device__ __host__ __half2 box_muller_half(unsigned short x, unsigned short y)

109 {

110  #if defined(ROCRAND_HALF_MATH_SUPPORTED)

111  __half u = __float2half(ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV));

112  __half v = __float2half(ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI));

113  __half s = hsqrt(__hmul(__float2half(-2.0f), hlog(u)));

114  return __half2 {

115  __hmul(hsin(v), s),

116  __hmul(hcos(v), s)

117  };

118  #else

119  float2 r;

120  float u = ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV);

121  float v = ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI);

122  float s = sqrtf(-2.0f * logf(u));

123  #ifdef __HIP_DEVICE_COMPILE__

124  __sincosf(v, &r.x, &r.y);

125  r.x *= s;

126  r.y *= s;

127  #else

128  r.x = sinf(v) * s;

129  r.y = cosf(v) * s;

130  #endif

131  return __half2 {

132  __float2half(r.x),

133  __float2half(r.y)

134  };

135  #endif

136 }

137 

138 template<typename state_type>

139 __forceinline__ __device__ __host__ float2 mrg_box_muller(unsigned int x, unsigned int y)

140 {

141  float2 result;

142  float u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);

143  float v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * ROCRAND_2PI;

144  float s = sqrtf(-2.0f * logf(u));

145  #ifdef __HIP_DEVICE_COMPILE__

146  __sincosf(v, &result.x, &result.y);

147  result.x *= s;

148  result.y *= s;

149  #else

150  result.x = sinf(v) * s;

151  result.y = cosf(v) * s;

152  #endif

153  return result;

154 }

155 

156 template<typename state_type>

157 __forceinline__ __device__ __host__ double2 mrg_box_muller_double(unsigned int x, unsigned int y)

158 {

159  double2 result;

160  double u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);

161  double v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * 2.0;

162  double s = sqrt(-2.0 * log(u));

163  #ifdef __HIP_DEVICE_COMPILE__

164  sincospi(v, &result.x, &result.y);

165  result.x *= s;

166  result.y *= s;

167  #else

168  result.x = sin(v * ROCRAND_PI_DOUBLE) * s;

169  result.y = cos(v * ROCRAND_PI_DOUBLE) * s;

170  #endif

171  return result;

172 }

173 

174 __forceinline__ __device__ __host__ float roc_f_erfinv(float x)

175 {

176  float tt1, tt2, lnx, sgn;

177  sgn = (x < 0.0f) ? -1.0f : 1.0f;

178 

179  x = (1.0f - x) * (1.0f + x);

180  lnx = logf(x);

181 

182  #ifdef __HIP_DEVICE_COMPILE__

183  if (isnan(lnx))

184  #else

185  if (std::isnan(lnx))

186  #endif

187  return 1.0f;

188  #ifdef __HIP_DEVICE_COMPILE__

189  else if (isinf(lnx))

190  #else

191  else if (std::isinf(lnx))

192  #endif

193  return 0.0f;

194 

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

196  tt2 = 1.0f / (0.147f) * lnx;

197 

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

199 }

200 

201 __forceinline__ __device__ __host__ double roc_d_erfinv(double x)

202 {

203  double tt1, tt2, lnx, sgn;

204  sgn = (x < 0.0) ? -1.0 : 1.0;

205 

206  x = (1.0 - x) * (1.0 + x);

207  lnx = log(x);

208 

209  #ifdef __HIP_DEVICE_COMPILE__

210  if (isnan(lnx))

211  #else

212  if (std::isnan(lnx))

213  #endif

214  return 1.0;

215  #ifdef __HIP_DEVICE_COMPILE__

216  else if (isinf(lnx))

217  #else

218  else if (std::isinf(lnx))

219  #endif

220  return 0.0;

221 

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

223  tt2 = 1.0 / (0.147) * lnx;

224 

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

226 }

227 

228 __forceinline__ __device__ __host__ float normal_distribution(unsigned int x)

229 {

230  float p = ::rocrand_device::detail::uniform_distribution(x);

231  float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);

232  return v;

233 }

234 

235 __forceinline__ __device__ __host__ float normal_distribution(unsigned long long int x)

236 {

237  float p = ::rocrand_device::detail::uniform_distribution(x);

238  float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);

239  return v;

240 }

241 

242 __forceinline__ __device__ __host__ float2 normal_distribution2(unsigned int v1, unsigned int v2)

243 {

244  return ::rocrand_device::detail::box_muller(v1, v2);

245 }

246 

247 __forceinline__ __device__ __host__ float2 normal_distribution2(uint2 v)

248 {

249  return ::rocrand_device::detail::box_muller(v.x, v.y);

250 }

251 

252 __forceinline__ __device__ __host__ float2 normal_distribution2(unsigned long long v)

253 {

254  return ::rocrand_device::detail::box_muller(v);

255 }

256 

257 __forceinline__ __device__ __host__ float4 normal_distribution4(uint4 v)

258 {

259  float2 r1 = ::rocrand_device::detail::box_muller(v.x, v.y);

260  float2 r2 = ::rocrand_device::detail::box_muller(v.z, v.w);

261  return float4{

262  r1.x,

263  r1.y,

264  r2.x,

265  r2.y

266  };

267 }

268 

269 __forceinline__ __device__ __host__ float4 normal_distribution4(longlong2 v)

270 {

271  float2 r1 = ::rocrand_device::detail::box_muller(v.x);

272  float2 r2 = ::rocrand_device::detail::box_muller(v.y);

273  return float4{r1.x, r1.y, r2.x, r2.y};

274 }

275 

276 __forceinline__ __device__ __host__ float4 normal_distribution4(unsigned long long v1,

277  unsigned long long v2)

278 {

279  float2 r1 = ::rocrand_device::detail::box_muller(v1);

280  float2 r2 = ::rocrand_device::detail::box_muller(v2);

281  return float4{r1.x, r1.y, r2.x, r2.y};

282 }

283 

284 __forceinline__ __device__ __host__ double normal_distribution_double(unsigned int x)

285 {

286  double p = ::rocrand_device::detail::uniform_distribution_double(x);

287  double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);

288  return v;

289 }

290 

291 __forceinline__ __device__ __host__ double normal_distribution_double(unsigned long long int x)

292 {

293  double p = ::rocrand_device::detail::uniform_distribution_double(x);

294  double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);

295  return v;

296 }

297 

298 __forceinline__ __device__ __host__ double2 normal_distribution_double2(uint4 v)

299 {

300  return ::rocrand_device::detail::box_muller_double(v);

301 }

302 

303 __forceinline__ __device__ __host__ double2 normal_distribution_double2(ulonglong2 v)

304 {

305  return ::rocrand_device::detail::box_muller_double(v);

306 }

307 

308 __forceinline__ __device__ __host__ __half2 normal_distribution_half2(unsigned int v)

309 {

310  return ::rocrand_device::detail::box_muller_half(

311  static_cast<unsigned short>(v),

312  static_cast<unsigned short>(v >> 16)

313  );

314 }

315 

316 __forceinline__ __device__ __host__ __half2 normal_distribution_half2(unsigned long long v)

317 {

318  return ::rocrand_device::detail::box_muller_half(static_cast<unsigned short>(v),

319  static_cast<unsigned short>(v >> 32));

320 }

321 

322 template<typename state_type>

323 __forceinline__ __device__ __host__ float2 mrg_normal_distribution2(unsigned int v1,

324  unsigned int v2)

325 {

326  return ::rocrand_device::detail::mrg_box_muller<state_type>(v1, v2);

327 }

328 

329 template<typename state_type>

330 __forceinline__ __device__ __host__ double2 mrg_normal_distribution_double2(unsigned int v1,

331  unsigned int v2)

332 {

333  return ::rocrand_device::detail::mrg_box_muller_double<state_type>(v1, v2);

334 }

335 

336 template<typename state_type>

337 __forceinline__ __device__ __host__ __half2 mrg_normal_distribution_half2(unsigned int v)

338 {

339  v = rocrand_device::detail::mrg_uniform_distribution_uint<state_type>(v);

340  return ::rocrand_device::detail::box_muller_half(

341  static_cast<unsigned short>(v),

342  static_cast<unsigned short>(v >> 16)

343  );

344 }

345 

346 } // end namespace detail

347 } // end namespace rocrand_device

348 

363 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

364 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_philox4x32_10* state)

365 {

366  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;

367 

368  if(bm_helper::has_float(state))

369  {

370  return bm_helper::get_float(state);

371  }

372 

373  auto state1 = rocrand(state);

374  auto state2 = rocrand(state);

375 

376  float2 r = rocrand_device::detail::normal_distribution2(state1, state2);

377  bm_helper::save_float(state, r.y);

378  return r.x;

379 }

380 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

381 

396 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_philox4x32_10* state)

397 {

398  auto state1 = rocrand(state);

399  auto state2 = rocrand(state);

400 

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

402 }

403 

418 __forceinline__ __device__ __host__ float4 rocrand_normal4(rocrand_state_philox4x32_10* state)

419 {

420  return rocrand_device::detail::normal_distribution4(rocrand4(state));

421 }

422 

437 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

438 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_philox4x32_10* state)

439 {

440  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;

441 

442  if(bm_helper::has_double(state))

443  {

444  return bm_helper::get_double(state);

445  }

446  double2 r = rocrand_device::detail::normal_distribution_double2(rocrand4(state));

447  bm_helper::save_double(state, r.y);

448  return r.x;

449 }

450 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

451 

466 __forceinline__ __device__ __host__ double2

467  rocrand_normal_double2(rocrand_state_philox4x32_10* state)

468 {

469  return rocrand_device::detail::normal_distribution_double2(rocrand4(state));

470 }

471 

486 __forceinline__ __device__ __host__ double4

487  rocrand_normal_double4(rocrand_state_philox4x32_10* state)

488 {

489  double2 r1, r2;

490  r1 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));

491  r2 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));

492  return double4 {

493  r1.x, r1.y, r2.x, r2.y

494  };

495 }

496 

511 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

512 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_mrg31k3p* state)

513 {

514  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;

515 

516  if(bm_helper::has_float(state))

517  {

518  return bm_helper::get_float(state);

519  }

520 

521  auto state1 = state->next();

522  auto state2 = state->next();

523 

524  float2 r

525  = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);

526  bm_helper::save_float(state, r.y);

527  return r.x;

528 }

529 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

530 

545 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_mrg31k3p* state)

546 {

547  auto state1 = state->next();

548  auto state2 = state->next();

549 

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

551 }

552 

567 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

568 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_mrg31k3p* state)

569 {

570  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;

571 

572  if(bm_helper::has_double(state))

573  {

574  return bm_helper::get_double(state);

575  }

576 

577  auto state1 = state->next();

578  auto state2 = state->next();

579 

580  double2 r

581  = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,

582  state2);

583  bm_helper::save_double(state, r.y);

584  return r.x;

585 }

586 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

587 

602 __forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_mrg31k3p* state)

603 {

604  auto state1 = state->next();

605  auto state2 = state->next();

606 

607  return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,

608  state2);

609 }

610 

625 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

626 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_mrg32k3a* state)

627 {

628  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;

629 

630  if(bm_helper::has_float(state))

631  {

632  return bm_helper::get_float(state);

633  }

634 

635  auto state1 = state->next();

636  auto state2 = state->next();

637 

638  float2 r

639  = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);

640  bm_helper::save_float(state, r.y);

641  return r.x;

642 }

643 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

644 

659 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_mrg32k3a* state)

660 {

661  auto state1 = state->next();

662  auto state2 = state->next();

663 

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

665 }

666 

681 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

682 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_mrg32k3a* state)

683 {

684  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;

685 

686  if(bm_helper::has_double(state))

687  {

688  return bm_helper::get_double(state);

689  }

690 

691  auto state1 = state->next();

692  auto state2 = state->next();

693 

694  double2 r

695  = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,

696  state2);

697  bm_helper::save_double(state, r.y);

698  return r.x;

699 }

700 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

701 

716 __forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_mrg32k3a* state)

717 {

718  auto state1 = state->next();

719  auto state2 = state->next();

720 

721  return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,

722  state2);

723 }

724 

739 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

740 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_xorwow* state)

741 {

742  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;

743 

744  if(bm_helper::has_float(state))

745  {

746  return bm_helper::get_float(state);

747  }

748  auto state1 = rocrand(state);

749  auto state2 = rocrand(state);

750  float2 r = rocrand_device::detail::normal_distribution2(state1, state2);

751  bm_helper::save_float(state, r.y);

752  return r.x;

753 }

754 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

755 

770 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_xorwow* state)

771 {

772  auto state1 = rocrand(state);

773  auto state2 = rocrand(state);

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

775 }

776 

791 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

792 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_xorwow* state)

793 {

794  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;

795 

796  if(bm_helper::has_double(state))

797  {

798  return bm_helper::get_double(state);

799  }

800 

801  auto state1 = rocrand(state);

802  auto state2 = rocrand(state);

803  auto state3 = rocrand(state);

804  auto state4 = rocrand(state);

805 

806  double2 r = rocrand_device::detail::normal_distribution_double2(

807  uint4 { state1, state2, state3, state4 }

808  );

809  bm_helper::save_double(state, r.y);

810  return r.x;

811 }

812 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

813 

828 __forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_xorwow* state)

829 {

830  auto state1 = rocrand(state);

831  auto state2 = rocrand(state);

832  auto state3 = rocrand(state);

833  auto state4 = rocrand(state);

834 

835  return rocrand_device::detail::normal_distribution_double2(

836  uint4 { state1, state2, state3, state4 }

837  );

838 }

839 

852 __forceinline__ __device__ float rocrand_normal(rocrand_state_mtgp32* state)

853 {

854  return rocrand_device::detail::normal_distribution(rocrand(state));

855 }

856 

871 __forceinline__ __device__ float2 rocrand_normal2(rocrand_state_mtgp32* state)

872 {

873  auto state1 = rocrand(state);

874  auto state2 = rocrand(state);

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

876 }

877 

890 __forceinline__ __device__ double rocrand_normal_double(rocrand_state_mtgp32* state)

891 {

892  return rocrand_device::detail::normal_distribution_double(rocrand(state));

893 }

894 

909 __forceinline__ __device__ double2 rocrand_normal_double2(rocrand_state_mtgp32* state)

910 {

911  auto state1 = rocrand(state);

912  auto state2 = rocrand(state);

913  auto state3 = rocrand(state);

914  auto state4 = rocrand(state);

915 

916  return rocrand_device::detail::normal_distribution_double2(

917  uint4{state1, state2, state3, state4});

918 }

919 

932 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_sobol32* state)

933 {

934  return rocrand_device::detail::normal_distribution(rocrand(state));

935 }

936 

949 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_sobol32* state)

950 {

951  return rocrand_device::detail::normal_distribution_double(rocrand(state));

952 }

953 

966 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_scrambled_sobol32* state)

967 {

968  return rocrand_device::detail::normal_distribution(rocrand(state));

969 }

970 

983 __forceinline__ __device__ __host__ double

984  rocrand_normal_double(rocrand_state_scrambled_sobol32* state)

985 {

986  return rocrand_device::detail::normal_distribution_double(rocrand(state));

987 }

988 

1001 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_sobol64* state)

1002 {

1003  return rocrand_device::detail::normal_distribution(rocrand(state));

1004 }

1005 

1018 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_sobol64* state)

1019 {

1020  return rocrand_device::detail::normal_distribution_double(rocrand(state));

1021 }

1022 

1035 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_scrambled_sobol64* state)

1036 {

1037  return rocrand_device::detail::normal_distribution(rocrand(state));

1038 }

1039 

1052 __forceinline__ __device__ __host__ double

1053  rocrand_normal_double(rocrand_state_scrambled_sobol64* state)

1054 {

1055  return rocrand_device::detail::normal_distribution_double(rocrand(state));

1056 }

1057 

1070 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_lfsr113* state)

1071 {

1072  return rocrand_device::detail::normal_distribution(rocrand(state));

1073 }

1074 

1089 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_lfsr113* state)

1090 {

1091  auto state1 = rocrand(state);

1092  auto state2 = rocrand(state);

1093 

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

1095 }

1096 

1109 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_lfsr113* state)

1110 {

1111  return rocrand_device::detail::normal_distribution_double(rocrand(state));

1112 }

1113 

1128 __forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_lfsr113* state)

1129 {

1130  auto state1 = rocrand(state);

1131  auto state2 = rocrand(state);

1132  auto state3 = rocrand(state);

1133  auto state4 = rocrand(state);

1134 

1135  return rocrand_device::detail::normal_distribution_double2(

1136  uint4{state1, state2, state3, state4});

1137 }

1138 

1151 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_threefry2x32_20* state)

1152 {

1153  return rocrand_device::detail::normal_distribution(rocrand(state));

1154 }

1155 

1170 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_threefry2x32_20* state)

1171 {

1172  return rocrand_device::detail::normal_distribution2(rocrand2(state));

1173 }

1174 

1187 __forceinline__ __device__ __host__ double

1188  rocrand_normal_double(rocrand_state_threefry2x32_20* state)

1189 {

1190  return rocrand_device::detail::normal_distribution_double(rocrand(state));

1191 }

1192 

1207 __forceinline__ __device__ __host__ double2

1208  rocrand_normal_double2(rocrand_state_threefry2x32_20* state)

1209 {

1210  auto state1 = rocrand2(state);

1211  auto state2 = rocrand2(state);

1212 

1213  return rocrand_device::detail::normal_distribution_double2(

1214  uint4{state1.x, state1.y, state2.x, state2.y});

1215 }

1216 

1229 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_threefry2x64_20* state)

1230 {

1231  return rocrand_device::detail::normal_distribution(rocrand(state));

1232 }

1233 

1248 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_threefry2x64_20* state)

1249 {

1250  return rocrand_device::detail::normal_distribution2(rocrand(state));

1251 }

1252 

1265 __forceinline__ __device__ __host__ double

1266  rocrand_normal_double(rocrand_state_threefry2x64_20* state)

1267 {

1268  return rocrand_device::detail::normal_distribution_double(rocrand(state));

1269 }

1270 

1285 __forceinline__ __device__ __host__ double2

1286  rocrand_normal_double2(rocrand_state_threefry2x64_20* state)

1287 {

1288  return rocrand_device::detail::normal_distribution_double2(rocrand2(state));

1289 }

1290 

1303 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_threefry4x32_20* state)

1304 {

1305  return rocrand_device::detail::normal_distribution(rocrand(state));

1306 }

1307 

1322 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_threefry4x32_20* state)

1323 {

1324  auto state1 = rocrand(state);

1325  auto state2 = rocrand(state);

1326 

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

1328 }

1329 

1342 __forceinline__ __device__ __host__ double

1343  rocrand_normal_double(rocrand_state_threefry4x32_20* state)

1344 {

1345  return rocrand_device::detail::normal_distribution_double(rocrand(state));

1346 }

1347 

1362 __forceinline__ __device__ __host__ double2

1363  rocrand_normal_double2(rocrand_state_threefry4x32_20* state)

1364 {

1365  return rocrand_device::detail::normal_distribution_double2(rocrand4(state));

1366 }

1367 

1380 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_threefry4x64_20* state)

1381 {

1382  return rocrand_device::detail::normal_distribution(rocrand(state));

1383 }

1384 

1399 __forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_threefry4x64_20* state)

1400 {

1401  auto state1 = rocrand(state);

1402  auto state2 = rocrand(state);

1403 

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

1405 }

1406 

1419 __forceinline__ __device__ __host__ double

1420  rocrand_normal_double(rocrand_state_threefry4x64_20* state)

1421 {

1422  return rocrand_device::detail::normal_distribution_double(rocrand(state));

1423 }

1424 

1439 __forceinline__ __device__ __host__ double2

1440  rocrand_normal_double2(rocrand_state_threefry4x64_20* state)

1441 {

1442  auto state1 = rocrand(state);

1443  auto state2 = rocrand(state);

1444 

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

1446 }

1447  // end of group rocranddevice

1449 

1450 #endif // ROCRAND_NORMAL_H_

rocrand_normal_double4

__forceinline__ __device__ __host__ double4 rocrand_normal_double4(rocrand_state_philox4x32_10 *state)

Returns four normally distributed double values.

Definition: rocrand_normal.h:487

rocrand_normal_double2

__forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_philox4x32_10 *state)

Returns two normally distributed double values.

Definition: rocrand_normal.h:467

rocrand_normal

__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_philox4x32_10 *state)

Returns a normally distributed float value.

Definition: rocrand_normal.h:364

rocrand4

__forceinline__ __device__ __host__ uint4 rocrand4(rocrand_state_philox4x32_10 *state)

Returns four uniformly distributed random unsigned int values from [0; 2^32 - 1] range.

Definition: rocrand_philox4x32_10.h:386

rocrand_normal_double

__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_philox4x32_10 *state)

Returns a normally distributed double value.

Definition: rocrand_normal.h:438

rocrand_normal4

__forceinline__ __device__ __host__ float4 rocrand_normal4(rocrand_state_philox4x32_10 *state)

Returns four normally distributed float values.

Definition: rocrand_normal.h:418

rocrand

__forceinline__ __device__ __host__ unsigned int rocrand(rocrand_state_lfsr113 *state)

Returns uniformly distributed random unsigned int value from [0; 2^32 - 1] range.

Definition: rocrand_lfsr113.h:274

rocrand_normal2

__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_philox4x32_10 *state)

Returns two normally distributed float values.

Definition: rocrand_normal.h:396