nv/m10100_a3.cu

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