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

API library

• library
• include
• rocrand

rocrand_mtgp32.h

1 // Copyright (c) 2017-2023 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 #ifndef FQUALIFIERS

61 #define FQUALIFIERS __forceinline__ __device__

62 #endif // FQUALIFIERS_

63 

64 #include "rocrand/rocrand.h"

65 #include "rocrand/rocrand_common.h"

66 

67 #define MTGP_MEXP 11213

68 #define MTGP_N 351

69 #define MTGP_FLOOR_2P 256

70 #define MTGP_CEIL_2P 512

71 #define MTGP_TN MTGP_FLOOR_2P

72 #define MTGP_LS (MTGP_TN * 3)

73 #define MTGP_BN_MAX 512

74 #define MTGP_TS 16

75 #define MTGP_STATE 1024

76 #define MTGP_MASK 1023

77 

78 // Source: https://github.com/MersenneTwister-Lab/MTGP/blob/master/mtgp32-fast.h

99 struct mtgp32_params_fast_t {

100  int mexp;

101  int pos;

102  int sh1;

103  int sh2;

104  uint32_t tbl[16];

105  uint32_t tmp_tbl[16];

106  uint32_t flt_tmp_tbl[16];

107  uint32_t mask;

108  unsigned char poly_sha1[21];

109 };

110 

