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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/docs-6.2.1/library/include/rocrand/rocrand_mtgp32.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_mtgp32.h Source File
API library
rocrand_mtgp32.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 /*
22  * Copyright (c) 2009, 2010 Mutsuo Saito, Makoto Matsumoto and Hiroshima
23  * University. All rights reserved.
24  * Copyright (c) 2011 Mutsuo Saito, Makoto Matsumoto, Hiroshima
25  * University and University of Tokyo. All rights reserved.
26  *
27  * Redistribution and use in source and binary forms, with or without
28  * modification, are permitted provided that the following conditions are
29  * met:
30  *
31  * * Redistributions of source code must retain the above copyright
32  * notice, this list of conditions and the following disclaimer.
33  * * Redistributions in binary form must reproduce the above
34  * copyright notice, this list of conditions and the following
35  * disclaimer in the documentation and/or other materials provided
36  * with the distribution.
37  * * Neither the name of the Hiroshima University nor the names of
38  * its contributors may be used to endorse or promote products
39  * derived from this software without specific prior written
40  * permission.
41  *
42  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
43  * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
44  * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
45  * A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
46  * OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
47  * SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
48  * LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
49  * DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
50  * THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
51  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
52  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
53  */
54 
55 #ifndef ROCRAND_MTGP32_H_
56 #define ROCRAND_MTGP32_H_
57 
58 #include <stdlib.h>
59 
60 #include "rocrand/rocrand.h"
61 #include "rocrand/rocrand_common.h"
62 
63 #define MTGP_MEXP 11213
64 #define MTGP_N 351
65 #define MTGP_FLOOR_2P 256
66 #define MTGP_CEIL_2P 512
67 #define MTGP_TN MTGP_FLOOR_2P
68 #define MTGP_LS (MTGP_TN * 3)
69 #define MTGP_BN_MAX 512
70 #define MTGP_TS 16
71 #define MTGP_STATE 1024
72 #define MTGP_MASK 1023
73 
74 // Source: https://github.com/MersenneTwister-Lab/MTGP/blob/master/mtgp32-fast.h
96  int mexp;
97  int pos;
98  int sh1;
99  int sh2;
100  uint32_t tbl[16];
101  uint32_t tmp_tbl[16];
102  uint32_t flt_tmp_tbl[16];
103  uint32_t mask;
104  unsigned char poly_sha1[21];
105 };
106 
107 namespace rocrand_device {
108 
109 struct mtgp32_params
110 {
111  unsigned int pos_tbl[MTGP_BN_MAX];
112  unsigned int param_tbl[MTGP_BN_MAX][MTGP_TS];
113  unsigned int temper_tbl[MTGP_BN_MAX][MTGP_TS];
114  unsigned int single_temper_tbl[MTGP_BN_MAX][MTGP_TS];
115  unsigned int sh1_tbl[MTGP_BN_MAX];
116  unsigned int sh2_tbl[MTGP_BN_MAX];
117  unsigned int mask[1];
118 };
119 
120 typedef mtgp32_params_fast_t mtgp32_fast_params;
121 
122 struct mtgp32_state
123 {
124  int offset;
125  int id;
126  unsigned int status[MTGP_STATE];
127 };
128 
129 inline
130 void rocrand_mtgp32_init_state(unsigned int array[],
131  const mtgp32_fast_params *para, unsigned int seed)
132 {
133  int i;
134  int size = para->mexp / 32 + 1;
135  unsigned int hidden_seed;
136  unsigned int tmp;
137  hidden_seed = para->tbl[4] ^ (para->tbl[8] << 16);
138  tmp = hidden_seed;
139  tmp += tmp >> 16;
140  tmp += tmp >> 8;
141  memset(array, tmp & 0xff, sizeof(unsigned int) * size);
142  array[0] = seed;
143  array[1] = hidden_seed;
144  for (i = 1; i < size; i++)
145  array[i] ^= (1812433253) * (array[i - 1] ^ (array[i - 1] >> 30)) + i;
146 }
147 
148 class mtgp32_engine
149 {
150 public:
151  __forceinline__ __device__ __host__
152  // Initialization is not supported for __shared__ variables
153  mtgp32_engine() // cppcheck-suppress uninitMemberVar
154  {
155 
156  }
157 
158  __forceinline__ __device__ __host__ mtgp32_engine(const mtgp32_state& m_state,
159  const mtgp32_params* params,
160  int bid)
161  {
162  this->m_state = m_state;
163  pos_tbl = params->pos_tbl[bid];
164  sh1_tbl = params->sh1_tbl[bid];
165  sh2_tbl = params->sh2_tbl[bid];
166  mask = params->mask[0];
167  for (int j = 0; j < MTGP_TS; j++) {
168  param_tbl[j] = params->param_tbl[bid][j];
169  temper_tbl[j] = params->temper_tbl[bid][j];
170  single_temper_tbl[j] = params->single_temper_tbl[bid][j];
171  }
172  }
173 
174  __forceinline__ __device__ __host__ void copy(const mtgp32_engine* m_engine)
175  {
176 #if defined(__HIP_DEVICE_COMPILE__)
177  const unsigned int thread_id = threadIdx.x;
178  for(int i = thread_id; i < MTGP_STATE; i += blockDim.x)
179  m_state.status[i] = m_engine->m_state.status[i];
180 
181  if (thread_id == 0)
182  {
183  m_state.offset = m_engine->m_state.offset;
184  m_state.id = m_engine->m_state.id;
185  pos_tbl = m_engine->pos_tbl;
186  sh1_tbl = m_engine->sh1_tbl;
187  sh2_tbl = m_engine->sh2_tbl;
188  mask = m_engine->mask;
189  }
190  if (thread_id < MTGP_TS)
191  {
192  param_tbl[thread_id] = m_engine->param_tbl[thread_id];
193  temper_tbl[thread_id] = m_engine->temper_tbl[thread_id];
194  single_temper_tbl[thread_id] = m_engine->single_temper_tbl[thread_id];
195  }
196  __syncthreads();
197 #else
198  this->m_state = m_engine->m_state;
199  pos_tbl = m_engine->pos_tbl;
200  sh1_tbl = m_engine->sh1_tbl;
201  sh2_tbl = m_engine->sh2_tbl;
202  mask = m_engine->mask;
203  for (int j = 0; j < MTGP_TS; j++) {
204  param_tbl[j] = m_engine->param_tbl[j];
205  temper_tbl[j] = m_engine->temper_tbl[j];
206  single_temper_tbl[j] = m_engine->single_temper_tbl[j];
207  }
208 #endif
209  }
210 
211  __forceinline__ __device__ __host__ void set_params(mtgp32_params* params)
212  {
213  pos_tbl = params->pos_tbl[m_state.id];
214  sh1_tbl = params->sh1_tbl[m_state.id];
215  sh2_tbl = params->sh2_tbl[m_state.id];
216  mask = params->mask[0];
217  for (int j = 0; j < MTGP_TS; j++) {
218  param_tbl[j] = params->param_tbl[m_state.id][j];
219  temper_tbl[j] = params->temper_tbl[m_state.id][j];
220  single_temper_tbl[j] = params->single_temper_tbl[m_state.id][j];
221  }
222  }
223 
224  __forceinline__ __device__ __host__ unsigned int operator()()
225  {
226  return this->next();
227  }
228 
229  __forceinline__ __device__ __host__ unsigned int next()
230  {
231 #ifdef __HIP_DEVICE_COMPILE__
232  unsigned int o = next_thread(threadIdx.x);
233  __syncthreads();
234  if(threadIdx.x == 0)
235  {
236  m_state.offset = (m_state.offset + blockDim.x) & MTGP_MASK;
237  }
238  __syncthreads();
239  return o;
240 #else
241  return 0;
242 #endif
243  }
244 
245  __forceinline__ __device__ __host__ unsigned int next_single()
246  {
247 #if defined(__HIP_DEVICE_COMPILE__)
248  unsigned int t = threadIdx.x;
249  unsigned int d = blockDim.x;
250  int pos = pos_tbl;
251  unsigned int r;
252  unsigned int o;
253 
254  r = para_rec(m_state.status[(t + m_state.offset) & MTGP_MASK],
255  m_state.status[(t + m_state.offset + 1) & MTGP_MASK],
256  m_state.status[(t + m_state.offset + pos) & MTGP_MASK]);
257  m_state.status[(t + m_state.offset + MTGP_N) & MTGP_MASK] = r;
258 
259  o = temper_single(r, m_state.status[(t + m_state.offset + pos - 1) & MTGP_MASK]);
260  __syncthreads();
261  if (t == 0)
262  m_state.offset = (m_state.offset + d) & MTGP_MASK;
263  __syncthreads();
264  return o;
265 #else
266  return 0;
267 #endif
268  }
269 
270 private:
271  __forceinline__ __device__ __host__ unsigned int
272  para_rec(unsigned int X1, unsigned int X2, unsigned int Y) const
273  {
274  unsigned int X = (X1 & mask) ^ X2;
275  unsigned int MAT;
276 
277  X ^= X << sh1_tbl;
278  Y = X ^ (Y >> sh2_tbl);
279  MAT = param_tbl[Y & 0x0f];
280  return Y ^ MAT;
281  }
282 
283  __forceinline__ __device__ __host__ unsigned int temper(unsigned int V, unsigned int T) const
284  {
285  unsigned int MAT;
286 
287  T ^= T >> 16;
288  T ^= T >> 8;
289  MAT = temper_tbl[T & 0x0f];
290  return V ^ MAT;
291  }
292 
293  __forceinline__ __device__ __host__ unsigned int temper_single(unsigned int V,
294  unsigned int T) const
295  {
296  unsigned int MAT;
297  unsigned int r;
298 
299  T ^= T >> 16;
300  T ^= T >> 8;
301  MAT = single_temper_tbl[T & 0x0f];
302  r = (V >> 9) ^ MAT;
303  return r;
304  }
305 
306 protected:
309  __forceinline__ __device__ __host__ unsigned int next_thread(unsigned int thread_idx)
310  {
311  const unsigned int r
312  = para_rec(m_state.status[(thread_idx + m_state.offset) & MTGP_MASK],
313  m_state.status[(thread_idx + m_state.offset + 1) & MTGP_MASK],
314  m_state.status[(thread_idx + m_state.offset + pos_tbl) & MTGP_MASK]);
315  m_state.status[(thread_idx + m_state.offset + MTGP_N) & MTGP_MASK] = r;
316  return temper(r, m_state.status[(thread_idx + m_state.offset + pos_tbl - 1) & MTGP_MASK]);
317  }
318 
319 public:
320  // State
321  mtgp32_state m_state;
322  // Parameters
323  unsigned int pos_tbl;
324  unsigned int param_tbl[MTGP_TS];
325  unsigned int temper_tbl[MTGP_TS];
326  unsigned int sh1_tbl;
327  unsigned int sh2_tbl;
328  unsigned int single_temper_tbl[MTGP_TS];
329  unsigned int mask;
330 
331 }; // mtgp32_engine class
332 
333 } // end namespace rocrand_device
334 
341 typedef rocrand_device::mtgp32_engine rocrand_state_mtgp32;
342 typedef rocrand_device::mtgp32_state mtgp32_state;
343 typedef rocrand_device::mtgp32_fast_params mtgp32_fast_params;
344 typedef rocrand_device::mtgp32_params mtgp32_params;
346 
362 __host__ inline rocrand_status rocrand_make_state_mtgp32(rocrand_state_mtgp32* state,
363  mtgp32_fast_params params[],
364  int n,
365  unsigned long long seed)
366 {
367  int i;
368  rocrand_state_mtgp32 * h_state = (rocrand_state_mtgp32 *) malloc(sizeof(rocrand_state_mtgp32) * n);
369  seed = seed ^ (seed >> 32);
370 
371  if (h_state == NULL)
373 
374  for (i = 0; i < n; i++) {
375  rocrand_device::rocrand_mtgp32_init_state(&(h_state[i].m_state.status[0]), &params[i], (unsigned int)seed + i + 1);
376  h_state[i].m_state.offset = 0;
377  h_state[i].m_state.id = i;
378  h_state[i].pos_tbl = params[i].pos;
379  h_state[i].sh1_tbl = params[i].sh1;
380  h_state[i].sh2_tbl = params[i].sh2;
381  h_state[i].mask = params[0].mask;
382  for (int j = 0; j < MTGP_TS; j++) {
383  h_state[i].param_tbl[j] = params[i].tbl[j];
384  h_state[i].temper_tbl[j] = params[i].tmp_tbl[j];
385  h_state[i].single_temper_tbl[j] = params[i].flt_tmp_tbl[j];
386  }
387  }
388 
389  const hipError_t error
390  = hipMemcpy(state, h_state, sizeof(rocrand_state_mtgp32) * n, hipMemcpyDefault);
391  free(h_state);
392 
393  if(error != hipSuccess)
395 
396  return ROCRAND_STATUS_SUCCESS;
397 }
398 
415 __host__ inline
416 rocrand_status rocrand_make_constant(const mtgp32_fast_params params[], mtgp32_params * p)
417 {
418  const int block_num = MTGP_BN_MAX;
419  const int size1 = sizeof(uint32_t) * block_num;
420  const int size2 = sizeof(uint32_t) * block_num * MTGP_TS;
421  uint32_t *h_pos_tbl;
422  uint32_t *h_sh1_tbl;
423  uint32_t *h_sh2_tbl;
424  uint32_t *h_param_tbl;
425  uint32_t *h_temper_tbl;
426  uint32_t *h_single_temper_tbl;
427  uint32_t *h_mask;
428  h_pos_tbl = (uint32_t *)malloc(size1);
429  h_sh1_tbl = (uint32_t *)malloc(size1);
430  h_sh2_tbl = (uint32_t *)malloc(size1);
431  h_param_tbl = (uint32_t *)malloc(size2);
432  h_temper_tbl = (uint32_t *)malloc(size2);
433  h_single_temper_tbl = (uint32_t *)malloc(size2);
434  h_mask = (uint32_t *)malloc(sizeof(uint32_t));
436 
437  if (h_pos_tbl == NULL || h_sh1_tbl == NULL || h_sh2_tbl == NULL
438  || h_param_tbl == NULL || h_temper_tbl == NULL || h_single_temper_tbl == NULL
439  || h_mask == NULL) {
440  printf("failure in allocating host memory for constant table.\n");
442  }
443  else {
444  h_mask[0] = params[0].mask;
445  for (int i = 0; i < block_num; i++) {
446  h_pos_tbl[i] = params[i].pos;
447  h_sh1_tbl[i] = params[i].sh1;
448  h_sh2_tbl[i] = params[i].sh2;
449  for (int j = 0; j < MTGP_TS; j++) {
450  h_param_tbl[i * MTGP_TS + j] = params[i].tbl[j];
451  h_temper_tbl[i * MTGP_TS + j] = params[i].tmp_tbl[j];
452  h_single_temper_tbl[i * MTGP_TS + j] = params[i].flt_tmp_tbl[j];
453  }
454  }
455 
456  if (hipMemcpy(p->pos_tbl, h_pos_tbl, size1, hipMemcpyHostToDevice) != hipSuccess)
458  if (hipMemcpy(p->sh1_tbl, h_sh1_tbl, size1, hipMemcpyHostToDevice) != hipSuccess)
460  if (hipMemcpy(p->sh2_tbl, h_sh2_tbl, size1, hipMemcpyHostToDevice) != hipSuccess)
462  if (hipMemcpy(p->param_tbl, h_param_tbl, size2, hipMemcpyHostToDevice) != hipSuccess)
464  if (hipMemcpy(p->temper_tbl, h_temper_tbl, size2, hipMemcpyHostToDevice) != hipSuccess)
466  if (hipMemcpy(p->single_temper_tbl, h_single_temper_tbl, size2, hipMemcpyHostToDevice) != hipSuccess)
468  if (hipMemcpy(p->mask, h_mask, sizeof(unsigned int), hipMemcpyHostToDevice) != hipSuccess)
470  }
471 
472  free(h_pos_tbl);
473  free(h_sh1_tbl);
474  free(h_sh2_tbl);
475  free(h_param_tbl);
476  free(h_temper_tbl);
477  free(h_single_temper_tbl);
478  free(h_mask);
479 
480  return status;
481 }
482 
495 __forceinline__ __device__ unsigned int rocrand(rocrand_state_mtgp32* state)
496 {
497  return state->next();
498 }
499 
531 __forceinline__ __device__ void rocrand_mtgp32_block_copy(rocrand_state_mtgp32* src,
532  rocrand_state_mtgp32* dest)
533 {
534  dest->copy(src);
535 }
536 
543 __forceinline__ __device__ void rocrand_mtgp32_set_params(rocrand_state_mtgp32* state,
544  mtgp32_params* params)
545 {
546  state->set_params(params);
547 }
548  // end of group rocranddevice
550 
551 #endif // ROCRAND_MTGP32_H_
__forceinline__ __device__ void rocrand_mtgp32_block_copy(rocrand_state_mtgp32 *src, rocrand_state_mtgp32 *dest)
Copies MTGP32 state to another state using block of threads.
Definition: rocrand_mtgp32.h:531
__host__ rocrand_status rocrand_make_state_mtgp32(rocrand_state_mtgp32 *state, mtgp32_fast_params params[], int n, unsigned long long seed)
Initializes MTGP32 states.
Definition: rocrand_mtgp32.h:362
__forceinline__ __device__ unsigned int rocrand(rocrand_state_mtgp32 *state)
Returns uniformly distributed random unsigned int value from [0; 2^32 - 1] range.
Definition: rocrand_mtgp32.h:495
__host__ rocrand_status rocrand_make_constant(const mtgp32_fast_params params[], mtgp32_params *p)
Loads parameters for MTGP32.
Definition: rocrand_mtgp32.h:416
__forceinline__ __device__ void rocrand_mtgp32_set_params(rocrand_state_mtgp32 *state, mtgp32_params *params)
Changes parameters of a MTGP32 state.
Definition: rocrand_mtgp32.h:543
rocrand_status
rocRAND function call status type
Definition: rocrand.h:59
@ ROCRAND_STATUS_SUCCESS
No errors.
Definition: rocrand.h:60
@ ROCRAND_STATUS_ALLOCATION_FAILED
Memory allocation failed during execution.
Definition: rocrand.h:63
Definition: rocrand_mtgp32.h:95
uint32_t tmp_tbl[16]
Definition: rocrand_mtgp32.h:101
int pos
Definition: rocrand_mtgp32.h:97
int mexp
Definition: rocrand_mtgp32.h:96
int sh2
Definition: rocrand_mtgp32.h:99
int sh1
Definition: rocrand_mtgp32.h:98
uint32_t mask
Definition: rocrand_mtgp32.h:103
unsigned char poly_sha1[21]
Definition: rocrand_mtgp32.h:104
uint32_t tbl[16]
Definition: rocrand_mtgp32.h:100
uint32_t flt_tmp_tbl[16]
Definition: rocrand_mtgp32.h:102