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

• library
• include
• rocrand

rocrand_normal.h

1 // Copyright (c) 2017-2025 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 "rocrand/rocrand_lfsr113.h"

30 #include "rocrand/rocrand_mrg31k3p.h"

31 #include "rocrand/rocrand_mrg32k3a.h"

32 #include "rocrand/rocrand_mtgp32.h"

33 #include "rocrand/rocrand_philox4x32_10.h"

34 #include "rocrand/rocrand_scrambled_sobol32.h"

35 #include "rocrand/rocrand_scrambled_sobol64.h"

36 #include "rocrand/rocrand_sobol32.h"

37 #include "rocrand/rocrand_sobol64.h"

38 #include "rocrand/rocrand_threefry2x32_20.h"

39 #include "rocrand/rocrand_threefry2x64_20.h"

40 #include "rocrand/rocrand_threefry4x32_20.h"

41 #include "rocrand/rocrand_threefry4x64_20.h"

42 #include "rocrand/rocrand_xorwow.h"

43 

44 #include "rocrand/rocrand_uniform.h"

45 

46 #include <hip/hip_runtime.h>

47 

48 #include <math.h>

49 

50 namespace rocrand_device {

51 namespace detail {

52 

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

54 {

55  float2 result;

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

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

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

59  #ifdef __HIP_DEVICE_COMPILE__

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

61  result.x *= s;

62  result.y *= s;

63  #else

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

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

66  #endif

67  return result;

68 }

69 

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

71 {

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

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

74 

75  return box_muller(x, y);

76 }

77 

78 __forceinline__ __device__ __host__ double2 box_muller_double(uint4 v)

79 {

80  double2 result;

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

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

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

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

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

86  double w = (ROCRAND_2POW53_INV_DOUBLE * 2.0) +

87  (v2 * (ROCRAND_2POW53_INV_DOUBLE * 2.0));

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

89  #ifdef __HIP_DEVICE_COMPILE__

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

91  result.x *= s;

92  result.y *= s;

93  #else

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

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

96  #endif

97  return result;

98 }

99 

100 __forceinline__ __device__ __host__ double2 box_muller_double(ulonglong2 v)

101 {

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

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

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

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

106 

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

108 }

109 

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

111 {

112  #if defined(ROCRAND_HALF_MATH_SUPPORTED)

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

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

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

116  return __half2 {

117  __hmul(hsin(v), s),

118  __hmul(hcos(v), s)

119  };

120  #else

121  float2 r;

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

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

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

125  #ifdef __HIP_DEVICE_COMPILE__

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

127  r.x *= s;

128  r.y *= s;

129  #else

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

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

132  #endif

133  return __half2 {

134  __float2half(r.x),

135  __float2half(r.y)

136  };

137  #endif

138 }

139 

140 template<typename state_type>

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

142 {

143  float2 result;

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

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

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

147  #ifdef __HIP_DEVICE_COMPILE__

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

149  result.x *= s;

150  result.y *= s;

151  #else

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

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

154  #endif

155  return result;

156 }

157 

158 template<typename state_type>

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

160 {

161  double2 result;

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

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

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

165  #ifdef __HIP_DEVICE_COMPILE__

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

167  result.x *= s;

168  result.y *= s;

169  #else

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

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

172  #endif

173  return result;

174 }

175 

176 __forceinline__ __device__ __host__ float roc_f_erfinv(float x)

177 {

178  float tt1, tt2, lnx, sgn;

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

180 

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

182  lnx = logf(x);

183 

184  #ifdef __HIP_DEVICE_COMPILE__

185  if (isnan(lnx))

186  #else

187  if (std::isnan(lnx))

188  #endif

189  return 1.0f;

190  #ifdef __HIP_DEVICE_COMPILE__

191  else if (isinf(lnx))

192  #else

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

194  #endif

195  return 0.0f;

196 

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

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

199 

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

201 }

202 

203 __forceinline__ __device__ __host__ double roc_d_erfinv(double x)

204 {

205  double tt1, tt2, lnx, sgn;

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

207 

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

209  lnx = log(x);

210 

211  #ifdef __HIP_DEVICE_COMPILE__

212  if (isnan(lnx))

213  #else

214  if (std::isnan(lnx))

215  #endif

216  return 1.0;

217  #ifdef __HIP_DEVICE_COMPILE__

218  else if (isinf(lnx))

219  #else

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

221  #endif

222  return 0.0;

223 

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

225  tt2 = 1.0 / (0.147) * lnx;

226 

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

228 }

229 

230 __forceinline__ __device__ __host__ float normal_distribution(unsigned int x)

231 {

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

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

234  return v;

235 }

236 

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

238 {

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

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

241  return v;

242 }

243 

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

245 {

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

247 }

248 

249 __forceinline__ __device__ __host__ float2 normal_distribution2(uint2 v)

250 {

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

252 }

253 

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

255 {

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

257 }

258 

259 __forceinline__ __device__ __host__ float4 normal_distribution4(uint4 v)

260 {

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

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

263  return float4{

264  r1.x,

265  r1.y,

266  r2.x,

267  r2.y

268  };

269 }

270 

271 __forceinline__ __device__ __host__ float4 normal_distribution4(longlong2 v)

272 {

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

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

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

276 }

277 

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

279  unsigned long long v2)

280 {

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

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

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

284 }

285 

286 __forceinline__ __device__ __host__ double normal_distribution_double(unsigned int x)

287 {

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

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

290  return v;

291 }

292 

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

294 {

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

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

297  return v;

298 }

299 

300 __forceinline__ __device__ __host__ double2 normal_distribution_double2(uint4 v)

301 {

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

303 }

304 

305 __forceinline__ __device__ __host__ double2 normal_distribution_double2(ulonglong2 v)

306 {

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

308 }

309 

310 __forceinline__ __device__ __host__ __half2 normal_distribution_half2(unsigned int v)

311 {

312  return ::rocrand_device::detail::box_muller_half(

313  static_cast<unsigned short>(v),

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

315  );

316 }

317 

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

319 {

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

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

322 }

323 

324 template<typename state_type>

325 __forceinline__ __device__ __host__ float2 mrg_normal_distribution2(unsigned int v1,

326  unsigned int v2)

327 {

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

329 }

330 

331 template<typename state_type>

332 __forceinline__ __device__ __host__ double2 mrg_normal_distribution_double2(unsigned int v1,

333  unsigned int v2)

334 {

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

336 }

337 

338 template<typename state_type>

339 __forceinline__ __device__ __host__ __half2 mrg_normal_distribution_half2(unsigned int v)

340 {

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

342  return ::rocrand_device::detail::box_muller_half(

343  static_cast<unsigned short>(v),

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

345  );

346 }

347 

348 } // end namespace detail