111 namespace rocrand_device {

112 

113 struct mtgp32_params

114 {

115  unsigned int pos_tbl[MTGP_BN_MAX];

116  unsigned int param_tbl[MTGP_BN_MAX][MTGP_TS];

117  unsigned int temper_tbl[MTGP_BN_MAX][MTGP_TS];

118  unsigned int single_temper_tbl[MTGP_BN_MAX][MTGP_TS];

119  unsigned int sh1_tbl[MTGP_BN_MAX];

120  unsigned int sh2_tbl[MTGP_BN_MAX];

121  unsigned int mask[1];

122 };

123 

124 typedef mtgp32_params_fast_t mtgp32_fast_params;

125 

126 struct mtgp32_state

127 {

128  int offset;

129  int id;

130  unsigned int status[MTGP_STATE];

131 };

132 

133 inline

134 void rocrand_mtgp32_init_state(unsigned int array[],

135  const mtgp32_fast_params *para, unsigned int seed)

136 {

137  int i;

138  int size = para->mexp / 32 + 1;

139  unsigned int hidden_seed;

140  unsigned int tmp;

141  hidden_seed = para->tbl[4] ^ (para->tbl[8] << 16);

142  tmp = hidden_seed;

143  tmp += tmp >> 16;

144  tmp += tmp >> 8;

145  memset(array, tmp & 0xff, sizeof(unsigned int) * size);

146  array[0] = seed;

147  array[1] = hidden_seed;

148  for (i = 1; i < size; i++)

149  array[i] ^= (1812433253) * (array[i - 1] ^ (array[i - 1] >> 30)) + i;

150 }

151 

152 class mtgp32_engine

153 {

154 public:

155  FQUALIFIERS

156  // Initialization is not supported for __shared__ variables

157  mtgp32_engine() // cppcheck-suppress uninitMemberVar

158  {

159 

160  }

161 

162  FQUALIFIERS

163  mtgp32_engine(const mtgp32_state &m_state,

164  const mtgp32_params * params,

165  int bid)

166  {

167  this->m_state = m_state;

168  pos_tbl = params->pos_tbl[bid];

169  sh1_tbl = params->sh1_tbl[bid];

170  sh2_tbl = params->sh2_tbl[bid];

171  mask = params->mask[0];

172  for (int j = 0; j < MTGP_TS; j++) {

173  param_tbl[j] = params->param_tbl[bid][j];

174  temper_tbl[j] = params->temper_tbl[bid][j];

175  single_temper_tbl[j] = params->single_temper_tbl[bid][j];

176  }

177  }

178 

179  FQUALIFIERS

180  void copy(const mtgp32_engine * m_engine)

181  {

182 #if defined(__HIP_DEVICE_COMPILE__) || defined(USE_HIP_CPU)

183  const unsigned int thread_id = threadIdx.x;

184  for(int i = thread_id; i < MTGP_STATE; i += blockDim.x)

185  m_state.status[i] = m_engine->m_state.status[i];

186 

187  if (thread_id == 0)

188  {

189  m_state.offset = m_engine->m_state.offset;

190  m_state.id = m_engine->m_state.id;

191  pos_tbl = m_engine->pos_tbl;

192  sh1_tbl = m_engine->sh1_tbl;

193  sh2_tbl = m_engine->sh2_tbl;

194  mask = m_engine->mask;

195  }

196  if (thread_id < MTGP_TS)

197  {

198  param_tbl[thread_id] = m_engine->param_tbl[thread_id];

199  temper_tbl[thread_id] = m_engine->temper_tbl[thread_id];

200  single_temper_tbl[thread_id] = m_engine->single_temper_tbl[thread_id];

201  }

202  __syncthreads();

203 #else

204  this->m_state = m_engine->m_state;

205  pos_tbl = m_engine->pos_tbl;

206  sh1_tbl = m_engine->sh1_tbl;

207  sh2_tbl = m_engine->sh2_tbl;

208  mask = m_engine->mask;

209  for (int j = 0; j < MTGP_TS; j++) {

210  param_tbl[j] = m_engine->param_tbl[j];

211  temper_tbl[j] = m_engine->temper_tbl[j];

212  single_temper_tbl[j] = m_engine->single_temper_tbl[j];

213  }

214 #endif

215  }

216 

217  FQUALIFIERS

218  void set_params(mtgp32_params * params)

219  {

220  pos_tbl = params->pos_tbl[m_state.id];

221  sh1_tbl = params->sh1_tbl[m_state.id];

222  sh2_tbl = params->sh2_tbl[m_state.id];

223  mask = params->mask[0];

224  for (int j = 0; j < MTGP_TS; j++) {

225  param_tbl[j] = params->param_tbl[m_state.id][j];

226  temper_tbl[j] = params->temper_tbl[m_state.id][j];

227  single_temper_tbl[j] = params->single_temper_tbl[m_state.id][j];

228  }

229  }

230 

231  FQUALIFIERS

232  unsigned int operator()()

233  {

234  return this->next();

235  }

236 

237  FQUALIFIERS

238  unsigned int next()

239  {

240 #if defined(__HIP_DEVICE_COMPILE__) || defined(USE_HIP_CPU)

241  unsigned int t = threadIdx.x;

242  unsigned int d = blockDim.x;

243  int pos = pos_tbl;

244  unsigned int r;

245  unsigned int o;

246 

247  r = para_rec(m_state.status[(t + m_state.offset) & MTGP_MASK],

248  m_state.status[(t + m_state.offset + 1) & MTGP_MASK],

249  m_state.status[(t + m_state.offset + pos) & MTGP_MASK]);

250  m_state.status[(t + m_state.offset + MTGP_N) & MTGP_MASK] = r;

251 

252  o = temper(r, m_state.status[(t + m_state.offset + pos - 1) & MTGP_MASK]);

253  __syncthreads();

254  if (t == 0)

255  m_state.offset = (m_state.offset + d) & MTGP_MASK;

256  __syncthreads();

257  return o;

258 #else

259  return 0;

260 #endif

261  }

262 

263  FQUALIFIERS

264  unsigned int next_single()

265  {

266 #if defined(__HIP_DEVICE_COMPILE__) || defined(USE_HIP_CPU)

267  unsigned int t = threadIdx.x;

268  unsigned int d = blockDim.x;

269  int pos = pos_tbl;

270  unsigned int r;

271  unsigned int o;

272 

273  r = para_rec(m_state.status[(t + m_state.offset) & MTGP_MASK],

274  m_state.status[(t + m_state.offset + 1) & MTGP_MASK],

275  m_state.status[(t + m_state.offset + pos) & MTGP_MASK]);

276  m_state.status[(t + m_state.offset + MTGP_N) & MTGP_MASK] = r;

277 

278  o = temper_single(r, m_state.status[(t + m_state.offset + pos - 1) & MTGP_MASK]);

279  __syncthreads();

280  if (t == 0)

281  m_state.offset = (m_state.offset + d) & MTGP_MASK;

282  __syncthreads();

283  return o;

284 #else

285  return 0;

286 #endif

287  }

288 

289 private:

290  FQUALIFIERS

291  unsigned int para_rec(unsigned int X1, unsigned int X2, unsigned int Y) const

292  {

293  unsigned int X = (X1 & mask) ^ X2;

294  unsigned int MAT;

295 

296  X ^= X << sh1_tbl;

297  Y = X ^ (Y >> sh2_tbl);

298  MAT = param_tbl[Y & 0x0f];

299  return Y ^ MAT;

300  }

301 

302  FQUALIFIERS

303  unsigned int temper(unsigned int V, unsigned int T) const

304  {

305  unsigned int MAT;

306 

307  T ^= T >> 16;

308  T ^= T >> 8;

309  MAT = temper_tbl[T & 0x0f];

310  return V ^ MAT;

311  }

312 

313  FQUALIFIERS

314  unsigned int temper_single(unsigned int V, unsigned int T) const

315  {

316  unsigned int MAT;

317  unsigned int r;

318 

319  T ^= T >> 16;

320  T ^= T >> 8;

321  MAT = single_temper_tbl[T & 0x0f];

322  r = (V >> 9) ^ MAT;

323  return r;

324  }

325 

326 public:

327  // State

328  mtgp32_state m_state;

329  // Parameters

330  unsigned int pos_tbl;

331  unsigned int param_tbl[MTGP_TS];

332  unsigned int temper_tbl[MTGP_TS];

333  unsigned int sh1_tbl;

334  unsigned int sh2_tbl;

335  unsigned int single_temper_tbl[MTGP_TS];

336  unsigned int mask;

337 

338 }; // mtgp32_engine class

339 

340 } // end namespace rocrand_device

