nv/m00200_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _MYSQL323_
   7 #define _SCALAR_
   8
   9 #include "include/constants.h"
  10 #include "include/kernel_vendor.h"
  11
  12 #ifdef  VLIW1
  13 #define VECT_SIZE1
  14 #endif
  15
  16 #ifdef  VLIW2
  17 #define VECT_SIZE4
  18 #endif
  19
  20 #define DGST_R0 0
  21 #define DGST_R1 1
  22 #define DGST_R2 2
  23 #define DGST_R3 3
  24
  25 #include "include/kernel_functions.c"
  26 #include "types_nv.c"
  27 #include "common_nv.c"
  28
  29 #ifdef  VECT_SIZE1
  30 #define VECT_COMPARE_S "check_single_vect1_comp4_warp.c"
  31 #define VECT_COMPARE_M "check_multi_vect1_comp4_warp.c"
  32 #endif
  33
  34 #ifdef  VECT_SIZE2
  35 #define VECT_COMPARE_S "check_single_vect2_comp4_warp.c"
  36 #define VECT_COMPARE_M "check_multi_vect2_comp4_warp.c"
  37 #endif
  38
  39 #ifdef  VECT_SIZE4
  40 #define VECT_COMPARE_S "check_single_vect4_comp4_warp.c"
  41 #define VECT_COMPARE_M "check_multi_vect4_comp4_warp.c"
  42 #endif
  43
  44 __device__ __constant__ u32x c_bfs[1024];
  45
  46 __device__ static void m00200m (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, 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)
  47 {
  48   /**
  49    * modifier
  50    */
  51
  52   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  53   const u32 lid = threadIdx.x;
  54
  55   /**
  56    * loop
  57    */
  58
  59   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
  60
  61   u32x w0l = w[0];
  62
  63   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
  64   {
  65     const u32x w0r = c_bfs[il_pos];
  66
  67     const u32x w0 = w0l | w0r;
  68
  69     u32x a = MYSQL323_A;
  70     u32x b = MYSQL323_B;
  71
  72     u32x add = 7;
  73
  74     #define ROUND(v)                              \
  75     {                                             \
  76       a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
  77       b += (b << 8) ^ a;                          \
  78       add += v;                                   \
  79     }
  80
  81     if (pw_len >= 4)
  82     {
  83       ROUND ((w0 >>  0) & 0xff);
  84       ROUND ((w0 >>  8) & 0xff);
  85       ROUND ((w0 >> 16) & 0xff);
  86       ROUND ((w0 >> 24) & 0xff);
  87     }
  88     else if (pw_len == 3)
  89     {
  90       ROUND ((w0 >>  0) & 0xff);
  91       ROUND ((w0 >>  8) & 0xff);
  92       ROUND ((w0 >> 16) & 0xff);
  93     }
  94     else if (pw_len == 2)
  95     {
  96       ROUND ((w0 >>  0) & 0xff);
  97       ROUND ((w0 >>  8) & 0xff);
  98     }
  99     else if (pw_len == 1)
 100     {
 101       ROUND ((w0 >>  0) & 0xff);
 102     }
 103
 104     int i;
 105     int j;
 106
 107     for (i = 4, j = 1; i <= (int) pw_len - 4; i += 4, j += 1)
 108     {
 109       const u32 wj = w[j];
 110
 111       ROUND ((wj >>  0) & 0xff);
 112       ROUND ((wj >>  8) & 0xff);
 113       ROUND ((wj >> 16) & 0xff);
 114       ROUND ((wj >> 24) & 0xff);
 115     }
 116
 117     const u32 wj = w[j];
 118
 119     const u32 left = pw_len - i;
 120
 121     if (left == 3)
 122     {
 123       ROUND ((wj >>  0) & 0xff);
 124       ROUND ((wj >>  8) & 0xff);
 125       ROUND ((wj >> 16) & 0xff);
 126     }
 127     else if (left == 2)
 128     {
 129       ROUND ((wj >>  0) & 0xff);
 130       ROUND ((wj >>  8) & 0xff);
 131     }
 132     else if (left == 1)
 133     {
 134       ROUND ((wj >>  0) & 0xff);
 135     }
 136
 137     a &= 0x7fffffff;
 138     b &= 0x7fffffff;
 139
 140     const u32x r0 = a;
 141     const u32x r1 = b;
 142     const u32x r2 = 0;
 143     const u32x r3 = 0;
 144
 145     #include VECT_COMPARE_M
 146   }
 147 }
 148
 149 __device__ static void m00200s (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, 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)
 150 {
 151   /**
 152    * modifier
 153    */
 154
 155   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 156   const u32 lid = threadIdx.x;
 157
 158   /**
 159    * digest
 160    */
 161
 162   const u32 search[4] =
 163   {
 164     digests_buf[digests_offset].digest_buf[DGST_R0],
 165     digests_buf[digests_offset].digest_buf[DGST_R1],
 166     digests_buf[digests_offset].digest_buf[DGST_R2],
 167     digests_buf[digests_offset].digest_buf[DGST_R3]
 168   };
 169
 170   /**
 171    * loop
 172    */
 173
 174   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 175
 176   u32x w0l = w[0];
 177
 178   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 179   {
 180     const u32x w0r = c_bfs[il_pos];
 181
 182     const u32x w0 = w0l | w0r;
 183
 184     u32x a = MYSQL323_A;
 185     u32x b = MYSQL323_B;
 186
 187     u32x add = 7;
 188
 189     #define ROUND(v)                              \
 190     {                                             \
 191       a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
 192       b += (b << 8) ^ a;                          \
 193       add += v;                                   \
 194     }
 195
 196     if (pw_len >= 4)
 197     {
 198       ROUND ((w0 >>  0) & 0xff);
 199       ROUND ((w0 >>  8) & 0xff);
 200       ROUND ((w0 >> 16) & 0xff);
 201       ROUND ((w0 >> 24) & 0xff);
 202     }
 203     else if (pw_len == 3)
 204     {
 205       ROUND ((w0 >>  0) & 0xff);
 206       ROUND ((w0 >>  8) & 0xff);
 207       ROUND ((w0 >> 16) & 0xff);
 208     }
 209     else if (pw_len == 2)
 210     {
 211       ROUND ((w0 >>  0) & 0xff);
 212       ROUND ((w0 >>  8) & 0xff);
 213     }
 214     else if (pw_len == 1)
 215     {
 216       ROUND ((w0 >>  0) & 0xff);
 217     }
 218
 219     int i;
 220     int j;
 221
 222     for (i = 4, j = 1; i <= (int) pw_len - 4; i += 4, j += 1)
 223     {
 224       const u32 wj = w[j];
 225
 226       ROUND ((wj >>  0) & 0xff);
 227       ROUND ((wj >>  8) & 0xff);
 228       ROUND ((wj >> 16) & 0xff);
 229       ROUND ((wj >> 24) & 0xff);
 230     }
 231
 232     const u32 wj = w[j];
 233
 234     const u32 left = pw_len - i;
 235
 236     if (left == 3)
 237     {
 238       ROUND ((wj >>  0) & 0xff);
 239       ROUND ((wj >>  8) & 0xff);
 240       ROUND ((wj >> 16) & 0xff);
 241     }
 242     else if (left == 2)
 243     {
 244       ROUND ((wj >>  0) & 0xff);
 245       ROUND ((wj >>  8) & 0xff);
 246     }
 247     else if (left == 1)
 248     {
 249       ROUND ((wj >>  0) & 0xff);
 250     }
 251
 252     a &= 0x7fffffff;
 253     b &= 0x7fffffff;
 254
 255     const u32x r0 = a;
 256     const u32x r1 = b;
 257     const u32x r2 = 0;
 258     const u32x r3 = 0;
 259
 260     #include VECT_COMPARE_S
 261   }
 262 }
 263
 264 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, 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)
 265 {
 266   /**
 267    * base
 268    */
 269
 270   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 271
 272   if (gid >= gid_max) return;
 273
 274   u32 w[16];
 275
 276   w[ 0] = pws[gid].i[ 0];
 277   w[ 1] = pws[gid].i[ 1];
 278   w[ 2] = pws[gid].i[ 2];
 279   w[ 3] = pws[gid].i[ 3];
 280   w[ 4] = 0;
 281   w[ 5] = 0;
 282   w[ 6] = 0;
 283   w[ 7] = 0;
 284   w[ 8] = 0;
 285   w[ 9] = 0;
 286   w[10] = 0;
 287   w[11] = 0;
 288   w[12] = 0;
 289   w[13] = 0;
 290   w[14] = 0;
 291   w[15] = 0;
 292
 293   const u32 pw_len = pws[gid].pw_len;
 294
 295   /**
 296    * main
 297    */
 298
 299   m00200m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 300 }
 301
 302 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, 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)
 303 {
 304   /**
 305    * base
 306    */
 307
 308   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 309
 310   if (gid >= gid_max) return;
 311
 312   u32 w[16];
 313
 314   w[ 0] = pws[gid].i[ 0];
 315   w[ 1] = pws[gid].i[ 1];
 316   w[ 2] = pws[gid].i[ 2];
 317   w[ 3] = pws[gid].i[ 3];
 318   w[ 4] = pws[gid].i[ 4];
 319   w[ 5] = pws[gid].i[ 5];
 320   w[ 6] = pws[gid].i[ 6];
 321   w[ 7] = pws[gid].i[ 7];
 322   w[ 8] = 0;
 323   w[ 9] = 0;
 324   w[10] = 0;
 325   w[11] = 0;
 326   w[12] = 0;
 327   w[13] = 0;
 328   w[14] = 0;
 329   w[15] = 0;
 330
 331   const u32 pw_len = pws[gid].pw_len;
 332
 333   /**
 334    * main
 335    */
 336
 337   m00200m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 338 }
 339
 340 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, 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)
 341 {
 342   /**
 343    * base
 344    */
 345
 346   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 347
 348   if (gid >= gid_max) return;
 349
 350   u32 w[16];
 351
 352   w[ 0] = pws[gid].i[ 0];
 353   w[ 1] = pws[gid].i[ 1];
 354   w[ 2] = pws[gid].i[ 2];
 355   w[ 3] = pws[gid].i[ 3];
 356   w[ 4] = pws[gid].i[ 4];
 357   w[ 5] = pws[gid].i[ 5];
 358   w[ 6] = pws[gid].i[ 6];
 359   w[ 7] = pws[gid].i[ 7];
 360   w[ 8] = pws[gid].i[ 8];
 361   w[ 9] = pws[gid].i[ 9];
 362   w[10] = pws[gid].i[10];
 363   w[11] = pws[gid].i[11];
 364   w[12] = pws[gid].i[12];
 365   w[13] = pws[gid].i[13];
 366   w[14] = pws[gid].i[14];
 367   w[15] = pws[gid].i[15];
 368
 369   const u32 pw_len = pws[gid].pw_len;
 370
 371   /**
 372    * main
 373    */
 374
 375   m00200m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 376 }
 377
 378 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, 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)
 379 {
 380   /**
 381    * base
 382    */
 383
 384   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 385
 386   if (gid >= gid_max) return;
 387
 388   u32 w[16];
 389
 390   w[ 0] = pws[gid].i[ 0];
 391   w[ 1] = pws[gid].i[ 1];
 392   w[ 2] = pws[gid].i[ 2];
 393   w[ 3] = pws[gid].i[ 3];
 394   w[ 4] = 0;
 395   w[ 5] = 0;
 396   w[ 6] = 0;
 397   w[ 7] = 0;
 398   w[ 8] = 0;
 399   w[ 9] = 0;
 400   w[10] = 0;
 401   w[11] = 0;
 402   w[12] = 0;
 403   w[13] = 0;
 404   w[14] = 0;
 405   w[15] = 0;
 406
 407   const u32 pw_len = pws[gid].pw_len;
 408
 409   /**
 410    * main
 411    */
 412
 413   m00200s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 414 }
 415
 416 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, 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)
 417 {
 418   /**
 419    * base
 420    */
 421
 422   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 423
 424   if (gid >= gid_max) return;
 425
 426   u32 w[16];
 427
 428   w[ 0] = pws[gid].i[ 0];
 429   w[ 1] = pws[gid].i[ 1];
 430   w[ 2] = pws[gid].i[ 2];
 431   w[ 3] = pws[gid].i[ 3];
 432   w[ 4] = pws[gid].i[ 4];
 433   w[ 5] = pws[gid].i[ 5];
 434   w[ 6] = pws[gid].i[ 6];
 435   w[ 7] = pws[gid].i[ 7];
 436   w[ 8] = 0;
 437   w[ 9] = 0;
 438   w[10] = 0;
 439   w[11] = 0;
 440   w[12] = 0;
 441   w[13] = 0;
 442   w[14] = 0;
 443   w[15] = 0;
 444
 445   const u32 pw_len = pws[gid].pw_len;
 446
 447   /**
 448    * main
 449    */
 450
 451   m00200s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 452 }
 453
 454 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, 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)
 455 {
 456   /**
 457    * base
 458    */
 459
 460   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 461
 462   if (gid >= gid_max) return;
 463
 464   u32 w[16];
 465
 466   w[ 0] = pws[gid].i[ 0];
 467   w[ 1] = pws[gid].i[ 1];
 468   w[ 2] = pws[gid].i[ 2];
 469   w[ 3] = pws[gid].i[ 3];
 470   w[ 4] = pws[gid].i[ 4];
 471   w[ 5] = pws[gid].i[ 5];
 472   w[ 6] = pws[gid].i[ 6];
 473   w[ 7] = pws[gid].i[ 7];
 474   w[ 8] = pws[gid].i[ 8];
 475   w[ 9] = pws[gid].i[ 9];
 476   w[10] = pws[gid].i[10];
 477   w[11] = pws[gid].i[11];
 478   w[12] = pws[gid].i[12];
 479   w[13] = pws[gid].i[13];
 480   w[14] = pws[gid].i[14];
 481   w[15] = pws[gid].i[15];
 482
 483   const u32 pw_len = pws[gid].pw_len;
 484
 485   /**
 486    * main
 487    */
 488
 489   m00200s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 490 }