/build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_mtgp32.h Source File

/build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_mtgp32.h Source File#

API library: /build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_mtgp32.h Source File
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
107namespace rocrand_device {
108
109struct 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
120typedef mtgp32_params_fast_t mtgp32_fast_params;
121
122struct mtgp32_state
123{
124 int offset;
125 int id;
126 unsigned int status[MTGP_STATE];
127};
128
129inline
130void 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
148class mtgp32_engine
149{
150public:
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
270private:
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
306protected:
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
319public:
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
341typedef rocrand_device::mtgp32_engine rocrand_state_mtgp32;
342typedef rocrand_device::mtgp32_state mtgp32_state;
343typedef rocrand_device::mtgp32_fast_params mtgp32_fast_params;
344typedef 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
397}
398
415__host__ inline
416rocrand_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