341 

348 typedef rocrand_device::mtgp32_engine rocrand_state_mtgp32;

349 typedef rocrand_device::mtgp32_state mtgp32_state;

350 typedef rocrand_device::mtgp32_fast_params mtgp32_fast_params;

351 typedef rocrand_device::mtgp32_params mtgp32_params;

353 

369 __host__ inline

370 rocrand_status rocrand_make_state_mtgp32(rocrand_state_mtgp32 * d_state,

371  mtgp32_fast_params params[],

372  int n,

373  unsigned long long seed)

374 {

375  int i;

376  rocrand_state_mtgp32 * h_state = (rocrand_state_mtgp32 *) malloc(sizeof(rocrand_state_mtgp32) * n);

377  seed = seed ^ (seed >> 32);

378 

379  if (h_state == NULL)

380  return ROCRAND_STATUS_ALLOCATION_FAILED;

381 

382  for (i = 0; i < n; i++) {

383  rocrand_device::rocrand_mtgp32_init_state(&(h_state[i].m_state.status[0]), &params[i], (unsigned int)seed + i + 1);

384  h_state[i].m_state.offset = 0;

385  h_state[i].m_state.id = i;

386  h_state[i].pos_tbl = params[i].pos;

387  h_state[i].sh1_tbl = params[i].sh1;

388  h_state[i].sh2_tbl = params[i].sh2;

389  h_state[i].mask = params[0].mask;

390  for (int j = 0; j < MTGP_TS; j++) {

391  h_state[i].param_tbl[j] = params[i].tbl[j];

392  h_state[i].temper_tbl[j] = params[i].tmp_tbl[j];

393  h_state[i].single_temper_tbl[j] = params[i].flt_tmp_tbl[j];

394  }

395  }

396 

397  hipMemcpy(d_state, h_state, sizeof(rocrand_state_mtgp32) * n, hipMemcpyHostToDevice);

398  free(h_state);

399 

400  if (hipGetLastError() != hipSuccess)

401  return ROCRAND_STATUS_ALLOCATION_FAILED;

402 

403  return ROCRAND_STATUS_SUCCESS;

404 }