349 } // end namespace rocrand_device

350 

365 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

366 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_philox4x32_10* state)

367 {

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

369 

370  if(bm_helper::has_float(state))

371  {

372  return bm_helper::get_float(state);

373  }

374 

375  auto state1 = rocrand(state);

376  auto state2 = rocrand(state);

377 

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

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

380  return r.x;

381 }

382 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

383 

398 __forceinline__ __device__ __host__

399 float2 rocrand_normal2(rocrand_state_philox4x32_10* state)

400 {

401  auto state1 = rocrand(state);

402  auto state2 = rocrand(state);

403 

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

405 }

406 

421 __forceinline__ __device__ __host__

422 float4 rocrand_normal4(rocrand_state_philox4x32_10* state)

423 {

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

425 }

426 

441 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

442 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_philox4x32_10* state)

443 {

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

445 

446  if(bm_helper::has_double(state))

447  {

448  return bm_helper::get_double(state);

449  }

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

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

452  return r.x;

453 }

454 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

455 

470 __forceinline__ __device__ __host__

471 double2 rocrand_normal_double2(rocrand_state_philox4x32_10* state)

472 {

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

474 }

475 

490 __forceinline__ __device__ __host__

491 double4 rocrand_normal_double4(rocrand_state_philox4x32_10* state)

492 {

493  double2 r1, r2;

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

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

496  return double4 {

497  r1.x, r1.y, r2.x, r2.y

498  };

499 }

500 

515 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

516 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_mrg31k3p* state)

517 {

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

519 

520  if(bm_helper::has_float(state))

521  {

522  return bm_helper::get_float(state);

523  }

524 

525  auto state1 = state->next();

526  auto state2 = state->next();

527 

528  float2 r

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

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

531  return r.x;

532 }

533 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

534 

549 __forceinline__ __device__ __host__

550 float2 rocrand_normal2(rocrand_state_mrg31k3p* state)

551 {

552  auto state1 = state->next();

553  auto state2 = state->next();

554 

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

556 }

557 

572 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

573 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_mrg31k3p* state)

574 {

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

576 

577  if(bm_helper::has_double(state))

578  {

579  return bm_helper::get_double(state);

580  }

581 

582  auto state1 = state->next();

583  auto state2 = state->next();

584 

585  double2 r

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

587  state2);

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

589  return r.x;

590 }

591 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

592 

607 __forceinline__ __device__ __host__

608 double2 rocrand_normal_double2(rocrand_state_mrg31k3p* state)

609 {

610  auto state1 = state->next();

611  auto state2 = state->next();

612 

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

614  state2);

615 }

616 

631 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

632 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_mrg32k3a* state)

