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

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

API library: /build/reproducible-path/rocrand-6.4.1/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/*
22Copyright 2010-2011, D. E. Shaw Research.
23All rights reserved.
24
25Redistribution and use in source and binary forms, with or without
26modification, are permitted provided that the following conditions are
27met:
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
40THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
41"AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
42LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
43A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
44OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
45SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
46LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
47DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
48THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
49(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
50OF 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
67namespace rocrand_device
68{
69
70template<class value>
71__forceinline__ __device__ __host__ int threefry_rotation_array(int index) = delete;
72
73template<>
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
84template<>
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
95template<typename state_value, typename value, unsigned int Nrounds>
96class threefry_engine2_base
97{
98public:
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
162protected:
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
256protected:
257 threefry_state_2 m_state;
258}; // threefry_engine2_base class
259
260} // end namespace rocrand_device
261
262#endif // ROCRAND_THREEFRY2_IMPL_H_