405 

422 __host__ inline

423 rocrand_status rocrand_make_constant(const mtgp32_fast_params params[], mtgp32_params * p)

424 {

425  const int block_num = MTGP_BN_MAX;

426  const int size1 = sizeof(uint32_t) * block_num;

427  const int size2 = sizeof(uint32_t) * block_num * MTGP_TS;

428  uint32_t *h_pos_tbl;

429  uint32_t *h_sh1_tbl;

430  uint32_t *h_sh2_tbl;

431  uint32_t *h_param_tbl;

432  uint32_t *h_temper_tbl;

433  uint32_t *h_single_temper_tbl;

434  uint32_t *h_mask;

435  h_pos_tbl = (uint32_t *)malloc(size1);

436  h_sh1_tbl = (uint32_t *)malloc(size1);

437  h_sh2_tbl = (uint32_t *)malloc(size1);

438  h_param_tbl = (uint32_t *)malloc(size2);

439  h_temper_tbl = (uint32_t *)malloc(size2);

440  h_single_temper_tbl = (uint32_t *)malloc(size2);

441  h_mask = (uint32_t *)malloc(sizeof(uint32_t));

442  rocrand_status status = ROCRAND_STATUS_SUCCESS;

443 

444  if (h_pos_tbl == NULL || h_sh1_tbl == NULL || h_sh2_tbl == NULL

445  || h_param_tbl == NULL || h_temper_tbl == NULL || h_single_temper_tbl == NULL

446  || h_mask == NULL) {

447  printf("failure in allocating host memory for constant table.\n");

448  status = ROCRAND_STATUS_ALLOCATION_FAILED;

449  }

450  else {

451  h_mask[0] = params[0].mask;

452  for (int i = 0; i < block_num; i++) {

453  h_pos_tbl[i] = params[i].pos;

454  h_sh1_tbl[i] = params[i].sh1;

455  h_sh2_tbl[i] = params[i].sh2;

456  for (int j = 0; j < MTGP_TS; j++) {

457  h_param_tbl[i * MTGP_TS + j] = params[i].tbl[j];

458  h_temper_tbl[i * MTGP_TS + j] = params[i].tmp_tbl[j];

459  h_single_temper_tbl[i * MTGP_TS + j] = params[i].flt_tmp_tbl[j];

460  }

461  }

462 

463  if (hipMemcpy(p->pos_tbl, h_pos_tbl, size1, hipMemcpyHostToDevice) != hipSuccess)

464  status = ROCRAND_STATUS_ALLOCATION_FAILED;

465  if (hipMemcpy(p->sh1_tbl, h_sh1_tbl, size1, hipMemcpyHostToDevice) != hipSuccess)

466  status = ROCRAND_STATUS_ALLOCATION_FAILED;

467  if (hipMemcpy(p->sh2_tbl, h_sh2_tbl, size1, hipMemcpyHostToDevice) != hipSuccess)

468  status = ROCRAND_STATUS_ALLOCATION_FAILED;

469  if (hipMemcpy(p->param_tbl, h_param_tbl, size2, hipMemcpyHostToDevice) != hipSuccess)

470  status = ROCRAND_STATUS_ALLOCATION_FAILED;

471  if (hipMemcpy(p->temper_tbl, h_temper_tbl, size2, hipMemcpyHostToDevice) != hipSuccess)

472  status = ROCRAND_STATUS_ALLOCATION_FAILED;

473  if (hipMemcpy(p->single_temper_tbl, h_single_temper_tbl, size2, hipMemcpyHostToDevice) != hipSuccess)

474  status = ROCRAND_STATUS_ALLOCATION_FAILED;

475  if (hipMemcpy(p->mask, h_mask, sizeof(unsigned int), hipMemcpyHostToDevice) != hipSuccess)

476  status = ROCRAND_STATUS_ALLOCATION_FAILED;

477  }

478 

479  free(h_pos_tbl);

480  free(h_sh1_tbl);

481  free(h_sh2_tbl);

482  free(h_param_tbl);

483  free(h_temper_tbl);

484  free(h_single_temper_tbl);

485  free(h_mask);

486 

487  return status;

488 }