633 {

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

635 

636  if(bm_helper::has_float(state))

637  {

638  return bm_helper::get_float(state);

639  }

640 

641  auto state1 = state->next();

642  auto state2 = state->next();

643 

644  float2 r

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

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

647  return r.x;

648 }

649 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

650 

665 __forceinline__ __device__ __host__

666 float2 rocrand_normal2(rocrand_state_mrg32k3a* state)

667 {

668  auto state1 = state->next();

669  auto state2 = state->next();

670 

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

672 }

673 

688 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

689 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_mrg32k3a* state)

690 {

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

692 

693  if(bm_helper::has_double(state))

694  {

695  return bm_helper::get_double(state);

696  }

697 

698  auto state1 = state->next();

699  auto state2 = state->next();

700 

701  double2 r

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

703  state2);

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

705  return r.x;

706 }

707 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

708 

723 __forceinline__ __device__ __host__

724 double2 rocrand_normal_double2(rocrand_state_mrg32k3a* state)

725 {

726  auto state1 = state->next();

727  auto state2 = state->next();

728 

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

730  state2);

731 }

732 

747 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

748 __forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_xorwow* state)

749 {

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

751 

752  if(bm_helper::has_float(state))

753  {

754  return bm_helper::get_float(state);

755  }

756  auto state1 = rocrand(state);

757  auto state2 = rocrand(state);

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

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

760  return r.x;

761 }

762 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

763 

778 __forceinline__ __device__ __host__

779 float2 rocrand_normal2(rocrand_state_xorwow* state)

780 {

781  auto state1 = rocrand(state);

782  auto state2 = rocrand(state);

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

784 }

785 

800 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE

801 __forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_xorwow* state)

802 {

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

804 

805  if(bm_helper::has_double(state))

806  {

807  return bm_helper::get_double(state);

808  }

809 

810  auto state1 = rocrand(state);

811  auto state2 = rocrand(state);

812  auto state3 = rocrand(state);

813  auto state4 = rocrand(state);

814 

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

816  uint4 { state1, state2, state3, state4 }

817  );

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

819  return r.x;

820 }

821 #endif // ROCRAND_DETAIL_BM_NOT_IN_STATE

822 

837 __forceinline__ __device__ __host__

838 double2 rocrand_normal_double2(rocrand_state_xorwow* state)

839 {

840  auto state1 = rocrand(state);

841  auto state2 = rocrand(state);

842  auto state3 = rocrand(state);

843  auto state4 = rocrand(state);

844 

845  return rocrand_device::detail::normal_distribution_double2(

846  uint4 { state1, state2, state3, state4 }

847  );

848 }

849 

862 __forceinline__ __device__

863 float rocrand_normal(rocrand_state_mtgp32* state)

864 {

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

866 }

867 

882 __forceinline__ __device__

883 float2 rocrand_normal2(rocrand_state_mtgp32* state)

884 {

885  auto state1 = rocrand(state);

886  auto state2 = rocrand(state);

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

888 }

889 

902 __forceinline__ __device__

903 double rocrand_normal_double(rocrand_state_mtgp32* state)

904 {

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

906 }

907 

922 __forceinline__ __device__

923 double2 rocrand_normal_double2(rocrand_state_mtgp32* state)

