nv/m05000_a1.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _KECCAK_
   7
   8 #include "include/constants.h"
   9 #include "include/kernel_vendor.h"
  10
  11 #ifdef  VLIW1
  12 #define VECT_SIZE1
  13 #endif
  14
  15 #ifdef  VLIW2
  16 #define VECT_SIZE1
  17 #endif
  18
  19 #define DGST_R0 2
  20 #define DGST_R1 3
  21 #define DGST_R2 4
  22 #define DGST_R3 5
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27
  28 #ifdef  VECT_SIZE1
  29 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  30 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  31 #endif
  32
  33 #ifdef  VECT_SIZE2
  34 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  35 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  36 #endif
  37
  38 #ifndef KECCAK_ROUNDS
  39 #define KECCAK_ROUNDS 24
  40 #endif
  41
  42 #define Theta1(s) (st[0 + s] ^ st[5 + s] ^ st[10 + s] ^ st[15 + s] ^ st[20 + s])
  43
  44 #define Theta2(s)               \
  45 {                               \
  46   st[ 0 + s] ^= t;              \
  47   st[ 5 + s] ^= t;              \
  48   st[10 + s] ^= t;              \
  49   st[15 + s] ^= t;              \
  50   st[20 + s] ^= t;              \
  51 }
  52
  53 #define Rho_Pi(s)               \
  54 {                               \
  55   u32 j = keccakf_piln[s];     \
  56   u32 k = keccakf_rotc[s];     \
  57   bc0 = st[j];                  \
  58   st[j] = rotl64 (t, k);        \
  59   t = bc0;                      \
  60 }
  61
  62 #define Chi(s)                  \
  63 {                               \
  64   bc0 = st[0 + s];              \
  65   bc1 = st[1 + s];              \
  66   bc2 = st[2 + s];              \
  67   bc3 = st[3 + s];              \
  68   bc4 = st[4 + s];              \
  69   st[0 + s] ^= ~bc1 & bc2;      \
  70   st[1 + s] ^= ~bc2 & bc3;      \
  71   st[2 + s] ^= ~bc3 & bc4;      \
  72   st[3 + s] ^= ~bc4 & bc0;      \
  73   st[4 + s] ^= ~bc0 & bc1;      \
  74 }
  75
  76 __device__ __constant__ comb_t c_combs[1024];
  77
  78 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  79 {
  80   /**
  81    * modifier
  82    */
  83
  84   const u32 lid = threadIdx.x;
  85
  86   /**
  87    * base
  88    */
  89
  90   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  91
  92   if (gid >= gid_max) return;
  93
  94   /**
  95    * constants
  96    */
  97
  98   const u64 keccakf_rndc[24] =
  99   {
 100     0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
 101     0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
 102     0x8000000080008081, 0x8000000000008009, 0x000000000000008a,
 103     0x0000000000000088, 0x0000000080008009, 0x000000008000000a,
 104     0x000000008000808b, 0x800000000000008b, 0x8000000000008089,
 105     0x8000000000008003, 0x8000000000008002, 0x8000000000000080,
 106     0x000000000000800a, 0x800000008000000a, 0x8000000080008081,
 107     0x8000000000008080, 0x0000000080000001, 0x8000000080008008
 108   };
 109
 110   const u32 keccakf_rotc[24] =
 111   {
 112      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
 113     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
 114   };
 115
 116   const u32 keccakf_piln[24] =
 117   {
 118     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
 119     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
 120   };
 121
 122   u32x wordl0[4];
 123
 124   wordl0[0] = pws[gid].i[ 0];
 125   wordl0[1] = pws[gid].i[ 1];
 126   wordl0[2] = pws[gid].i[ 2];
 127   wordl0[3] = pws[gid].i[ 3];
 128
 129   u32x wordl1[4];
 130
 131   wordl1[0] = pws[gid].i[ 4];
 132   wordl1[1] = pws[gid].i[ 5];
 133   wordl1[2] = pws[gid].i[ 6];
 134   wordl1[3] = pws[gid].i[ 7];
 135
 136   u32x wordl2[4];
 137
 138   wordl2[0] = 0;
 139   wordl2[1] = 0;
 140   wordl2[2] = 0;
 141   wordl2[3] = 0;
 142
 143   u32x wordl3[4];
 144
 145   wordl3[0] = 0;
 146   wordl3[1] = 0;
 147   wordl3[2] = 0;
 148   wordl3[3] = 0;
 149
 150   const u32 pw_l_len = pws[gid].pw_len;
 151
 152   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 153   {
 154     append_0x01_2 (wordl0, wordl1, pw_l_len);
 155
 156     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 157   }
 158
 159   /**
 160    * 0x80 keccak, very special
 161    */
 162
 163   const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
 164
 165   const u32 rsiz = 200 - (2 * mdlen);
 166
 167   const u32 add80w = (rsiz - 1) / 8;
 168
 169   /**
 170    * loop
 171    */
 172
 173   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 174   {
 175     const u32 pw_r_len = c_combs[il_pos].pw_len;
 176
 177     const u32 pw_len = pw_l_len + pw_r_len;
 178
 179     u32 wordr0[4];
 180
 181     wordr0[0] = c_combs[il_pos].i[0];
 182     wordr0[1] = c_combs[il_pos].i[1];
 183     wordr0[2] = c_combs[il_pos].i[2];
 184     wordr0[3] = c_combs[il_pos].i[3];
 185
 186     u32 wordr1[4];
 187
 188     wordr1[0] = c_combs[il_pos].i[4];
 189     wordr1[1] = c_combs[il_pos].i[5];
 190     wordr1[2] = c_combs[il_pos].i[6];
 191     wordr1[3] = c_combs[il_pos].i[7];
 192
 193     u32 wordr2[4];
 194
 195     wordr2[0] = 0;
 196     wordr2[1] = 0;
 197     wordr2[2] = 0;
 198     wordr2[3] = 0;
 199
 200     u32 wordr3[4];
 201
 202     wordr3[0] = 0;
 203     wordr3[1] = 0;
 204     wordr3[2] = 0;
 205     wordr3[3] = 0;
 206
 207     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 208     {
 209       append_0x01_2 (wordr0, wordr1, pw_r_len);
 210
 211       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 212     }
 213
 214     u32x w0[4];
 215
 216     w0[0] = wordl0[0] | wordr0[0];
 217     w0[1] = wordl0[1] | wordr0[1];
 218     w0[2] = wordl0[2] | wordr0[2];
 219     w0[3] = wordl0[3] | wordr0[3];
 220
 221     u32x w1[4];
 222
 223     w1[0] = wordl1[0] | wordr1[0];
 224     w1[1] = wordl1[1] | wordr1[1];
 225     w1[2] = wordl1[2] | wordr1[2];
 226     w1[3] = wordl1[3] | wordr1[3];
 227
 228     u32x w2[4];
 229
 230     w2[0] = wordl2[0] | wordr2[0];
 231     w2[1] = wordl2[1] | wordr2[1];
 232     w2[2] = wordl2[2] | wordr2[2];
 233     w2[3] = wordl2[3] | wordr2[3];
 234
 235     u32x w3[4];
 236
 237     w3[0] = wordl3[0] | wordr3[0];
 238     w3[1] = wordl3[1] | wordr3[1];
 239     w3[2] = pw_len * 8;
 240     w3[3] = 0;
 241
 242     u64x st[25];
 243
 244     st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
 245     st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
 246     st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
 247     st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
 248     st[ 4] = 0;
 249     st[ 5] = 0;
 250     st[ 6] = 0;
 251     st[ 7] = 0;
 252     st[ 8] = 0;
 253     st[ 9] = 0;
 254     st[10] = 0;
 255     st[11] = 0;
 256     st[12] = 0;
 257     st[13] = 0;
 258     st[14] = 0;
 259     st[15] = 0;
 260     st[16] = 0;
 261     st[17] = 0;
 262     st[18] = 0;
 263     st[19] = 0;
 264     st[20] = 0;
 265     st[21] = 0;
 266     st[22] = 0;
 267     st[23] = 0;
 268     st[24] = 0;
 269
 270     st[add80w] |= 0x8000000000000000;
 271
 272     int round;
 273
 274     for (round = 0; round < KECCAK_ROUNDS; round++)
 275     {
 276       // Theta
 277
 278       u64x bc0 = Theta1 (0);
 279       u64x bc1 = Theta1 (1);
 280       u64x bc2 = Theta1 (2);
 281       u64x bc3 = Theta1 (3);
 282       u64x bc4 = Theta1 (4);
 283
 284       u64x t;
 285
 286       t = bc4 ^ rotl64 (bc1, 1); Theta2 (0);
 287       t = bc0 ^ rotl64 (bc2, 1); Theta2 (1);
 288       t = bc1 ^ rotl64 (bc3, 1); Theta2 (2);
 289       t = bc2 ^ rotl64 (bc4, 1); Theta2 (3);
 290       t = bc3 ^ rotl64 (bc0, 1); Theta2 (4);
 291
 292       // Rho Pi
 293
 294       t = st[1];
 295
 296       Rho_Pi (0);
 297       Rho_Pi (1);
 298       Rho_Pi (2);
 299       Rho_Pi (3);
 300       Rho_Pi (4);
 301       Rho_Pi (5);
 302       Rho_Pi (6);
 303       Rho_Pi (7);
 304       Rho_Pi (8);
 305       Rho_Pi (9);
 306       Rho_Pi (10);
 307       Rho_Pi (11);
 308       Rho_Pi (12);
 309       Rho_Pi (13);
 310       Rho_Pi (14);
 311       Rho_Pi (15);
 312       Rho_Pi (16);
 313       Rho_Pi (17);
 314       Rho_Pi (18);
 315       Rho_Pi (19);
 316       Rho_Pi (20);
 317       Rho_Pi (21);
 318       Rho_Pi (22);
 319       Rho_Pi (23);
 320
 321       //  Chi
 322
 323       Chi (0);
 324       Chi (5);
 325       Chi (10);
 326       Chi (15);
 327       Chi (20);
 328
 329       //  Iota
 330
 331       st[0] ^= keccakf_rndc[round];
 332     }
 333
 334     const u32x r0 = l32_from_64 (st[1]);
 335     const u32x r1 = h32_from_64 (st[1]);
 336     const u32x r2 = l32_from_64 (st[2]);
 337     const u32x r3 = h32_from_64 (st[2]);
 338
 339     #include VECT_COMPARE_M
 340   }
 341 }
 342
 343 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 344 {
 345 }
 346
 347 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 348 {
 349 }
 350
 351 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 352 {
 353   /**
 354    * modifier
 355    */
 356
 357   const u32 lid = threadIdx.x;
 358
 359   /**
 360    * base
 361    */
 362
 363   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 364
 365   if (gid >= gid_max) return;
 366
 367   /**
 368    * digest
 369    */
 370
 371   const u32 search[4] =
 372   {
 373     digests_buf[digests_offset].digest_buf[DGST_R0],
 374     digests_buf[digests_offset].digest_buf[DGST_R1],
 375     digests_buf[digests_offset].digest_buf[DGST_R2],
 376     digests_buf[digests_offset].digest_buf[DGST_R3]
 377   };
 378
 379   /**
 380    * constants
 381    */
 382
 383   const u64 keccakf_rndc[24] =
 384   {
 385     0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
 386     0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
 387     0x8000000080008081, 0x8000000000008009, 0x000000000000008a,
 388     0x0000000000000088, 0x0000000080008009, 0x000000008000000a,
 389     0x000000008000808b, 0x800000000000008b, 0x8000000000008089,
 390     0x8000000000008003, 0x8000000000008002, 0x8000000000000080,
 391     0x000000000000800a, 0x800000008000000a, 0x8000000080008081,
 392     0x8000000000008080, 0x0000000080000001, 0x8000000080008008
 393   };
 394
 395   const u32 keccakf_rotc[24] =
 396   {
 397      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
 398     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
 399   };
 400
 401   const u32 keccakf_piln[24] =
 402   {
 403     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
 404     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
 405   };
 406
 407   u32x wordl0[4];
 408
 409   wordl0[0] = pws[gid].i[ 0];
 410   wordl0[1] = pws[gid].i[ 1];
 411   wordl0[2] = pws[gid].i[ 2];
 412   wordl0[3] = pws[gid].i[ 3];
 413
 414   u32x wordl1[4];
 415
 416   wordl1[0] = pws[gid].i[ 4];
 417   wordl1[1] = pws[gid].i[ 5];
 418   wordl1[2] = pws[gid].i[ 6];
 419   wordl1[3] = pws[gid].i[ 7];
 420
 421   u32x wordl2[4];
 422
 423   wordl2[0] = 0;
 424   wordl2[1] = 0;
 425   wordl2[2] = 0;
 426   wordl2[3] = 0;
 427
 428   u32x wordl3[4];
 429
 430   wordl3[0] = 0;
 431   wordl3[1] = 0;
 432   wordl3[2] = 0;
 433   wordl3[3] = 0;
 434
 435   const u32 pw_l_len = pws[gid].pw_len;
 436
 437   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 438   {
 439     append_0x01_2 (wordl0, wordl1, pw_l_len);
 440
 441     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 442   }
 443
 444   /**
 445    * 0x80 keccak, very special
 446    */
 447
 448   const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
 449
 450   const u32 rsiz = 200 - (2 * mdlen);
 451
 452   const u32 add80w = (rsiz - 1) / 8;
 453
 454   /**
 455    * loop
 456    */
 457
 458   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 459   {
 460     const u32 pw_r_len = c_combs[il_pos].pw_len;
 461
 462     const u32 pw_len = pw_l_len + pw_r_len;
 463
 464     u32 wordr0[4];
 465
 466     wordr0[0] = c_combs[il_pos].i[0];
 467     wordr0[1] = c_combs[il_pos].i[1];
 468     wordr0[2] = c_combs[il_pos].i[2];
 469     wordr0[3] = c_combs[il_pos].i[3];
 470
 471     u32 wordr1[4];
 472
 473     wordr1[0] = c_combs[il_pos].i[4];
 474     wordr1[1] = c_combs[il_pos].i[5];
 475     wordr1[2] = c_combs[il_pos].i[6];
 476     wordr1[3] = c_combs[il_pos].i[7];
 477
 478     u32 wordr2[4];
 479
 480     wordr2[0] = 0;
 481     wordr2[1] = 0;
 482     wordr2[2] = 0;
 483     wordr2[3] = 0;
 484
 485     u32 wordr3[4];
 486
 487     wordr3[0] = 0;
 488     wordr3[1] = 0;
 489     wordr3[2] = 0;
 490     wordr3[3] = 0;
 491
 492     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 493     {
 494       append_0x01_2 (wordr0, wordr1, pw_r_len);
 495
 496       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 497     }
 498
 499     u32x w0[4];
 500
 501     w0[0] = wordl0[0] | wordr0[0];
 502     w0[1] = wordl0[1] | wordr0[1];
 503     w0[2] = wordl0[2] | wordr0[2];
 504     w0[3] = wordl0[3] | wordr0[3];
 505
 506     u32x w1[4];
 507
 508     w1[0] = wordl1[0] | wordr1[0];
 509     w1[1] = wordl1[1] | wordr1[1];
 510     w1[2] = wordl1[2] | wordr1[2];
 511     w1[3] = wordl1[3] | wordr1[3];
 512
 513     u32x w2[4];
 514
 515     w2[0] = wordl2[0] | wordr2[0];
 516     w2[1] = wordl2[1] | wordr2[1];
 517     w2[2] = wordl2[2] | wordr2[2];
 518     w2[3] = wordl2[3] | wordr2[3];
 519
 520     u32x w3[4];
 521
 522     w3[0] = wordl3[0] | wordr3[0];
 523     w3[1] = wordl3[1] | wordr3[1];
 524     w3[2] = pw_len * 8;
 525     w3[3] = 0;
 526
 527     u64x st[25];
 528
 529     st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
 530     st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
 531     st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
 532     st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
 533     st[ 4] = 0;
 534     st[ 5] = 0;
 535     st[ 6] = 0;
 536     st[ 7] = 0;
 537     st[ 8] = 0;
 538     st[ 9] = 0;
 539     st[10] = 0;
 540     st[11] = 0;
 541     st[12] = 0;
 542     st[13] = 0;
 543     st[14] = 0;
 544     st[15] = 0;
 545     st[16] = 0;
 546     st[17] = 0;
 547     st[18] = 0;
 548     st[19] = 0;
 549     st[20] = 0;
 550     st[21] = 0;
 551     st[22] = 0;
 552     st[23] = 0;
 553     st[24] = 0;
 554
 555     st[add80w] |= 0x8000000000000000;
 556
 557     int round;
 558
 559     for (round = 0; round < KECCAK_ROUNDS; round++)
 560     {
 561       // Theta
 562
 563       u64x bc0 = Theta1 (0);
 564       u64x bc1 = Theta1 (1);
 565       u64x bc2 = Theta1 (2);
 566       u64x bc3 = Theta1 (3);
 567       u64x bc4 = Theta1 (4);
 568
 569       u64x t;
 570
 571       t = bc4 ^ rotl64 (bc1, 1); Theta2 (0);
 572       t = bc0 ^ rotl64 (bc2, 1); Theta2 (1);
 573       t = bc1 ^ rotl64 (bc3, 1); Theta2 (2);
 574       t = bc2 ^ rotl64 (bc4, 1); Theta2 (3);
 575       t = bc3 ^ rotl64 (bc0, 1); Theta2 (4);
 576
 577       // Rho Pi
 578
 579       t = st[1];
 580
 581       Rho_Pi (0);
 582       Rho_Pi (1);
 583       Rho_Pi (2);
 584       Rho_Pi (3);
 585       Rho_Pi (4);
 586       Rho_Pi (5);
 587       Rho_Pi (6);
 588       Rho_Pi (7);
 589       Rho_Pi (8);
 590       Rho_Pi (9);
 591       Rho_Pi (10);
 592       Rho_Pi (11);
 593       Rho_Pi (12);
 594       Rho_Pi (13);
 595       Rho_Pi (14);
 596       Rho_Pi (15);
 597       Rho_Pi (16);
 598       Rho_Pi (17);
 599       Rho_Pi (18);
 600       Rho_Pi (19);
 601       Rho_Pi (20);
 602       Rho_Pi (21);
 603       Rho_Pi (22);
 604       Rho_Pi (23);
 605
 606       //  Chi
 607
 608       Chi (0);
 609       Chi (5);
 610       Chi (10);
 611       Chi (15);
 612       Chi (20);
 613
 614       //  Iota
 615
 616       st[0] ^= keccakf_rndc[round];
 617     }
 618
 619     const u32x r0 = l32_from_64 (st[1]);
 620     const u32x r1 = h32_from_64 (st[1]);
 621     const u32x r2 = l32_from_64 (st[2]);
 622     const u32x r3 = h32_from_64 (st[2]);
 623
 624     #include VECT_COMPARE_S
 625   }
 626 }
 627
 628 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 629 {
 630 }
 631
 632 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 633 {
 634 }