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

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

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/docs-6.4.2/library/include/rocrand/rocrand_threefry2_impl.h Source File
rocrand_threefry2_impl.h
1 // Copyright (c) 2022-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 2010-2011, D. E. Shaw Research.
23 All rights reserved.
24 
25 Redistribution and use in source and binary forms, with or without
26 modification, are permitted provided that the following conditions are
27 met:
28 
29 * Redistributions of source code must retain the above copyright
30  notice, this list of conditions, and the following disclaimer.
31 
32 * Redistributions in binary form must reproduce the above copyright
33  notice, this list of conditions, and the following disclaimer in the
34  documentation and/or other materials provided with the distribution.
35 
36 * Neither the name of D. E. Shaw Research nor the names of its
37  contributors may be used to endorse or promote products derived from
38  this software without specific prior written permission.
39 
40 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
41 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
42 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
43 A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
44 OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
45 SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
46 LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
47 DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
48 THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
49 (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
50 OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
51 */
52 
53 #ifndef ROCRAND_THREEFRY2_IMPL_H_
54 #define ROCRAND_THREEFRY2_IMPL_H_
55 
56 #include "rocrand/rocrand_threefry_common.h"
57 #include <rocrand/rocrand_common.h>
58 
59 #ifndef THREEFRY2x32_DEFAULT_ROUNDS
60  #define THREEFRY2x32_DEFAULT_ROUNDS 20
61 #endif
62 
63 #ifndef THREEFRY2x64_DEFAULT_ROUNDS
64  #define THREEFRY2x64_DEFAULT_ROUNDS 20
65 #endif
66 
67 namespace rocrand_device
68 {
69 
70 template<class value>
71 __forceinline__ __device__ __host__ int threefry_rotation_array(int index) = delete;
72 
73 template<>
74 __forceinline__ __device__ __host__ int threefry_rotation_array<unsigned int>(int index)
75 {
76  // Output from skein_rot_search (srs32x2-X5000.out)
77  // Random seed = 1. BlockSize = 64 bits. sampleCnt = 1024. rounds = 8, minHW_or=28
78  // Start: Tue Jul 12 11:11:33 2011
79  // rMin = 0.334. #0206[*07] [CRC=1D9765C0. hw_OR=32. cnt=16384. blkSize= 64].format
80  static constexpr int THREEFRY_ROTATION_32_2[8] = {13, 15, 26, 6, 17, 29, 16, 24};
81  return THREEFRY_ROTATION_32_2[index];
82 }
83 
84 template<>
85 __forceinline__ __device__ __host__ int threefry_rotation_array<unsigned long long>(int index)
86 {
87  // Output from skein_rot_search: (srs64_B64-X1000)
88  // Random seed = 1. BlockSize = 128 bits. sampleCnt = 1024. rounds = 8, minHW_or=57
89  // Start: Tue Mar 1 10:07:48 2011
90  // rMin = 0.136. #0325[*15] [CRC=455A682F. hw_OR=64. cnt=16384. blkSize= 128].format
91  static constexpr int THREEFRY_ROTATION_64_2[8] = {16, 42, 12, 31, 16, 32, 24, 21};
92  return THREEFRY_ROTATION_64_2[index];
93 }
94 
95 template<typename state_value, typename value, unsigned int Nrounds>
96 class threefry_engine2_base
97 {
98 public:
99  struct threefry_state_2
100  {
101  state_value counter;
102  state_value key;
103  state_value result;
104  unsigned int substate;
105  };
106  using state_type = threefry_state_2;
107  using state_vector_type = state_value;
108 
109  __forceinline__ __device__ __host__ void discard(unsigned long long offset)
110  {
111  this->discard_impl(offset);
112  m_state.result = this->threefry_rounds(m_state.counter, m_state.key);
113  }
114 
115  __forceinline__ __device__ __host__ void discard()
116  {
117  m_state.result = this->threefry_rounds(m_state.counter, m_state.key);
118  }
119 
125  __forceinline__ __device__ __host__ void discard_subsequence(unsigned long long subsequence)
126  {
127  this->discard_subsequence_impl(subsequence);
128  m_state.result = this->threefry_rounds(m_state.counter, m_state.key);
129  }
130 
131  __forceinline__ __device__ __host__ value operator()()
132  {
133  return this->next();
134  }
135 
136  __forceinline__ __device__ __host__ value next()
137  {
138 #if defined(__HIP_PLATFORM_AMD__)
139  value ret = m_state.result.data[m_state.substate];
140 #else
141  value ret = (&m_state.result.x)[m_state.substate];
142 #endif
143  m_state.substate++;
144  if(m_state.substate == 2)
145  {
146  m_state.substate = 0;
147  m_state.counter = this->bump_counter(m_state.counter);
148  m_state.result = this->threefry_rounds(m_state.counter, m_state.key);
149  }
150  return ret;
151  }
152 
153  __forceinline__ __device__ __host__ state_value next2()
154  {
155  state_value ret = m_state.result;
156  m_state.counter = this->bump_counter(m_state.counter);
157  m_state.result = this->threefry_rounds(m_state.counter, m_state.key);
158 
159  return this->interleave(ret, m_state.result);
160  }
161 
162 protected:
163  __forceinline__ __device__ __host__ static state_value threefry_rounds(state_value counter,
164  state_value key)
165  {
166  state_value X;
167  value ks[2 + 1];
168 
169  static_assert(Nrounds <= 32, "32 or less only supported in threefry rounds");
170 
171  ks[2] = skein_ks_parity<value>();
172 
173  ks[0] = key.x;
174  ks[1] = key.y;
175 
176  X.x = counter.x;
177  X.y = counter.y;
178 
179  ks[2] ^= key.x;
180  ks[2] ^= key.y;
181 
182  /* Insert initial key before round 0 */
183  X.x += ks[0];
184  X.y += ks[1];
185 
186  for(unsigned int round_idx = 0; round_idx < Nrounds; round_idx++)
187  {
188  X.x += X.y;
189  X.y = rotl<value>(X.y, threefry_rotation_array<value>(round_idx & 7u));
190  X.y ^= X.x;
191 
192  if((round_idx & 3u) == 3)
193  {
194  unsigned int inject_idx = round_idx / 4;
195  // InjectKey(r = 1 + inject_idx)
196  X.x += ks[(1 + inject_idx) % 3];
197  X.y += ks[(2 + inject_idx) % 3];
198  X.y += 1 + inject_idx;
199  }
200  }
201 
202  return X;
203  }
204 
207  __forceinline__ __device__ __host__ void discard_impl(unsigned long long offset)
208  {
209  // Adjust offset for subset
210  m_state.substate += offset & 1;
211  unsigned long long counter_offset = offset / 2;
212  counter_offset += m_state.substate < 2 ? 0 : 1;
213  m_state.substate += m_state.substate < 2 ? 0 : -2;
214  // Discard states
215  this->discard_state(counter_offset);
216  }
217 
219  __forceinline__ __device__ __host__ void
220  discard_subsequence_impl(unsigned long long subsequence)
221  {
222  m_state.counter.y += subsequence;
223  }
224 
227  __forceinline__ __device__ __host__ void discard_state(unsigned long long offset)
228  {
229  value lo, hi;
230  ::rocrand_device::detail::split_ull(lo, hi, offset);
231 
232  value old_counter = m_state.counter.x;
233  m_state.counter.x += lo;
234  m_state.counter.y += hi + (m_state.counter.x < old_counter ? 1 : 0);
235  }
236 
237  __forceinline__ __device__ __host__ static state_value bump_counter(state_value counter)
238  {
239  counter.x++;
240  value add = counter.x == 0 ? 1 : 0;
241  counter.y += add;
242  return counter;
243  }
244 
245  __forceinline__ __device__ __host__ state_value interleave(const state_value prev,
246  const state_value next) const
247  {
248  switch(m_state.substate)
249  {
250  case 0: return prev;
251  case 1: return state_value{prev.y, next.x};
252  }
253  __builtin_unreachable();
254  }
255 
256 protected:
257  threefry_state_2 m_state;
258 }; // threefry_engine2_base class
259 
260 } // end namespace rocrand_device
261 
262 #endif // ROCRAND_THREEFRY2_IMPL_H_