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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/docs-6.2.1/library/include/rocrand/rocrand_uniform.h Source File#

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/docs-6.2.1/library/include/rocrand/rocrand_uniform.h Source File
API library
rocrand_uniform.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 
26 #ifndef ROCRAND_UNIFORM_H_
27 #define ROCRAND_UNIFORM_H_
28 
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_common.h"
45 
46 namespace rocrand_device {
47 namespace detail {
48 
49 struct two_uints
50 {
51  unsigned int x;
52  unsigned int y;
53 };
54 
55 union two_uints_to_ulong
56 {
57  two_uints uint2_value;
58  unsigned long long int ulong_value;
59 };
60 
61 // For unsigned integer between 0 and UINT_MAX, returns value between
62 // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
63 __forceinline__ __device__ __host__ float uniform_distribution(unsigned int v)
64 {
65  return ROCRAND_2POW32_INV + (v * ROCRAND_2POW32_INV);
66 }
67 
68 // For unsigned integer between 0 and ULLONG_MAX, returns value between
69 // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
70 __forceinline__ __device__ __host__ float uniform_distribution(unsigned long long int v)
71 {
72  return ROCRAND_2POW32_INV + (v >> 32) * ROCRAND_2POW32_INV;
73 }
74 
75 __forceinline__ __device__ __host__ float4 uniform_distribution4(uint4 v)
76 {
77  return float4 {
78  ROCRAND_2POW32_INV + (v.x * ROCRAND_2POW32_INV),
79  ROCRAND_2POW32_INV + (v.y * ROCRAND_2POW32_INV),
80  ROCRAND_2POW32_INV + (v.z * ROCRAND_2POW32_INV),
81  ROCRAND_2POW32_INV + (v.w * ROCRAND_2POW32_INV)
82  };
83 }
84 
85 __forceinline__ __device__ __host__ float4 uniform_distribution4(ulonglong4 v)
86 {
87  return float4{ROCRAND_2POW64_INV + (v.x * ROCRAND_2POW64_INV),
88  ROCRAND_2POW64_INV + (v.y * ROCRAND_2POW64_INV),
89  ROCRAND_2POW64_INV + (v.z * ROCRAND_2POW64_INV),
90  ROCRAND_2POW64_INV + (v.w * ROCRAND_2POW64_INV)};
91 }
92 
93 // For unsigned integer between 0 and UINT_MAX, returns value between
94 // 0.0 and 1.0, excluding 0.0 and including 1.0.
95 __forceinline__ __device__ __host__ double uniform_distribution_double(unsigned int v)
96 {
97  return ROCRAND_2POW32_INV_DOUBLE + (v * ROCRAND_2POW32_INV_DOUBLE);
98 }
99 
100 __forceinline__ __device__ __host__ double uniform_distribution_double(unsigned int v1,
101  unsigned int v2)
102 {
103  two_uints_to_ulong v;
104  v.uint2_value.x = v1;
105  v.uint2_value.y = (v2 >> 11);
106  return ROCRAND_2POW53_INV_DOUBLE + (v.ulong_value * ROCRAND_2POW53_INV_DOUBLE);
107 }
108 
109 __forceinline__ __device__ __host__ double uniform_distribution_double(unsigned long long int v)
110 {
111  return ROCRAND_2POW53_INV_DOUBLE + (
112  // 2^53 is the biggest int that can be stored in double, such
113  // that it and all smaller integers can be stored in double
114  (v >> 11) * ROCRAND_2POW53_INV_DOUBLE
115  );
116 }
117 
118 __forceinline__ __device__ __host__ double2 uniform_distribution_double2(uint4 v)
119 {
120  return double2 {
121  uniform_distribution_double(v.x, v.y),
122  uniform_distribution_double(v.z, v.w)
123  };
124 }
125 
126 __forceinline__ __device__ __host__ double4 uniform_distribution_double4(uint4 v1, uint4 v2)
127 {
128  return double4 {
129  uniform_distribution_double(v1.x, v1.y),
130  uniform_distribution_double(v1.z, v1.w),
131  uniform_distribution_double(v2.x, v2.y),
132  uniform_distribution_double(v2.z, v2.w)
133  };
134 }
135 
136 __forceinline__ __device__ __host__ double2 uniform_distribution_double2(ulonglong2 v)
137 {
138  return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
139 }
140 
141 __forceinline__ __device__ __host__ double2 uniform_distribution_double2(ulonglong4 v)
142 {
143  return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
144 }
145 
146 __forceinline__ __device__ __host__ double4 uniform_distribution_double4(ulonglong4 v)
147 {
148  return double4{uniform_distribution_double(v.x),
149  uniform_distribution_double(v.z),
150  uniform_distribution_double(v.x),
151  uniform_distribution_double(v.z)};
152 }
153 
154 __forceinline__ __device__ __host__ __half uniform_distribution_half(unsigned short v)
155 {
156  return __float2half(ROCRAND_2POW16_INV + (v * ROCRAND_2POW16_INV));
157 }
158 
159 // For an unsigned integer produced by an MRG-based engine, returns a value
160 // in range [0, UINT32_MAX].
161 template<typename state_type>
162 __forceinline__ __device__ __host__ unsigned int mrg_uniform_distribution_uint(unsigned int v)
163  = delete;
164 
165 template<>
166 __forceinline__ __device__ __host__ unsigned int
167  mrg_uniform_distribution_uint<rocrand_state_mrg31k3p>(unsigned int v)
168 {
169  return static_cast<unsigned int>((v - 1) * ROCRAND_MRG31K3P_UINT32_NORM);
170 }
171 
172 template<>
173 __forceinline__ __device__ __host__ unsigned int
174  mrg_uniform_distribution_uint<rocrand_state_mrg32k3a>(unsigned int v)
175 {
176  return static_cast<unsigned int>((v - 1) * ROCRAND_MRG32K3A_UINT_NORM);
177 }
178 
179 // For an unsigned integer produced by an MRG-based engine, returns value between
180 // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
181 template<typename state_type>
182 __forceinline__ __device__ __host__ float mrg_uniform_distribution(unsigned int v) = delete;
183 
184 template<>
185 __forceinline__ __device__ __host__ float
186  mrg_uniform_distribution<rocrand_state_mrg31k3p>(unsigned int v)
187 {
188  double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
189  return static_cast<float>(ret);
190 }
191 
192 template<>
193 __forceinline__ __device__ __host__ float
194  mrg_uniform_distribution<rocrand_state_mrg32k3a>(unsigned int v)
195 {
196  double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
197  return static_cast<float>(ret);
198 }
199 
200 // For an unsigned integer produced by an MRG generator, returns value between
201 // 0.0 and 1.0, excluding 0.0 and including 1.0.
202 template<typename state_type>
203 __forceinline__ __device__ __host__ double mrg_uniform_distribution_double(unsigned int v) = delete;
204 
205 template<>
206 __forceinline__ __device__ __host__ double
207  mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(unsigned int v)
208 {
209  double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
210  return ret;
211 }
212 
213 template<>
214 __forceinline__ __device__ __host__ double
215  mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(unsigned int v)
216 {
217  double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
218  return ret;
219 }
220 
221 } // end namespace detail
222 } // end namespace rocrand_device
223 
236 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_philox4x32_10* state)
237 {
238  return rocrand_device::detail::uniform_distribution(rocrand(state));
239 }
240 
253 __forceinline__ __device__ __host__ float2 rocrand_uniform2(rocrand_state_philox4x32_10* state)
254 {
255  auto state1 = rocrand(state);
256  auto state2 = rocrand(state);
257 
258  return float2 {
259  rocrand_device::detail::uniform_distribution(state1),
260  rocrand_device::detail::uniform_distribution(state2)
261  };
262 }
263 
276 __forceinline__ __device__ __host__ float4 rocrand_uniform4(rocrand_state_philox4x32_10* state)
277 {
278  return rocrand_device::detail::uniform_distribution4(rocrand4(state));
279 }
280 
293 __forceinline__ __device__ __host__ double
294  rocrand_uniform_double(rocrand_state_philox4x32_10* state)
295 {
296  auto state1 = rocrand(state);
297  auto state2 = rocrand(state);
298 
299  return rocrand_device::detail::uniform_distribution_double(state1, state2);
300 }
301 
314 __forceinline__ __device__ __host__ double2
315  rocrand_uniform_double2(rocrand_state_philox4x32_10* state)
316 {
317  return rocrand_device::detail::uniform_distribution_double2(rocrand4(state));
318 }
319 
332 __forceinline__ __device__ __host__ double4
333  rocrand_uniform_double4(rocrand_state_philox4x32_10* state)
334 {
335  return rocrand_device::detail::uniform_distribution_double4(rocrand4(state), rocrand4(state));
336 }
337 
350 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_mrg31k3p* state)
351 {
352  return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg31k3p>(state->next());
353 }
354 
370 __forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_mrg31k3p* state)
371 {
372  return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(
373  state->next());
374 }
375 
388 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_mrg32k3a* state)
389 {
390  return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg32k3a>(state->next());
391 }
392 
408 __forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_mrg32k3a* state)
409 {
410  return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(
411  state->next());
412 }
413 
426 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_xorwow* state)
427 {
428  return rocrand_device::detail::uniform_distribution(rocrand(state));
429 }
430 
443 __forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_xorwow* state)
444 {
445  auto state1 = rocrand(state);
446  auto state2 = rocrand(state);
447 
448  return rocrand_device::detail::uniform_distribution_double(state1, state2);
449 }
450 
463 __forceinline__ __device__ float rocrand_uniform(rocrand_state_mtgp32* state)
464 {
465  return rocrand_device::detail::uniform_distribution(rocrand(state));
466 }
467 
483 __forceinline__ __device__ double rocrand_uniform_double(rocrand_state_mtgp32* state)
484 {
485  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
486 }
487 
500 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_sobol32* state)
501 {
502  return rocrand_device::detail::uniform_distribution(rocrand(state));
503 }
504 
520 __forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_sobol32* state)
521 {
522  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
523 }
524 
537 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_scrambled_sobol32* state)
538 {
539  return rocrand_device::detail::uniform_distribution(rocrand(state));
540 }
541 
557 __forceinline__ __device__ __host__ double
558  rocrand_uniform_double(rocrand_state_scrambled_sobol32* state)
559 {
560  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
561 }
562 
575 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_sobol64* state)
576 {
577  return rocrand_device::detail::uniform_distribution(rocrand(state));
578 }
579 
592 __forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_sobol64* state)
593 {
594  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
595 }
596 
609 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_scrambled_sobol64* state)
610 {
611  return rocrand_device::detail::uniform_distribution(rocrand(state));
612 }
613 
626 __forceinline__ __device__ __host__ double
627  rocrand_uniform_double(rocrand_state_scrambled_sobol64* state)
628 {
629  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
630 }
631 
644 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_lfsr113* state)
645 {
646  return rocrand_device::detail::uniform_distribution(rocrand(state));
647 }
648 
664 __forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_lfsr113* state)
665 {
666  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
667 }
668 
681 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_threefry2x32_20* state)
682 {
683  return rocrand_device::detail::uniform_distribution(rocrand(state));
684 }
685 
701 __forceinline__ __device__ __host__ double
702  rocrand_uniform_double(rocrand_state_threefry2x32_20* state)
703 {
704  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
705 }
706 
719 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_threefry2x64_20* state)
720 {
721  return rocrand_device::detail::uniform_distribution(rocrand(state));
722 }
723 
739 __forceinline__ __device__ __host__ double
740  rocrand_uniform_double(rocrand_state_threefry2x64_20* state)
741 {
742  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
743 }
744 
757 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_threefry4x32_20* state)
758 {
759  return rocrand_device::detail::uniform_distribution(rocrand(state));
760 }
761 
777 __forceinline__ __device__ __host__ double
778  rocrand_uniform_double(rocrand_state_threefry4x32_20* state)
779 {
780  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
781 }
782 
795 __forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_threefry4x64_20* state)
796 {
797  return rocrand_device::detail::uniform_distribution(rocrand(state));
798 }
799 
815 __forceinline__ __device__ __host__ double
816  rocrand_uniform_double(rocrand_state_threefry4x64_20* state)
817 {
818  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
819 }
820  // end of group rocranddevice
822 
823 #endif // ROCRAND_UNIFORM_H_
__forceinline__ __device__ __host__ float4 rocrand_uniform4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random float values from (0; 1] range.
Definition: rocrand_uniform.h:276
__forceinline__ __device__ __host__ double4 rocrand_uniform_double4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random double values from (0; 1] range.
Definition: rocrand_uniform.h:333
__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
__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random double value from (0; 1] range.
Definition: rocrand_uniform.h:294
__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random float value from (0; 1] range.
Definition: rocrand_uniform.h:236
__forceinline__ __device__ __host__ float2 rocrand_uniform2(rocrand_state_philox4x32_10 *state)
Returns two uniformly distributed random float values from (0; 1] range.
Definition: rocrand_uniform.h:253
__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
__forceinline__ __device__ __host__ double2 rocrand_uniform_double2(rocrand_state_philox4x32_10 *state)
Returns two uniformly distributed random double values from (0; 1] range.
Definition: rocrand_uniform.h:315