/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/develop/projects/rocrand/library/include/rocrand/rocrand_threefry_common.h Source File

• library
• include
• rocrand

rocrand_threefry_common.h

1 // Copyright (c) 2022-2026 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_THREEFRY_COMMON_H_

54 #define ROCRAND_THREEFRY_COMMON_H_

55 

56 #include <hip/hip_runtime.h>

57 

58 // C240 constant for Skein Hash function Threefish

59 #define SKEIN_MK_64(hi32, lo32) ((lo32) + (((unsigned long long)(hi32)) << 32))

60 #define SKEIN_KS_PARITY64 SKEIN_MK_64(0x1BD11BDA, 0xA9FC1A22)

61 #define SKEIN_KS_PARITY32 0x1BD11BDA

62 

63 namespace rocrand_device

64 {

65 

66 template<typename value>

67 __forceinline__ __device__ __host__

68 value rotl(value x, int d);

69 

70 template<>

71 __forceinline__ __device__ __host__

72 unsigned long long rotl<unsigned long long>(unsigned long long x, int d)

73 {

74  return (x << (d & 63)) | (x >> ((64 - d) & 63));

75 };

76 

77 template<>

78 __forceinline__ __device__ __host__

79 unsigned int rotl<unsigned int>(unsigned int x, int d)

80 {

81  return (x << (d & 31)) | (x >> ((32 - d) & 31));

82 };

83 

84 template<typename value>

85 __forceinline__ __device__ __host__

86 value skein_ks_parity();

87 

88 template<>

89 __forceinline__ __device__ __host__

90 unsigned int skein_ks_parity<unsigned int>()

91 {

92  return SKEIN_KS_PARITY32;

93 }

94 

95 template<>

96 __forceinline__ __device__ __host__

97 unsigned long long skein_ks_parity<unsigned long long>()

98 {

99  return SKEIN_KS_PARITY64;

100 }

101 

102 template<class value>

103 __forceinline__ __device__ __host__

104 int threefry_rotation_array(int index)

105  = delete;

106 

107 template<>

108 __forceinline__ __device__ __host__

109 int threefry_rotation_array<unsigned int>(int index)

110 {

111  // Output from skein_rot_search (srs32x2-X5000.out)

112  // Random seed = 1. BlockSize = 64 bits. sampleCnt = 1024. rounds = 8, minHW_or=28

113  // Start: Tue Jul 12 11:11:33 2011

114  // rMin = 0.334. #0206[*07] [CRC=1D9765C0. hw_OR=32. cnt=16384. blkSize= 64].format

115  static constexpr int THREEFRY_ROTATION_32_2[8] = {13, 15, 26, 6, 17, 29, 16, 24};

116  return THREEFRY_ROTATION_32_2[index];

117 }

118 

119 template<>

120 __forceinline__ __device__ __host__

121 int threefry_rotation_array<unsigned long long>(int index)

122 {

123  // Output from skein_rot_search: (srs64_B64-X1000)

124  // Random seed = 1. BlockSize = 128 bits. sampleCnt = 1024. rounds = 8, minHW_or=57

125  // Start: Tue Mar 1 10:07:48 2011

126  // rMin = 0.136. #0325[*15] [CRC=455A682F. hw_OR=64. cnt=16384. blkSize= 128].format

127  static constexpr int THREEFRY_ROTATION_64_2[8] = {16, 42, 12, 31, 16, 32, 24, 21};

128  return THREEFRY_ROTATION_64_2[index];

129 }

130 

131 template<typename state_value, typename value, unsigned int Nrounds>

132 __forceinline__ __device__ __host__

133 auto rounds_2(state_value X, value ks[3]) -> std::enable_if_t<Nrounds % 4 != 0, state_value>

134 {

135  for(unsigned int round_idx = 0; round_idx < Nrounds; round_idx++)

136  {

137  X.x += X.y;

138  X.y = rotl<value>(X.y, threefry_rotation_array<value>(round_idx & 7u));

139  X.y ^= X.x;

140 

141  if((round_idx & 3u) == 3)

142  {

143  unsigned int inject_idx = round_idx / 4;

144  // InjectKey(r = 1 + inject_idx)

145  X.x += ks[(1 + inject_idx) % 3];

146  X.y += ks[(2 + inject_idx) % 3];

147  X.y += 1 + inject_idx;

148  }

149  }

150  return X;

151 }

152 

153 template<typename state_value, typename value, unsigned int Nrounds>

154 __forceinline__ __device__ __host__

155 auto rounds_2(state_value X, value ks[3]) -> std::enable_if_t<Nrounds % 4 == 0, state_value>

156 {

157 #pragma unroll

158  for(unsigned int i = 0; i < Nrounds / 4; i++)

159  {

160  unsigned int round_idx = 4 * i;

161  for(unsigned int j = 0; j < 4; j++)

162  {

163  X.x += X.y;

164  X.y = rotl<value>(X.y, threefry_rotation_array<value>((round_idx + j) & 7u));

165  X.y ^= X.x;

166  }

167 

168  // (round_idx & 3u) == 3

169  // InjectKey(r = 1 + i)

170  X.x += ks[(1 + i) % 3];

171  X.y += ks[(2 + i) % 3];

172  X.y += 1 + i;

173  }

174  return X;

175 }

176 

177 template<class value>

178 __forceinline__ __device__ __host__

179 int threefry_rotation_array(int indexX, int indexY)

180  = delete;

181 

182 template<>

183 __forceinline__ __device__ __host__

184 int threefry_rotation_array<unsigned int>(int indexX, int indexY)

185 {

186  // Output from skein_rot_search: (srs-B128-X5000.out)

187  // Random seed = 1. BlockSize = 64 bits. sampleCnt = 1024. rounds = 8, minHW_or=28

188  // Start: Mon Aug 24 22:41:36 2009

189  // ...

190  // rMin = 0.472. #0A4B[*33] [CRC=DD1ECE0F. hw_OR=31. cnt=16384. blkSize= 128].format

191  static constexpr int THREEFRY_ROTATION_32_4[8][2] = {

192  {10, 26},

193  {11, 21},

194  {13, 27},

195  {23, 5},

196  { 6, 20},

197  {17, 11},

198  {25, 10},

199  {18, 20}

200  };

201  return THREEFRY_ROTATION_32_4[indexX][indexY];

202 }

203 

204 template<>

205 __forceinline__ __device__ __host__

206 int threefry_rotation_array<unsigned long long>(int indexX, int indexY)

207 {

208  // These are the R_256 constants from the Threefish reference sources

209  // with names changed to R_64x4... */

210  static constexpr int THREEFRY_ROTATION_64_4[8][2] = {

211  {14, 16},

212  {52, 57},

213  {23, 40},

214  { 5, 37},

215  {25, 33},

216  {46, 12},

217  {58, 22},

218  {32, 32}

219  };

220  return THREEFRY_ROTATION_64_4[indexX][indexY];

221 }

222 

223 template<typename state_value, typename value, unsigned int Nrounds>

224 __forceinline__ __device__ __host__

225 auto rounds_4(state_value X, value ks[5]) -> std::enable_if_t<Nrounds % 4 != 0, state_value>

226 {

227  for(unsigned int round_idx = 0; round_idx < Nrounds; round_idx++)

228  {

229  int rot_0 = threefry_rotation_array<value>(round_idx & 7u, 0);

230  int rot_1 = threefry_rotation_array<value>(round_idx & 7u, 1);

231  if((round_idx & 2u) == 0)

232  {

233  X.x += X.y;

234  X.y = rotl<value>(X.y, rot_0);

235  X.y ^= X.x;

236  X.z += X.w;

237  X.w = rotl<value>(X.w, rot_1);

238  X.w ^= X.z;

239  }

240  else

241  {

242  X.x += X.w;

243  X.w = rotl<value>(X.w, rot_0);

244  X.w ^= X.x;

245  X.z += X.y;

246  X.y = rotl<value>(X.y, rot_1);

247  X.y ^= X.z;

248  }

249 

250  if((round_idx & 3u) == 3)

251  {

252  unsigned int inject_idx = round_idx / 4;

253  // InjectKey(r = 1 + inject_idx)

254  X.x += ks[(1 + inject_idx) % 5];

255  X.y += ks[(2 + inject_idx) % 5];

256  X.z += ks[(3 + inject_idx) % 5];

257  X.w += ks[(4 + inject_idx) % 5];

258  X.w += 1 + inject_idx;

259  }

260  }

261  return X;

262 }

263 

264 template<typename state_value, typename value, unsigned int Nrounds>

265 __forceinline__ __device__ __host__

266 auto rounds_4(state_value X, value ks[5]) -> std::enable_if_t<Nrounds % 4 == 0, state_value>

267 {

268 #pragma unroll

269  for(unsigned int i = 0; i < Nrounds / 4; i++)

270  {

271  unsigned int round_idx = 4 * i;

272 

273  // (round_idx & 2u) == 0

274  unsigned int j = 0;

275  for(; j < 2; j++)

276  {

277  X.x += X.y;

278  X.y = rotl<value>(X.y, threefry_rotation_array<value>((round_idx + j) & 7u, 0));

279  X.y ^= X.x;

280  X.z += X.w;

281  X.w = rotl<value>(X.w, threefry_rotation_array<value>((round_idx + j) & 7u, 1));

282  X.w ^= X.z;

283  }

284 

285  // (round_idx & 2u) != 0

286  for(; j < 4; j++)

287  {

288  X.x += X.w;

289  X.w = rotl<value>(X.w, threefry_rotation_array<value>((round_idx + j) & 7u, 0));

290  X.w ^= X.x;

291  X.z += X.y;

292  X.y = rotl<value>(X.y, threefry_rotation_array<value>((round_idx + j) & 7u, 1));

293  X.y ^= X.z;

294  }

295 

296  // (round_idx & 3u) == 3

297  // InjectKey(r = 1 + i)

298  X.x += ks[(1 + i) % 5];

299  X.y += ks[(2 + i) % 5];

300  X.z += ks[(3 + i) % 5];

301  X.w += ks[(4 + i) % 5];

302  X.w += 1 + i;

303  }

304  return X;

305 }

306 

307 } // end namespace rocrand_device

308 

309 #endif // ROCRAND_THREEFRY_COMMON_H_