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