924 {

925  auto state1 = rocrand(state);

926  auto state2 = rocrand(state);

927  auto state3 = rocrand(state);

928  auto state4 = rocrand(state);

929 

930  return rocrand_device::detail::normal_distribution_double2(

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

932 }

933 

946 __forceinline__ __device__ __host__

947 float rocrand_normal(rocrand_state_sobol32* state)

948 {

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

950 }

951 

964 __forceinline__ __device__ __host__

965 double rocrand_normal_double(rocrand_state_sobol32* state)

966 {

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

968 }

969 

982 __forceinline__ __device__ __host__

983 float rocrand_normal(rocrand_state_scrambled_sobol32* state)

984 {

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

986 }

987 

1000 __forceinline__ __device__ __host__

1001 double rocrand_normal_double(rocrand_state_scrambled_sobol32* state)

1002 {

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

1004 }

1005 

1018 __forceinline__ __device__ __host__

1019 float rocrand_normal(rocrand_state_sobol64* state)

1020 {

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

1022 }

1023 

1036 __forceinline__ __device__ __host__

1037 double rocrand_normal_double(rocrand_state_sobol64* state)

1038 {

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

1040 }

1041 

1054 __forceinline__ __device__ __host__

1055 float rocrand_normal(rocrand_state_scrambled_sobol64* state)

1056 {

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

1058 }

1059 

1072 __forceinline__ __device__ __host__

1073 double rocrand_normal_double(rocrand_state_scrambled_sobol64* state)

1074 {

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

1076 }

1077 

1090 __forceinline__ __device__ __host__

1091 float rocrand_normal(rocrand_state_lfsr113* state)

1092 {

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

1094 }

1095 

1110 __forceinline__ __device__ __host__

1111 float2 rocrand_normal2(rocrand_state_lfsr113* state)

1112 {

1113  auto state1 = rocrand(state);

1114  auto state2 = rocrand(state);

1115 

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

1117 }

1118 

1131 __forceinline__ __device__ __host__

1132 double rocrand_normal_double(rocrand_state_lfsr113* state)

1133 {

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

1135 }

1136 

1151 __forceinline__ __device__ __host__

1152 double2 rocrand_normal_double2(rocrand_state_lfsr113* state)

1153 {

1154  auto state1 = rocrand(state);

1155  auto state2 = rocrand(state);

1156  auto state3 = rocrand(state);

1157  auto state4 = rocrand(state);

1158 

1159  return rocrand_device::detail::normal_distribution_double2(

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

1161 }

1162 

1175 __forceinline__ __device__ __host__

1176 float rocrand_normal(rocrand_state_threefry2x32_20* state)

1177 {

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

1179 }

1180 

1195 __forceinline__ __device__ __host__

1196 float2 rocrand_normal2(rocrand_state_threefry2x32_20* state)

1197 {

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

1199 }

1200 

1213 __forceinline__ __device__ __host__

1214 double rocrand_normal_double(rocrand_state_threefry2x32_20* state)

1215 {

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

1217 }

1218 

1233 __forceinline__ __device__ __host__

1234 double2 rocrand_normal_double2(rocrand_state_threefry2x32_20* state)

1235 {

1236  auto state1 = rocrand2(state);

1237  auto state2 = rocrand2(state);

1238 

1239  return rocrand_device::detail::normal_distribution_double2(

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

1241 }

1242 

1255 __forceinline__ __device__ __host__

1256 float rocrand_normal(rocrand_state_threefry2x64_20* state)

1257 {

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

1259 }

1260 

1275 __forceinline__ __device__ __host__

1276 float2 rocrand_normal2(rocrand_state_threefry2x64_20* state)

1277 {

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

1279 }

1280 

1293 __forceinline__ __device__ __host__

1294 double rocrand_normal_double(rocrand_state_threefry2x64_20* state)

1295 {

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

1297 }

1298 

1313 __forceinline__ __device__ __host__

1314 double2 rocrand_normal_double2(rocrand_state_threefry2x64_20* state)

1315 {

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

1317 }

1318 

1331 __forceinline__ __device__ __host__

1332 float rocrand_normal(rocrand_state_threefry4x32_20* state)

1333 {

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

1335 }

1336 

1351 __forceinline__ __device__ __host__

1352 float2 rocrand_normal2(rocrand_state_threefry4x32_20* state)

1353 {

1354  auto state1 = rocrand(state);

1355  auto state2 = rocrand(state);

1356 

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

1358 }

1359 

1372 __forceinline__ __device__ __host__

1373 double rocrand_normal_double(rocrand_state_threefry4x32_20* state)

1374 {

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

1376 }

1377 

1392 __forceinline__ __device__ __host__

1393 double2 rocrand_normal_double2(rocrand_state_threefry4x32_20* state)

1394 {

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

1396 }

1397 

1410 __forceinline__ __device__ __host__

1411 float rocrand_normal(rocrand_state_threefry4x64_20* state)

1412 {

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

1414 }

1415 

1430 __forceinline__ __device__ __host__

1431 float2 rocrand_normal2(rocrand_state_threefry4x64_20* state)

1432 {

1433  auto state1 = rocrand(state);

1434  auto state2 = rocrand(state);

1435 

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

1437 }

1438 

1451 __forceinline__ __device__ __host__

1452 double rocrand_normal_double(rocrand_state_threefry4x64_20* state)

1453 {

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

1455 }

1456 

1471 __forceinline__ __device__ __host__

1472 double2 rocrand_normal_double2(rocrand_state_threefry4x64_20* state)

1473 {

1474  auto state1 = rocrand(state);

1475  auto state2 = rocrand(state);

1476 

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

1478 }

1479  // end of group rocranddevice

1481 

1482 #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:491

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:471

rocrand_normal

__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_philox4x32_10 *state)

Returns a normally distributed float value.

Definition: rocrand_normal.h:366

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:379

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:442

rocrand_normal4

__forceinline__ __device__ __host__ float4 rocrand_normal4(rocrand_state_philox4x32_10 *state)

Returns four normally distributed float values.

Definition: rocrand_normal.h:422

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:277

rocrand_normal2

__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_philox4x32_10 *state)

Returns two normally distributed float values.

Definition: rocrand_normal.h:399