nv/m11500_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _CRC32_
   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__ const u32 crc32tab[0x100] =
  45 {
  46   0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
  47   0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
  48   0x0edb8832, 0x79dcb8a4, 0xe0d5e91e, 0x97d2d988,
  49   0x09b64c2b, 0x7eb17cbd, 0xe7b82d07, 0x90bf1d91,
  50   0x1db71064, 0x6ab020f2, 0xf3b97148, 0x84be41de,
  51   0x1adad47d, 0x6ddde4eb, 0xf4d4b551, 0x83d385c7,
  52   0x136c9856, 0x646ba8c0, 0xfd62f97a, 0x8a65c9ec,
  53   0x14015c4f, 0x63066cd9, 0xfa0f3d63, 0x8d080df5,
  54   0x3b6e20c8, 0x4c69105e, 0xd56041e4, 0xa2677172,
  55   0x3c03e4d1, 0x4b04d447, 0xd20d85fd, 0xa50ab56b,
  56   0x35b5a8fa, 0x42b2986c, 0xdbbbc9d6, 0xacbcf940,
  57   0x32d86ce3, 0x45df5c75, 0xdcd60dcf, 0xabd13d59,
  58   0x26d930ac, 0x51de003a, 0xc8d75180, 0xbfd06116,
  59   0x21b4f4b5, 0x56b3c423, 0xcfba9599, 0xb8bda50f,
  60   0x2802b89e, 0x5f058808, 0xc60cd9b2, 0xb10be924,
  61   0x2f6f7c87, 0x58684c11, 0xc1611dab, 0xb6662d3d,
  62   0x76dc4190, 0x01db7106, 0x98d220bc, 0xefd5102a,
  63   0x71b18589, 0x06b6b51f, 0x9fbfe4a5, 0xe8b8d433,
  64   0x7807c9a2, 0x0f00f934, 0x9609a88e, 0xe10e9818,
  65   0x7f6a0dbb, 0x086d3d2d, 0x91646c97, 0xe6635c01,
  66   0x6b6b51f4, 0x1c6c6162, 0x856530d8, 0xf262004e,
  67   0x6c0695ed, 0x1b01a57b, 0x8208f4c1, 0xf50fc457,
  68   0x65b0d9c6, 0x12b7e950, 0x8bbeb8ea, 0xfcb9887c,
  69   0x62dd1ddf, 0x15da2d49, 0x8cd37cf3, 0xfbd44c65,
  70   0x4db26158, 0x3ab551ce, 0xa3bc0074, 0xd4bb30e2,
  71   0x4adfa541, 0x3dd895d7, 0xa4d1c46d, 0xd3d6f4fb,
  72   0x4369e96a, 0x346ed9fc, 0xad678846, 0xda60b8d0,
  73   0x44042d73, 0x33031de5, 0xaa0a4c5f, 0xdd0d7cc9,
  74   0x5005713c, 0x270241aa, 0xbe0b1010, 0xc90c2086,
  75   0x5768b525, 0x206f85b3, 0xb966d409, 0xce61e49f,
  76   0x5edef90e, 0x29d9c998, 0xb0d09822, 0xc7d7a8b4,
  77   0x59b33d17, 0x2eb40d81, 0xb7bd5c3b, 0xc0ba6cad,
  78   0xedb88320, 0x9abfb3b6, 0x03b6e20c, 0x74b1d29a,
  79   0xead54739, 0x9dd277af, 0x04db2615, 0x73dc1683,
  80   0xe3630b12, 0x94643b84, 0x0d6d6a3e, 0x7a6a5aa8,
  81   0xe40ecf0b, 0x9309ff9d, 0x0a00ae27, 0x7d079eb1,
  82   0xf00f9344, 0x8708a3d2, 0x1e01f268, 0x6906c2fe,
  83   0xf762575d, 0x806567cb, 0x196c3671, 0x6e6b06e7,
  84   0xfed41b76, 0x89d32be0, 0x10da7a5a, 0x67dd4acc,
  85   0xf9b9df6f, 0x8ebeeff9, 0x17b7be43, 0x60b08ed5,
  86   0xd6d6a3e8, 0xa1d1937e, 0x38d8c2c4, 0x4fdff252,
  87   0xd1bb67f1, 0xa6bc5767, 0x3fb506dd, 0x48b2364b,
  88   0xd80d2bda, 0xaf0a1b4c, 0x36034af6, 0x41047a60,
  89   0xdf60efc3, 0xa867df55, 0x316e8eef, 0x4669be79,
  90   0xcb61b38c, 0xbc66831a, 0x256fd2a0, 0x5268e236,
  91   0xcc0c7795, 0xbb0b4703, 0x220216b9, 0x5505262f,
  92   0xc5ba3bbe, 0xb2bd0b28, 0x2bb45a92, 0x5cb36a04,
  93   0xc2d7ffa7, 0xb5d0cf31, 0x2cd99e8b, 0x5bdeae1d,
  94   0x9b64c2b0, 0xec63f226, 0x756aa39c, 0x026d930a,
  95   0x9c0906a9, 0xeb0e363f, 0x72076785, 0x05005713,
  96   0x95bf4a82, 0xe2b87a14, 0x7bb12bae, 0x0cb61b38,
  97   0x92d28e9b, 0xe5d5be0d, 0x7cdcefb7, 0x0bdbdf21,
  98   0x86d3d2d4, 0xf1d4e242, 0x68ddb3f8, 0x1fda836e,
  99   0x81be16cd, 0xf6b9265b, 0x6fb077e1, 0x18b74777,
 100   0x88085ae6, 0xff0f6a70, 0x66063bca, 0x11010b5c,
 101   0x8f659eff, 0xf862ae69, 0x616bffd3, 0x166ccf45,
 102   0xa00ae278, 0xd70dd2ee, 0x4e048354, 0x3903b3c2,
 103   0xa7672661, 0xd06016f7, 0x4969474d, 0x3e6e77db,
 104   0xaed16a4a, 0xd9d65adc, 0x40df0b66, 0x37d83bf0,
 105   0xa9bcae53, 0xdebb9ec5, 0x47b2cf7f, 0x30b5ffe9,
 106   0xbdbdf21c, 0xcabac28a, 0x53b39330, 0x24b4a3a6,
 107   0xbad03605, 0xcdd70693, 0x54de5729, 0x23d967bf,
 108   0xb3667a2e, 0xc4614ab8, 0x5d681b02, 0x2a6f2b94,
 109   0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d
 110 };
 111
 112 __device__ static u32x round_crc32 (u32x a, const u32x v)
 113 {
 114   const u32x k = (a ^ v) & 0xff;
 115
 116   const u32x s = a >> 8;
 117
 118   #ifdef VECT_SIZE1
 119   a    = crc32tab[k];
 120   #endif
 121
 122   #ifdef VECT_SIZE2
 123   a.x = crc32tab[k.x];
 124   a.y = crc32tab[k.y];
 125   #endif
 126
 127   #ifdef VECT_SIZE4
 128   a.x = crc32tab[k.x];
 129   a.y = crc32tab[k.y];
 130   a.z = crc32tab[k.z];
 131   a.w = crc32tab[k.w];
 132   #endif
 133
 134   a ^= s;
 135
 136   return a;
 137 }
 138
 139 __device__ static u32x crc32 (const u32x w[16], const u32 pw_len, const u32 iv)
 140 {
 141   u32x a = iv ^ ~0;
 142
 143   if (pw_len >=  1) a = round_crc32 (a, w[0] >>  0);
 144   if (pw_len >=  2) a = round_crc32 (a, w[0] >>  8);
 145   if (pw_len >=  3) a = round_crc32 (a, w[0] >> 16);
 146   if (pw_len >=  4) a = round_crc32 (a, w[0] >> 24);
 147
 148   for (u32 i = 4, j = 1; i < pw_len; i += 4, j += 1)
 149   {
 150     if (pw_len >= (i + 1)) a = round_crc32 (a, w[j] >>  0);
 151     if (pw_len >= (i + 2)) a = round_crc32 (a, w[j] >>  8);
 152     if (pw_len >= (i + 3)) a = round_crc32 (a, w[j] >> 16);
 153     if (pw_len >= (i + 4)) a = round_crc32 (a, w[j] >> 24);
 154   }
 155
 156   return ~a;
 157 }
 158
 159 __device__ __constant__ u32x c_bfs[1024];
 160
 161 __device__ static void m11500m (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, 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)
 162 {
 163   /**
 164    * modifier
 165    */
 166
 167   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 168   const u32 lid = threadIdx.x;
 169
 170
 171   /**
 172    * digest
 173    */
 174
 175   const u32 iv = salt_bufs[salt_pos].salt_buf[0];
 176
 177   /**
 178    * loop
 179    */
 180
 181   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 182
 183   u32x w0l = w[0];
 184
 185   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 186   {
 187     const u32x w0r = c_bfs[il_pos];
 188
 189     const u32x w0 = w0l | w0r;
 190
 191     u32x w_t[16];
 192
 193     w_t[ 0] = w0;
 194     w_t[ 1] = w[ 1];
 195     w_t[ 2] = w[ 2];
 196     w_t[ 3] = w[ 3];
 197     w_t[ 4] = w[ 4];
 198     w_t[ 5] = w[ 5];
 199     w_t[ 6] = w[ 6];
 200     w_t[ 7] = w[ 7];
 201     w_t[ 8] = w[ 8];
 202     w_t[ 9] = w[ 9];
 203     w_t[10] = w[10];
 204     w_t[11] = w[11];
 205     w_t[12] = w[12];
 206     w_t[13] = w[13];
 207     w_t[14] = w[14];
 208     w_t[15] = w[15];
 209
 210     u32x a = crc32 (w_t, pw_len, iv);
 211     u32x b = 0;
 212
 213     const u32x r0 = a;
 214     const u32x r1 = b;
 215     const u32x r2 = 0;
 216     const u32x r3 = 0;
 217
 218     #include VECT_COMPARE_M
 219   }
 220 }
 221
 222 __device__ static void m11500s (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, 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)
 223 {
 224   /**
 225    * modifier
 226    */
 227
 228   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 229   const u32 lid = threadIdx.x;
 230
 231   /**
 232    * digest
 233    */
 234
 235   const u32 iv = salt_bufs[salt_pos].salt_buf[0];
 236
 237   const u32 search[4] =
 238   {
 239     digests_buf[digests_offset].digest_buf[DGST_R0],
 240     digests_buf[digests_offset].digest_buf[DGST_R1],
 241     digests_buf[digests_offset].digest_buf[DGST_R2],
 242     digests_buf[digests_offset].digest_buf[DGST_R3]
 243   };
 244
 245   /**
 246    * loop
 247    */
 248
 249   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 250
 251   u32x w0l = w[0];
 252
 253   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 254   {
 255     const u32x w0r = c_bfs[il_pos];
 256
 257     const u32x w0 = w0l | w0r;
 258
 259     u32x w_t[16];
 260
 261     w_t[ 0] = w0;
 262     w_t[ 1] = w[ 1];
 263     w_t[ 2] = w[ 2];
 264     w_t[ 3] = w[ 3];
 265     w_t[ 4] = w[ 4];
 266     w_t[ 5] = w[ 5];
 267     w_t[ 6] = w[ 6];
 268     w_t[ 7] = w[ 7];
 269     w_t[ 8] = w[ 8];
 270     w_t[ 9] = w[ 9];
 271     w_t[10] = w[10];
 272     w_t[11] = w[11];
 273     w_t[12] = w[12];
 274     w_t[13] = w[13];
 275     w_t[14] = w[14];
 276     w_t[15] = w[15];
 277
 278     u32x a = crc32 (w_t, pw_len, iv);
 279     u32x b = 0;
 280
 281     const u32x r0 = a;
 282     const u32x r1 = b;
 283     const u32x r2 = 0;
 284     const u32x r3 = 0;
 285
 286     #include VECT_COMPARE_S
 287   }
 288 }
 289
 290 extern "C" __global__ void __launch_bounds__ (256, 1) m11500_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 291 {
 292   /**
 293    * base
 294    */
 295
 296   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 297
 298   if (gid >= gid_max) return;
 299
 300   u32 w[16];
 301
 302   w[ 0] = pws[gid].i[ 0];
 303   w[ 1] = pws[gid].i[ 1];
 304   w[ 2] = pws[gid].i[ 2];
 305   w[ 3] = pws[gid].i[ 3];
 306   w[ 4] = 0;
 307   w[ 5] = 0;
 308   w[ 6] = 0;
 309   w[ 7] = 0;
 310   w[ 8] = 0;
 311   w[ 9] = 0;
 312   w[10] = 0;
 313   w[11] = 0;
 314   w[12] = 0;
 315   w[13] = 0;
 316   w[14] = 0;
 317   w[15] = 0;
 318
 319   const u32 pw_len = pws[gid].pw_len;
 320
 321   /**
 322    * main
 323    */
 324
 325   m11500m (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);
 326 }
 327
 328 extern "C" __global__ void __launch_bounds__ (256, 1) m11500_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 329 {
 330   /**
 331    * base
 332    */
 333
 334   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 335
 336   if (gid >= gid_max) return;
 337
 338   u32 w[16];
 339
 340   w[ 0] = pws[gid].i[ 0];
 341   w[ 1] = pws[gid].i[ 1];
 342   w[ 2] = pws[gid].i[ 2];
 343   w[ 3] = pws[gid].i[ 3];
 344   w[ 4] = pws[gid].i[ 4];
 345   w[ 5] = pws[gid].i[ 5];
 346   w[ 6] = pws[gid].i[ 6];
 347   w[ 7] = pws[gid].i[ 7];
 348   w[ 8] = 0;
 349   w[ 9] = 0;
 350   w[10] = 0;
 351   w[11] = 0;
 352   w[12] = 0;
 353   w[13] = 0;
 354   w[14] = 0;
 355   w[15] = 0;
 356
 357   const u32 pw_len = pws[gid].pw_len;
 358
 359   /**
 360    * main
 361    */
 362
 363   m11500m (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);
 364 }
 365
 366 extern "C" __global__ void __launch_bounds__ (256, 1) m11500_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 367 {
 368   /**
 369    * base
 370    */
 371
 372   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 373
 374   if (gid >= gid_max) return;
 375
 376   u32 w[16];
 377
 378   w[ 0] = pws[gid].i[ 0];
 379   w[ 1] = pws[gid].i[ 1];
 380   w[ 2] = pws[gid].i[ 2];
 381   w[ 3] = pws[gid].i[ 3];
 382   w[ 4] = pws[gid].i[ 4];
 383   w[ 5] = pws[gid].i[ 5];
 384   w[ 6] = pws[gid].i[ 6];
 385   w[ 7] = pws[gid].i[ 7];
 386   w[ 8] = pws[gid].i[ 8];
 387   w[ 9] = pws[gid].i[ 9];
 388   w[10] = pws[gid].i[10];
 389   w[11] = pws[gid].i[11];
 390   w[12] = pws[gid].i[12];
 391   w[13] = pws[gid].i[13];
 392   w[14] = pws[gid].i[14];
 393   w[15] = pws[gid].i[15];
 394
 395   const u32 pw_len = pws[gid].pw_len;
 396
 397   /**
 398    * main
 399    */
 400
 401   m11500m (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);
 402 }
 403
 404 extern "C" __global__ void __launch_bounds__ (256, 1) m11500_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 405 {
 406   /**
 407    * base
 408    */
 409
 410   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 411
 412   if (gid >= gid_max) return;
 413
 414   u32 w[16];
 415
 416   w[ 0] = pws[gid].i[ 0];
 417   w[ 1] = pws[gid].i[ 1];
 418   w[ 2] = pws[gid].i[ 2];
 419   w[ 3] = pws[gid].i[ 3];
 420   w[ 4] = 0;
 421   w[ 5] = 0;
 422   w[ 6] = 0;
 423   w[ 7] = 0;
 424   w[ 8] = 0;
 425   w[ 9] = 0;
 426   w[10] = 0;
 427   w[11] = 0;
 428   w[12] = 0;
 429   w[13] = 0;
 430   w[14] = 0;
 431   w[15] = 0;
 432
 433   const u32 pw_len = pws[gid].pw_len;
 434
 435   /**
 436    * main
 437    */
 438
 439   m11500s (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);
 440 }
 441
 442 extern "C" __global__ void __launch_bounds__ (256, 1) m11500_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 443 {
 444   /**
 445    * base
 446    */
 447
 448   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 449
 450   if (gid >= gid_max) return;
 451
 452   u32 w[16];
 453
 454   w[ 0] = pws[gid].i[ 0];
 455   w[ 1] = pws[gid].i[ 1];
 456   w[ 2] = pws[gid].i[ 2];
 457   w[ 3] = pws[gid].i[ 3];
 458   w[ 4] = pws[gid].i[ 4];
 459   w[ 5] = pws[gid].i[ 5];
 460   w[ 6] = pws[gid].i[ 6];
 461   w[ 7] = pws[gid].i[ 7];
 462   w[ 8] = 0;
 463   w[ 9] = 0;
 464   w[10] = 0;
 465   w[11] = 0;
 466   w[12] = 0;
 467   w[13] = 0;
 468   w[14] = 0;
 469   w[15] = 0;
 470
 471   const u32 pw_len = pws[gid].pw_len;
 472
 473   /**
 474    * main
 475    */
 476
 477   m11500s (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);
 478 }
 479
 480 extern "C" __global__ void __launch_bounds__ (256, 1) m11500_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 481 {
 482   /**
 483    * base
 484    */
 485
 486   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 487
 488   if (gid >= gid_max) return;
 489
 490   u32 w[16];
 491
 492   w[ 0] = pws[gid].i[ 0];
 493   w[ 1] = pws[gid].i[ 1];
 494   w[ 2] = pws[gid].i[ 2];
 495   w[ 3] = pws[gid].i[ 3];
 496   w[ 4] = pws[gid].i[ 4];
 497   w[ 5] = pws[gid].i[ 5];
 498   w[ 6] = pws[gid].i[ 6];
 499   w[ 7] = pws[gid].i[ 7];
 500   w[ 8] = pws[gid].i[ 8];
 501   w[ 9] = pws[gid].i[ 9];
 502   w[10] = pws[gid].i[10];
 503   w[11] = pws[gid].i[11];
 504   w[12] = pws[gid].i[12];
 505   w[13] = pws[gid].i[13];
 506   w[14] = pws[gid].i[14];
 507   w[15] = pws[gid].i[15];
 508
 509   const u32 pw_len = pws[gid].pw_len;
 510
 511   /**
 512    * main
 513    */
 514
 515   m11500s (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);
 516 }