489 

502 FQUALIFIERS

503 unsigned int rocrand(rocrand_state_mtgp32 * state)

504 {

505  return state->next();

506 }

507 

539 FQUALIFIERS

540 void rocrand_mtgp32_block_copy(rocrand_state_mtgp32 * src, rocrand_state_mtgp32 * dest)

541 {

542  dest->copy(src);

543 }

544 

551 FQUALIFIERS

552 void rocrand_mtgp32_set_params(rocrand_state_mtgp32 * state, mtgp32_params * params)

553 {

554  state->set_params(params);

555 }

556  // end of group rocranddevice

558 

559 #endif // ROCRAND_MTGP32_H_

rocrand_mtgp32_block_copy

FQUALIFIERS 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:540

rocrand_make_state_mtgp32

__host__ rocrand_status rocrand_make_state_mtgp32(rocrand_state_mtgp32 *d_state, mtgp32_fast_params params[], int n, unsigned long long seed)

Initializes MTGP32 states.

Definition: rocrand_mtgp32.h:370

rocrand_make_constant

__host__ rocrand_status rocrand_make_constant(const mtgp32_fast_params params[], mtgp32_params *p)

Loads parameters for MTGP32.

Definition: rocrand_mtgp32.h:423

rocrand

FQUALIFIERS unsigned int rocrand(rocrand_state_mtgp32 *state)

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

Definition: rocrand_mtgp32.h:503

rocrand_mtgp32_set_params

FQUALIFIERS void rocrand_mtgp32_set_params(rocrand_state_mtgp32 *state, mtgp32_params *params)

Changes parameters of a MTGP32 state.

Definition: rocrand_mtgp32.h:552

FQUALIFIERS

#define FQUALIFIERS

Shorthand for commonly used function qualifiers.

Definition: rocrand_uniform.h:31

rocrand_status

rocrand_status

rocRAND function call status type

Definition: rocrand.h:60

ROCRAND_STATUS_SUCCESS

@ ROCRAND_STATUS_SUCCESS

No errors.

Definition: rocrand.h:61

ROCRAND_STATUS_ALLOCATION_FAILED

@ ROCRAND_STATUS_ALLOCATION_FAILED

Memory allocation failed during execution.

Definition: rocrand.h:64

mtgp32_params_fast_t

Definition: rocrand_mtgp32.h:99

mtgp32_params_fast_t::tmp_tbl

uint32_t tmp_tbl[16]

Definition: rocrand_mtgp32.h:105

mtgp32_params_fast_t::pos

int pos

Definition: rocrand_mtgp32.h:101

mtgp32_params_fast_t::mexp

int mexp

Definition: rocrand_mtgp32.h:100

mtgp32_params_fast_t::sh2

int sh2

Definition: rocrand_mtgp32.h:103

mtgp32_params_fast_t::sh1

int sh1

Definition: rocrand_mtgp32.h:102

mtgp32_params_fast_t::mask

uint32_t mask

Definition: rocrand_mtgp32.h:107

mtgp32_params_fast_t::poly_sha1

unsigned char poly_sha1[21]

Definition: rocrand_mtgp32.h:108

mtgp32_params_fast_t::tbl

uint32_t tbl[16]

Definition: rocrand_mtgp32.h:104

mtgp32_params_fast_t::flt_tmp_tbl

uint32_t flt_tmp_tbl[16]

Definition: rocrand_mtgp32.h:106

---

Generated by  1.9.1