nv/m11400_a1.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _MD5_
   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 0
  20 #define DGST_R1 3
  21 #define DGST_R2 2
  22 #define DGST_R3 1
  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 #ifdef VECT_SIZE1
  44 #define uint_to_hex_lower8(i) l_bin2asc[(i)]
  45 #endif
  46
  47 #ifdef VECT_SIZE2
  48 #define uint_to_hex_lower8(i) u32x (l_bin2asc[(i).x], l_bin2asc[(i).y])
  49 #endif
  50
  51 #ifdef VECT_SIZE4
  52 #define uint_to_hex_lower8(i) u32x (l_bin2asc[(i).x], l_bin2asc[(i).y], l_bin2asc[(i).z], l_bin2asc[(i).w])
  53 #endif
  54
  55 __device__ static u32 memcat32 (u32x block0[16], u32x block1[16], const u32 block_len, const u32x append0[4], const u32x append1[4], const u32x append2[4], const u32x append3[4], const u32 append_len)
  56 {
  57   const u32 mod = block_len & 3;
  58   const u32 div = block_len / 4;
  59
  60   const int offset_minus_4 = 4 - mod;
  61
  62   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
  63
  64   u32x append0_t[4];
  65
  66   append0_t[0] = __byte_perm (         0, append0[0], selector);
  67   append0_t[1] = __byte_perm (append0[0], append0[1], selector);
  68   append0_t[2] = __byte_perm (append0[1], append0[2], selector);
  69   append0_t[3] = __byte_perm (append0[2], append0[3], selector);
  70
  71   u32x append1_t[4];
  72
  73   append1_t[0] = __byte_perm (append0[3], append1[0], selector);
  74   append1_t[1] = __byte_perm (append1[0], append1[1], selector);
  75   append1_t[2] = __byte_perm (append1[1], append1[2], selector);
  76   append1_t[3] = __byte_perm (append1[2], append1[3], selector);
  77
  78   u32x append2_t[4];
  79
  80   append2_t[0] = __byte_perm (append1[3], append2[0], selector);
  81   append2_t[1] = __byte_perm (append2[0], append2[1], selector);
  82   append2_t[2] = __byte_perm (append2[1], append2[2], selector);
  83   append2_t[3] = __byte_perm (append2[2], append2[3], selector);
  84
  85   u32x append3_t[4];
  86
  87   append3_t[0] = __byte_perm (append2[3], append3[0], selector);
  88   append3_t[1] = __byte_perm (append3[0], append3[1], selector);
  89   append3_t[2] = __byte_perm (append3[1], append3[2], selector);
  90   append3_t[3] = __byte_perm (append3[2], append3[3], selector);
  91
  92   u32x append4_t[4];
  93
  94   append4_t[0] = __byte_perm (append3[3],          0, selector);
  95   append4_t[1] = 0;
  96   append4_t[2] = 0;
  97   append4_t[3] = 0;
  98
  99   switch (div)
 100   {
 101     case  0:  block0[ 0] |= append0_t[0];
 102               block0[ 1]  = append0_t[1];
 103               block0[ 2]  = append0_t[2];
 104               block0[ 3]  = append0_t[3];
 105
 106               block0[ 4]  = append1_t[0];
 107               block0[ 5]  = append1_t[1];
 108               block0[ 6]  = append1_t[2];
 109               block0[ 7]  = append1_t[3];
 110
 111               block0[ 8]  = append2_t[0];
 112               block0[ 9]  = append2_t[1];
 113               block0[10]  = append2_t[2];
 114               block0[11]  = append2_t[3];
 115
 116               block0[12]  = append3_t[0];
 117               block0[13]  = append3_t[1];
 118               block0[14]  = append3_t[2];
 119               block0[15]  = append3_t[3];
 120
 121               block1[ 0]  = append4_t[0];
 122               block1[ 1]  = append4_t[1];
 123               block1[ 2]  = append4_t[2];
 124               block1[ 3]  = append4_t[3];
 125               break;
 126
 127     case  1:  block0[ 1] |= append0_t[0];
 128               block0[ 2]  = append0_t[1];
 129               block0[ 3]  = append0_t[2];
 130               block0[ 4]  = append0_t[3];
 131
 132               block0[ 5]  = append1_t[0];
 133               block0[ 6]  = append1_t[1];
 134               block0[ 7]  = append1_t[2];
 135               block0[ 8]  = append1_t[3];
 136
 137               block0[ 9]  = append2_t[0];
 138               block0[10]  = append2_t[1];
 139               block0[11]  = append2_t[2];
 140               block0[12]  = append2_t[3];
 141
 142               block0[13]  = append3_t[0];
 143               block0[14]  = append3_t[1];
 144               block0[15]  = append3_t[2];
 145               block1[ 0]  = append3_t[3];
 146
 147               block1[ 1]  = append4_t[0];
 148               block1[ 2]  = append4_t[1];
 149               block1[ 3]  = append4_t[2];
 150               block1[ 4]  = append4_t[3];
 151               break;
 152
 153     case  2:  block0[ 2] |= append0_t[0];
 154               block0[ 3]  = append0_t[1];
 155               block0[ 4]  = append0_t[2];
 156               block0[ 5]  = append0_t[3];
 157
 158               block0[ 6]  = append1_t[0];
 159               block0[ 7]  = append1_t[1];
 160               block0[ 8]  = append1_t[2];
 161               block0[ 9]  = append1_t[3];
 162
 163               block0[10]  = append2_t[0];
 164               block0[11]  = append2_t[1];
 165               block0[12]  = append2_t[2];
 166               block0[13]  = append2_t[3];
 167
 168               block0[14]  = append3_t[0];
 169               block0[15]  = append3_t[1];
 170               block1[ 0]  = append3_t[2];
 171               block1[ 1]  = append3_t[3];
 172
 173               block1[ 2]  = append4_t[0];
 174               block1[ 3]  = append4_t[1];
 175               block1[ 4]  = append4_t[2];
 176               block1[ 5]  = append4_t[3];
 177               break;
 178
 179     case  3:  block0[ 3] |= append0_t[0];
 180               block0[ 4]  = append0_t[1];
 181               block0[ 5]  = append0_t[2];
 182               block0[ 6]  = append0_t[3];
 183
 184               block0[ 7]  = append1_t[0];
 185               block0[ 8]  = append1_t[1];
 186               block0[ 9]  = append1_t[2];
 187               block0[10]  = append1_t[3];
 188
 189               block0[11]  = append2_t[0];
 190               block0[12]  = append2_t[1];
 191               block0[13]  = append2_t[2];
 192               block0[14]  = append2_t[3];
 193
 194               block0[15]  = append3_t[0];
 195               block1[ 0]  = append3_t[1];
 196               block1[ 1]  = append3_t[2];
 197               block1[ 2]  = append3_t[3];
 198
 199               block1[ 3]  = append4_t[0];
 200               block1[ 4]  = append4_t[1];
 201               block1[ 5]  = append4_t[2];
 202               block1[ 6]  = append4_t[3];
 203               break;
 204
 205     case  4:  block0[ 4] |= append0_t[0];
 206               block0[ 5]  = append0_t[1];
 207               block0[ 6]  = append0_t[2];
 208               block0[ 7]  = append0_t[3];
 209
 210               block0[ 8]  = append1_t[0];
 211               block0[ 9]  = append1_t[1];
 212               block0[10]  = append1_t[2];
 213               block0[11]  = append1_t[3];
 214
 215               block0[12]  = append2_t[0];
 216               block0[13]  = append2_t[1];
 217               block0[14]  = append2_t[2];
 218               block0[15]  = append2_t[3];
 219
 220               block1[ 0]  = append3_t[0];
 221               block1[ 1]  = append3_t[1];
 222               block1[ 2]  = append3_t[2];
 223               block1[ 3]  = append3_t[3];
 224
 225               block1[ 4]  = append4_t[0];
 226               block1[ 5]  = append4_t[1];
 227               block1[ 6]  = append4_t[2];
 228               block1[ 7]  = append4_t[3];
 229               break;
 230
 231     case  5:  block0[ 5] |= append0_t[0];
 232               block0[ 6]  = append0_t[1];
 233               block0[ 7]  = append0_t[2];
 234               block0[ 8]  = append0_t[3];
 235
 236               block0[ 9]  = append1_t[0];
 237               block0[10]  = append1_t[1];
 238               block0[11]  = append1_t[2];
 239               block0[12]  = append1_t[3];
 240
 241               block0[13]  = append2_t[0];
 242               block0[14]  = append2_t[1];
 243               block0[15]  = append2_t[2];
 244               block1[ 0]  = append2_t[3];
 245
 246               block1[ 1]  = append3_t[0];
 247               block1[ 2]  = append3_t[1];
 248               block1[ 3]  = append3_t[2];
 249               block1[ 4]  = append3_t[3];
 250
 251               block1[ 5]  = append4_t[0];
 252               block1[ 6]  = append4_t[1];
 253               block1[ 7]  = append4_t[2];
 254               block1[ 8]  = append4_t[3];
 255               break;
 256
 257     case  6:  block0[ 6] |= append0_t[0];
 258               block0[ 7]  = append0_t[1];
 259               block0[ 8]  = append0_t[2];
 260               block0[ 9]  = append0_t[3];
 261
 262               block0[10]  = append1_t[0];
 263               block0[11]  = append1_t[1];
 264               block0[12]  = append1_t[2];
 265               block0[13]  = append1_t[3];
 266
 267               block0[14]  = append2_t[0];
 268               block0[15]  = append2_t[1];
 269               block1[ 0]  = append2_t[2];
 270               block1[ 1]  = append2_t[3];
 271
 272               block1[ 2]  = append3_t[0];
 273               block1[ 3]  = append3_t[1];
 274               block1[ 4]  = append3_t[2];
 275               block1[ 5]  = append3_t[3];
 276
 277               block1[ 6]  = append4_t[0];
 278               block1[ 7]  = append4_t[1];
 279               block1[ 8]  = append4_t[2];
 280               block1[ 9]  = append4_t[3];
 281               break;
 282
 283     case  7:  block0[ 7] |= append0_t[0];
 284               block0[ 8]  = append0_t[1];
 285               block0[ 9]  = append0_t[2];
 286               block0[10]  = append0_t[3];
 287
 288               block0[11]  = append1_t[0];
 289               block0[12]  = append1_t[1];
 290               block0[13]  = append1_t[2];
 291               block0[14]  = append1_t[3];
 292
 293               block0[15]  = append2_t[0];
 294               block1[ 0]  = append2_t[1];
 295               block1[ 1]  = append2_t[2];
 296               block1[ 2]  = append2_t[3];
 297
 298               block1[ 3]  = append3_t[0];
 299               block1[ 4]  = append3_t[1];
 300               block1[ 5]  = append3_t[2];
 301               block1[ 6]  = append3_t[3];
 302
 303               block1[ 7]  = append4_t[0];
 304               block1[ 8]  = append4_t[1];
 305               block1[ 9]  = append4_t[2];
 306               block1[10]  = append4_t[3];
 307               break;
 308
 309     case  8:  block0[ 8] |= append0_t[0];
 310               block0[ 9]  = append0_t[1];
 311               block0[10]  = append0_t[2];
 312               block0[11]  = append0_t[3];
 313
 314               block0[12]  = append1_t[0];
 315               block0[13]  = append1_t[1];
 316               block0[14]  = append1_t[2];
 317               block0[15]  = append1_t[3];
 318
 319               block1[ 0]  = append2_t[0];
 320               block1[ 1]  = append2_t[1];
 321               block1[ 2]  = append2_t[2];
 322               block1[ 3]  = append2_t[3];
 323
 324               block1[ 4]  = append3_t[0];
 325               block1[ 5]  = append3_t[1];
 326               block1[ 6]  = append3_t[2];
 327               block1[ 7]  = append3_t[3];
 328
 329               block1[ 8]  = append4_t[0];
 330               block1[ 9]  = append4_t[1];
 331               block1[10]  = append4_t[2];
 332               block1[11]  = append4_t[3];
 333               break;
 334
 335     case  9:  block0[ 9] |= append0_t[0];
 336               block0[10]  = append0_t[1];
 337               block0[11]  = append0_t[2];
 338               block0[12]  = append0_t[3];
 339
 340               block0[13]  = append1_t[0];
 341               block0[14]  = append1_t[1];
 342               block0[15]  = append1_t[2];
 343               block1[ 0]  = append1_t[3];
 344
 345               block1[ 1]  = append2_t[0];
 346               block1[ 2]  = append2_t[1];
 347               block1[ 3]  = append2_t[2];
 348               block1[ 4]  = append2_t[3];
 349
 350               block1[ 5]  = append3_t[0];
 351               block1[ 6]  = append3_t[1];
 352               block1[ 7]  = append3_t[2];
 353               block1[ 8]  = append3_t[3];
 354
 355               block1[ 9]  = append4_t[0];
 356               block1[10]  = append4_t[1];
 357               block1[11]  = append4_t[2];
 358               block1[12]  = append4_t[3];
 359               break;
 360
 361     case 10:  block0[10] |= append0_t[0];
 362               block0[11]  = append0_t[1];
 363               block0[12]  = append0_t[2];
 364               block0[13]  = append0_t[3];
 365
 366               block0[14]  = append1_t[0];
 367               block0[15]  = append1_t[1];
 368               block1[ 0]  = append1_t[2];
 369               block1[ 1]  = append1_t[3];
 370
 371               block1[ 2]  = append2_t[0];
 372               block1[ 3]  = append2_t[1];
 373               block1[ 4]  = append2_t[2];
 374               block1[ 5]  = append2_t[3];
 375
 376               block1[ 6]  = append3_t[0];
 377               block1[ 7]  = append3_t[1];
 378               block1[ 8]  = append3_t[2];
 379               block1[ 9]  = append3_t[3];
 380
 381               block1[10]  = append4_t[0];
 382               block1[11]  = append4_t[1];
 383               block1[12]  = append4_t[2];
 384               block1[13]  = append4_t[3];
 385               break;
 386
 387     case 11:  block0[11] |= append0_t[0];
 388               block0[12]  = append0_t[1];
 389               block0[13]  = append0_t[2];
 390               block0[14]  = append0_t[3];
 391
 392               block0[15]  = append1_t[0];
 393               block1[ 0]  = append1_t[1];
 394               block1[ 1]  = append1_t[2];
 395               block1[ 2]  = append1_t[3];
 396
 397               block1[ 3]  = append2_t[0];
 398               block1[ 4]  = append2_t[1];
 399               block1[ 5]  = append2_t[2];
 400               block1[ 6]  = append2_t[3];
 401
 402               block1[ 7]  = append3_t[0];
 403               block1[ 8]  = append3_t[1];
 404               block1[ 9]  = append3_t[2];
 405               block1[10]  = append3_t[3];
 406
 407               block1[11]  = append4_t[0];
 408               block1[12]  = append4_t[1];
 409               block1[13]  = append4_t[2];
 410               block1[14]  = append4_t[3];
 411               break;
 412
 413     case 12:  block0[12] |= append0_t[0];
 414               block0[13]  = append0_t[1];
 415               block0[14]  = append0_t[2];
 416               block0[15]  = append0_t[3];
 417
 418               block1[ 0]  = append1_t[0];
 419               block1[ 1]  = append1_t[1];
 420               block1[ 2]  = append1_t[2];
 421               block1[ 3]  = append1_t[3];
 422
 423               block1[ 4]  = append2_t[0];
 424               block1[ 5]  = append2_t[1];
 425               block1[ 6]  = append2_t[2];
 426               block1[ 7]  = append2_t[3];
 427
 428               block1[ 8]  = append3_t[0];
 429               block1[ 9]  = append3_t[1];
 430               block1[10]  = append3_t[2];
 431               block1[11]  = append3_t[3];
 432
 433               block1[12]  = append4_t[0];
 434               block1[13]  = append4_t[1];
 435               block1[14]  = append4_t[2];
 436               block1[15]  = append4_t[3];
 437               break;
 438
 439     case 13:  block0[13] |= append0_t[0];
 440               block0[14]  = append0_t[1];
 441               block0[15]  = append0_t[2];
 442               block1[ 0]  = append0_t[3];
 443
 444               block1[ 1]  = append1_t[0];
 445               block1[ 2]  = append1_t[1];
 446               block1[ 3]  = append1_t[2];
 447               block1[ 4]  = append1_t[3];
 448
 449               block1[ 5]  = append2_t[0];
 450               block1[ 6]  = append2_t[1];
 451               block1[ 7]  = append2_t[2];
 452               block1[ 8]  = append2_t[3];
 453
 454               block1[ 9]  = append3_t[0];
 455               block1[10]  = append3_t[1];
 456               block1[11]  = append3_t[2];
 457               block1[12]  = append3_t[3];
 458
 459               block1[13]  = append4_t[0];
 460               block1[14]  = append4_t[1];
 461               block1[15]  = append4_t[2];
 462               break;
 463
 464     case 14:  block0[14] |= append0_t[0];
 465               block0[15]  = append0_t[1];
 466               block1[ 0]  = append0_t[2];
 467               block1[ 1]  = append0_t[3];
 468
 469               block1[ 2]  = append1_t[0];
 470               block1[ 3]  = append1_t[1];
 471               block1[ 4]  = append1_t[2];
 472               block1[ 5]  = append1_t[3];
 473
 474               block1[ 6]  = append2_t[0];
 475               block1[ 7]  = append2_t[1];
 476               block1[ 8]  = append2_t[2];
 477               block1[ 9]  = append2_t[3];
 478
 479               block1[10]  = append3_t[0];
 480               block1[11]  = append3_t[1];
 481               block1[12]  = append3_t[2];
 482               block1[13]  = append3_t[3];
 483
 484               block1[14]  = append4_t[0];
 485               block1[15]  = append4_t[1];
 486               break;
 487
 488     case 15:  block0[15] |= append0_t[0];
 489               block1[ 0]  = append0_t[1];
 490               block1[ 1]  = append0_t[2];
 491               block1[ 2]  = append0_t[3];
 492
 493               block1[ 3]  = append1_t[1];
 494               block1[ 4]  = append1_t[2];
 495               block1[ 5]  = append1_t[3];
 496               block1[ 6]  = append1_t[0];
 497
 498               block1[ 7]  = append2_t[0];
 499               block1[ 8]  = append2_t[1];
 500               block1[ 9]  = append2_t[2];
 501               block1[10]  = append2_t[3];
 502
 503               block1[11]  = append3_t[0];
 504               block1[12]  = append3_t[1];
 505               block1[13]  = append3_t[2];
 506               block1[14]  = append3_t[3];
 507
 508               block1[15]  = append4_t[0];
 509               break;
 510
 511     case 16:  block1[ 0] |= append0_t[0];
 512               block1[ 1]  = append0_t[1];
 513               block1[ 2]  = append0_t[2];
 514               block1[ 3]  = append0_t[3];
 515
 516               block1[ 4]  = append1_t[0];
 517               block1[ 5]  = append1_t[1];
 518               block1[ 6]  = append1_t[2];
 519               block1[ 7]  = append1_t[3];
 520
 521               block1[ 8]  = append2_t[0];
 522               block1[ 9]  = append2_t[1];
 523               block1[10]  = append2_t[2];
 524               block1[11]  = append2_t[3];
 525
 526               block1[12]  = append3_t[0];
 527               block1[13]  = append3_t[1];
 528               block1[14]  = append3_t[2];
 529               block1[15]  = append3_t[3];
 530               break;
 531
 532     case 17:  block1[ 1] |= append0_t[0];
 533               block1[ 2]  = append0_t[1];
 534               block1[ 3]  = append0_t[2];
 535               block1[ 4]  = append0_t[3];
 536
 537               block1[ 5]  = append1_t[0];
 538               block1[ 6]  = append1_t[1];
 539               block1[ 7]  = append1_t[2];
 540               block1[ 8]  = append1_t[3];
 541
 542               block1[ 9]  = append2_t[0];
 543               block1[10]  = append2_t[1];
 544               block1[11]  = append2_t[2];
 545               block1[12]  = append2_t[3];
 546
 547               block1[13]  = append3_t[0];
 548               block1[14]  = append3_t[1];
 549               block1[15]  = append3_t[2];
 550               break;
 551
 552     case 18:  block1[ 2] |= append0_t[0];
 553               block1[ 3]  = append0_t[1];
 554               block1[ 4]  = append0_t[2];
 555               block1[ 5]  = append0_t[3];
 556
 557               block1[ 6]  = append1_t[0];
 558               block1[ 7]  = append1_t[1];
 559               block1[ 8]  = append1_t[2];
 560               block1[ 9]  = append1_t[3];
 561
 562               block1[10]  = append2_t[0];
 563               block1[11]  = append2_t[1];
 564               block1[12]  = append2_t[2];
 565               block1[13]  = append2_t[3];
 566
 567               block1[14]  = append3_t[0];
 568               block1[15]  = append3_t[1];
 569               break;
 570
 571     case 19:  block1[ 3] |= append0_t[0];
 572               block1[ 4]  = append0_t[1];
 573               block1[ 5]  = append0_t[2];
 574               block1[ 6]  = append0_t[3];
 575
 576               block1[ 7]  = append1_t[0];
 577               block1[ 8]  = append1_t[1];
 578               block1[ 9]  = append1_t[2];
 579               block1[10]  = append1_t[3];
 580
 581               block1[11]  = append2_t[0];
 582               block1[12]  = append2_t[1];
 583               block1[13]  = append2_t[2];
 584               block1[14]  = append2_t[3];
 585
 586               block1[15]  = append3_t[0];
 587               break;
 588
 589     case 20:  block1[ 4] |= append0_t[0];
 590               block1[ 5]  = append0_t[1];
 591               block1[ 6]  = append0_t[2];
 592               block1[ 7]  = append0_t[3];
 593
 594               block1[ 8]  = append1_t[0];
 595               block1[ 9]  = append1_t[1];
 596               block1[10]  = append1_t[2];
 597               block1[11]  = append1_t[3];
 598
 599               block1[12]  = append2_t[0];
 600               block1[13]  = append2_t[1];
 601               block1[14]  = append2_t[2];
 602               block1[15]  = append2_t[3];
 603               break;
 604
 605     case 21:  block1[ 5] |= append0_t[0];
 606               block1[ 6]  = append0_t[1];
 607               block1[ 7]  = append0_t[2];
 608               block1[ 8]  = append0_t[3];
 609
 610               block1[ 9]  = append1_t[0];
 611               block1[10]  = append1_t[1];
 612               block1[11]  = append1_t[2];
 613               block1[12]  = append1_t[3];
 614
 615               block1[13]  = append2_t[0];
 616               block1[14]  = append2_t[1];
 617               block1[15]  = append2_t[2];
 618               break;
 619
 620     case 22:  block1[ 6] |= append0_t[0];
 621               block1[ 7]  = append0_t[1];
 622               block1[ 8]  = append0_t[2];
 623               block1[ 9]  = append0_t[3];
 624
 625               block1[10]  = append1_t[0];
 626               block1[11]  = append1_t[1];
 627               block1[12]  = append1_t[2];
 628               block1[13]  = append1_t[3];
 629
 630               block1[14]  = append2_t[0];
 631               block1[15]  = append2_t[1];
 632               break;
 633
 634     case 23:  block1[ 7] |= append0_t[0];
 635               block1[ 8]  = append0_t[1];
 636               block1[ 9]  = append0_t[2];
 637               block1[10]  = append0_t[3];
 638
 639               block1[11]  = append1_t[0];
 640               block1[12]  = append1_t[1];
 641               block1[13]  = append1_t[2];
 642               block1[14]  = append1_t[3];
 643
 644               block1[15]  = append2_t[0];
 645               break;
 646
 647     case 24:  block1[ 8] |= append0_t[0];
 648               block1[ 9]  = append0_t[1];
 649               block1[10]  = append0_t[2];
 650               block1[11]  = append0_t[3];
 651
 652               block1[12]  = append1_t[0];
 653               block1[13]  = append1_t[1];
 654               block1[14]  = append1_t[2];
 655               block1[15]  = append1_t[3];
 656               break;
 657
 658     case 25:  block1[ 9] |= append0_t[0];
 659               block1[10]  = append0_t[1];
 660               block1[11]  = append0_t[2];
 661               block1[12]  = append0_t[3];
 662
 663               block1[13]  = append1_t[0];
 664               block1[14]  = append1_t[1];
 665               block1[15]  = append1_t[2];
 666               break;
 667
 668     case 26:  block1[10] |= append0_t[0];
 669               block1[11]  = append0_t[1];
 670               block1[12]  = append0_t[2];
 671               block1[13]  = append0_t[3];
 672
 673               block1[14]  = append1_t[0];
 674               block1[15]  = append1_t[1];
 675               break;
 676
 677     case 27:  block1[11] |= append0_t[0];
 678               block1[12]  = append0_t[1];
 679               block1[13]  = append0_t[2];
 680               block1[14]  = append0_t[3];
 681
 682               block1[15]  = append1_t[0];
 683               break;
 684
 685     case 28:  block1[12] |= append0_t[0];
 686               block1[13]  = append0_t[1];
 687               block1[14]  = append0_t[2];
 688               block1[15]  = append0_t[3];
 689               break;
 690
 691     case 29:  block1[13] |= append0_t[0];
 692               block1[14]  = append0_t[1];
 693               block1[15]  = append0_t[2];
 694               break;
 695
 696     case 30:  block1[14] |= append0_t[0];
 697               block1[15]  = append0_t[1];
 698               break;
 699   }
 700
 701   u32 new_len = block_len + append_len;
 702
 703   return new_len;
 704 }
 705
 706 __device__ __constant__ char c_bin2asc[16] = { '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', 'a', 'b', 'c', 'd', 'e', 'f' };
 707
 708 __device__ __shared__ short l_bin2asc[256];
 709
 710 __device__ __constant__ comb_t c_combs[1024];
 711
 712 extern "C" __global__ void __launch_bounds__ (256, 1) m11400_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 sip_t *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)
 713 {
 714   /**
 715    * modifier
 716    */
 717
 718   const u32 lid = threadIdx.x;
 719
 720   /**
 721    * base
 722    */
 723
 724   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 725
 726   u32x wordl0[4];
 727
 728   wordl0[0] = pws[gid].i[ 0];
 729   wordl0[1] = pws[gid].i[ 1];
 730   wordl0[2] = pws[gid].i[ 2];
 731   wordl0[3] = pws[gid].i[ 3];
 732
 733   u32x wordl1[4];
 734
 735   wordl1[0] = pws[gid].i[ 4];
 736   wordl1[1] = pws[gid].i[ 5];
 737   wordl1[2] = pws[gid].i[ 6];
 738   wordl1[3] = pws[gid].i[ 7];
 739
 740   u32x wordl2[4];
 741
 742   wordl2[0] = 0;
 743   wordl2[1] = 0;
 744   wordl2[2] = 0;
 745   wordl2[3] = 0;
 746
 747   u32x wordl3[4];
 748
 749   wordl3[0] = 0;
 750   wordl3[1] = 0;
 751   wordl3[2] = 0;
 752   wordl3[3] = 0;
 753
 754   /**
 755    * bin2asc table
 756    */
 757
 758   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
 759                  | c_bin2asc[(lid >> 4) & 15] << 0;
 760
 761   __syncthreads ();
 762
 763   if (gid >= gid_max) return;
 764
 765   const u32 pw_l_len = pws[gid].pw_len;
 766
 767   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 768   {
 769     append_0x80_2 (wordl0, wordl1, pw_l_len);
 770
 771     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 772   }
 773
 774   /**
 775    * salt
 776    */
 777
 778   const u32 salt_len = esalt_bufs[salt_pos].salt_len; // not a bug, we need to get it from the esalt
 779
 780   u32 salt_buf0[16];
 781
 782   salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0];
 783   salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1];
 784   salt_buf0[ 2] = esalt_bufs[salt_pos].salt_buf[ 2];
 785   salt_buf0[ 3] = esalt_bufs[salt_pos].salt_buf[ 3];
 786   salt_buf0[ 4] = esalt_bufs[salt_pos].salt_buf[ 4];
 787   salt_buf0[ 5] = esalt_bufs[salt_pos].salt_buf[ 5];
 788   salt_buf0[ 6] = esalt_bufs[salt_pos].salt_buf[ 6];
 789   salt_buf0[ 7] = esalt_bufs[salt_pos].salt_buf[ 7];
 790   salt_buf0[ 8] = esalt_bufs[salt_pos].salt_buf[ 8];
 791   salt_buf0[ 9] = esalt_bufs[salt_pos].salt_buf[ 9];
 792   salt_buf0[10] = esalt_bufs[salt_pos].salt_buf[10];
 793   salt_buf0[11] = esalt_bufs[salt_pos].salt_buf[11];
 794   salt_buf0[12] = esalt_bufs[salt_pos].salt_buf[12];
 795   salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13];
 796   salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14];
 797   salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15];
 798
 799   u32 salt_buf1[16];
 800
 801   salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16];
 802   salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17];
 803   salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18];
 804   salt_buf1[ 3] = esalt_bufs[salt_pos].salt_buf[19];
 805   salt_buf1[ 4] = esalt_bufs[salt_pos].salt_buf[20];
 806   salt_buf1[ 5] = esalt_bufs[salt_pos].salt_buf[21];
 807   salt_buf1[ 6] = esalt_bufs[salt_pos].salt_buf[22];
 808   salt_buf1[ 7] = esalt_bufs[salt_pos].salt_buf[23];
 809   salt_buf1[ 8] = esalt_bufs[salt_pos].salt_buf[24];
 810   salt_buf1[ 9] = esalt_bufs[salt_pos].salt_buf[25];
 811   salt_buf1[10] = esalt_bufs[salt_pos].salt_buf[26];
 812   salt_buf1[11] = esalt_bufs[salt_pos].salt_buf[27];
 813   salt_buf1[12] = esalt_bufs[salt_pos].salt_buf[28];
 814   salt_buf1[13] = esalt_bufs[salt_pos].salt_buf[29];
 815   salt_buf1[14] = 0;
 816   salt_buf1[15] = 0;
 817
 818   /**
 819    * esalt
 820    */
 821
 822   const u32 esalt_len = esalt_bufs[salt_pos].esalt_len;
 823
 824   u32 esalt_buf0[16];
 825
 826   esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0];
 827   esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1];
 828   esalt_buf0[ 2] = esalt_bufs[salt_pos].esalt_buf[ 2];
 829   esalt_buf0[ 3] = esalt_bufs[salt_pos].esalt_buf[ 3];
 830   esalt_buf0[ 4] = esalt_bufs[salt_pos].esalt_buf[ 4];
 831   esalt_buf0[ 5] = esalt_bufs[salt_pos].esalt_buf[ 5];
 832   esalt_buf0[ 6] = esalt_bufs[salt_pos].esalt_buf[ 6];
 833   esalt_buf0[ 7] = esalt_bufs[salt_pos].esalt_buf[ 7];
 834   esalt_buf0[ 8] = esalt_bufs[salt_pos].esalt_buf[ 8];
 835   esalt_buf0[ 9] = esalt_bufs[salt_pos].esalt_buf[ 9];
 836   esalt_buf0[10] = esalt_bufs[salt_pos].esalt_buf[10];
 837   esalt_buf0[11] = esalt_bufs[salt_pos].esalt_buf[11];
 838   esalt_buf0[12] = esalt_bufs[salt_pos].esalt_buf[12];
 839   esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13];
 840   esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14];
 841   esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15];
 842
 843   u32 esalt_buf1[16];
 844
 845   esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16];
 846   esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17];
 847   esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18];
 848   esalt_buf1[ 3] = esalt_bufs[salt_pos].esalt_buf[19];
 849   esalt_buf1[ 4] = esalt_bufs[salt_pos].esalt_buf[20];
 850   esalt_buf1[ 5] = esalt_bufs[salt_pos].esalt_buf[21];
 851   esalt_buf1[ 6] = esalt_bufs[salt_pos].esalt_buf[22];
 852   esalt_buf1[ 7] = esalt_bufs[salt_pos].esalt_buf[23];
 853   esalt_buf1[ 8] = esalt_bufs[salt_pos].esalt_buf[24];
 854   esalt_buf1[ 9] = esalt_bufs[salt_pos].esalt_buf[25];
 855   esalt_buf1[10] = esalt_bufs[salt_pos].esalt_buf[26];
 856   esalt_buf1[11] = esalt_bufs[salt_pos].esalt_buf[27];
 857   esalt_buf1[12] = esalt_bufs[salt_pos].esalt_buf[28];
 858   esalt_buf1[13] = esalt_bufs[salt_pos].esalt_buf[29];
 859   esalt_buf1[14] = esalt_bufs[salt_pos].esalt_buf[30];
 860   esalt_buf1[15] = esalt_bufs[salt_pos].esalt_buf[31];
 861
 862   u32 esalt_buf2[16];
 863
 864   esalt_buf2[ 0] = esalt_bufs[salt_pos].esalt_buf[32];
 865   esalt_buf2[ 1] = esalt_bufs[salt_pos].esalt_buf[33];
 866   esalt_buf2[ 2] = esalt_bufs[salt_pos].esalt_buf[34];
 867   esalt_buf2[ 3] = esalt_bufs[salt_pos].esalt_buf[35];
 868   esalt_buf2[ 4] = esalt_bufs[salt_pos].esalt_buf[36];
 869   esalt_buf2[ 5] = esalt_bufs[salt_pos].esalt_buf[37];
 870   esalt_buf2[ 6] = 0;
 871   esalt_buf2[ 7] = 0;
 872   esalt_buf2[ 8] = 0;
 873   esalt_buf2[ 9] = 0;
 874   esalt_buf2[10] = 0;
 875   esalt_buf2[11] = 0;
 876   esalt_buf2[12] = 0;
 877   esalt_buf2[13] = 0;
 878   esalt_buf2[14] = 0;
 879   esalt_buf2[15] = 0;
 880
 881   const u32 digest_esalt_len = 32 + esalt_len;
 882   const u32 remaining_bytes  = digest_esalt_len + 1 - 64; // substract previous block
 883
 884   /**
 885    * loop
 886    */
 887
 888   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 889   {
 890     const u32 pw_r_len = c_combs[il_pos].pw_len;
 891
 892     const u32 pw_len = pw_l_len + pw_r_len;
 893
 894     u32 wordr0[4];
 895
 896     wordr0[0] = c_combs[il_pos].i[0];
 897     wordr0[1] = c_combs[il_pos].i[1];
 898     wordr0[2] = c_combs[il_pos].i[2];
 899     wordr0[3] = c_combs[il_pos].i[3];
 900
 901     u32 wordr1[4];
 902
 903     wordr1[0] = c_combs[il_pos].i[4];
 904     wordr1[1] = c_combs[il_pos].i[5];
 905     wordr1[2] = c_combs[il_pos].i[6];
 906     wordr1[3] = c_combs[il_pos].i[7];
 907
 908     u32 wordr2[4];
 909
 910     wordr2[0] = 0;
 911     wordr2[1] = 0;
 912     wordr2[2] = 0;
 913     wordr2[3] = 0;
 914
 915     u32 wordr3[4];
 916
 917     wordr3[0] = 0;
 918     wordr3[1] = 0;
 919     wordr3[2] = 0;
 920     wordr3[3] = 0;
 921
 922     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 923     {
 924       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 925     }
 926
 927     u32x w0[4];
 928
 929     w0[0] = wordl0[0] | wordr0[0];
 930     w0[1] = wordl0[1] | wordr0[1];
 931     w0[2] = wordl0[2] | wordr0[2];
 932     w0[3] = wordl0[3] | wordr0[3];
 933
 934     u32x w1[4];
 935
 936     w1[0] = wordl1[0] | wordr1[0];
 937     w1[1] = wordl1[1] | wordr1[1];
 938     w1[2] = wordl1[2] | wordr1[2];
 939     w1[3] = wordl1[3] | wordr1[3];
 940
 941     u32x w2[4];
 942
 943     w2[0] = wordl2[0] | wordr2[0];
 944     w2[1] = wordl2[1] | wordr2[1];
 945     w2[2] = wordl2[2] | wordr2[2];
 946     w2[3] = wordl2[3] | wordr2[3];
 947
 948     u32x w3[4];
 949
 950     w3[0] = wordl3[0] | wordr3[0];
 951     w3[1] = wordl3[1] | wordr3[1];
 952     w3[2] = wordl3[2] | wordr3[2];
 953     w3[3] = wordl3[3] | wordr3[3];
 954
 955     const u32 pw_salt_len = salt_len + pw_len;
 956
 957     /*
 958      * HA1 = md5 ($salt . $pass)
 959      */
 960
 961     // append the pass to the salt
 962
 963     u32x block0[16];
 964
 965     block0[ 0] = salt_buf0[ 0];
 966     block0[ 1] = salt_buf0[ 1];
 967     block0[ 2] = salt_buf0[ 2];
 968     block0[ 3] = salt_buf0[ 3];
 969     block0[ 4] = salt_buf0[ 4];
 970     block0[ 5] = salt_buf0[ 5];
 971     block0[ 6] = salt_buf0[ 6];
 972     block0[ 7] = salt_buf0[ 7];
 973     block0[ 8] = salt_buf0[ 8];
 974     block0[ 9] = salt_buf0[ 9];
 975     block0[10] = salt_buf0[10];
 976     block0[11] = salt_buf0[11];
 977     block0[12] = salt_buf0[12];
 978     block0[13] = salt_buf0[13];
 979     block0[14] = salt_buf0[14];
 980     block0[15] = salt_buf0[15];
 981
 982     u32x block1[16];
 983
 984     block1[ 0] = salt_buf1[ 0];
 985     block1[ 1] = salt_buf1[ 1];
 986     block1[ 2] = salt_buf1[ 2];
 987     block1[ 3] = salt_buf1[ 3];
 988     block1[ 4] = salt_buf1[ 4];
 989     block1[ 5] = salt_buf1[ 5];
 990     block1[ 6] = salt_buf1[ 6];
 991     block1[ 7] = salt_buf1[ 7];
 992     block1[ 8] = salt_buf1[ 8];
 993     block1[ 9] = salt_buf1[ 9];
 994     block1[10] = salt_buf1[10];
 995     block1[11] = salt_buf1[11];
 996     block1[12] = salt_buf1[12];
 997     block1[13] = salt_buf1[13];
 998     block1[14] = salt_buf1[14];
 999     block1[15] = salt_buf1[15];
1000
1001     u32 block_len = 0;
1002
1003     block_len = memcat32 (block0, block1, salt_len, w0, w1, w2, w3, pw_len);
1004
1005     u32x w0_t[4];
1006
1007     w0_t[0] = block0[ 0];
1008     w0_t[1] = block0[ 1];
1009     w0_t[2] = block0[ 2];
1010     w0_t[3] = block0[ 3];
1011
1012     u32x w1_t[4];
1013
1014     w1_t[0] = block0[ 4];
1015     w1_t[1] = block0[ 5];
1016     w1_t[2] = block0[ 6];
1017     w1_t[3] = block0[ 7];
1018
1019     u32x w2_t[4];
1020
1021     w2_t[0] = block0[ 8];
1022     w2_t[1] = block0[ 9];
1023     w2_t[2] = block0[10];
1024     w2_t[3] = block0[11];
1025
1026     u32x w3_t[4];
1027
1028     w3_t[0] = block0[12];
1029     w3_t[1] = block0[13];
1030     w3_t[2] = block0[14];
1031     w3_t[3] = block0[15];
1032
1033     if (block_len < 56)
1034     {
1035       w3_t[2] = pw_salt_len * 8;
1036     }
1037
1038     // md5
1039
1040     u32x tmp2;
1041
1042     u32x a = MD5M_A;
1043     u32x b = MD5M_B;
1044     u32x c = MD5M_C;
1045     u32x d = MD5M_D;
1046
1047     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1048     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1049     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1050     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1051     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1052     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1053     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1054     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1055     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1056     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1057     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1058     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1059     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1060     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1061     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1062     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1063
1064     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1065     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1066     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1067     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1068     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1069     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1070     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1071     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1072     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1073     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1074     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1075     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1076     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1077     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1078     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1079     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1080
1081     MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1082     MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1083     MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1084     MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1085     MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1086     MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1087     MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1088     MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1089     MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1090     MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1091     MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1092     MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1093     MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1094     MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1095     MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1096     MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1097
1098     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1099     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1100     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1101     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1102     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1103     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1104     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1105     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1106     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1107     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1108     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1109     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1110     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1111     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1112     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1113     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1114
1115     a += MD5M_A;
1116     b += MD5M_B;
1117     c += MD5M_C;
1118     d += MD5M_D;
1119
1120     if (block_len > 55)
1121     {
1122       u32x r_a = a;
1123       u32x r_b = b;
1124       u32x r_c = c;
1125       u32x r_d = d;
1126
1127       w0_t[0] = block1[ 0];
1128       w0_t[1] = block1[ 1];
1129       w0_t[2] = block1[ 2];
1130       w0_t[3] = block1[ 3];
1131
1132       w1_t[0] = block1[ 4];
1133       w1_t[1] = block1[ 5];
1134       w1_t[2] = block1[ 6];
1135       w1_t[3] = block1[ 7];
1136
1137       w2_t[0] = block1[ 8];
1138       w2_t[1] = block1[ 9];
1139       w2_t[2] = block1[10];
1140       w2_t[3] = block1[11];
1141
1142       w3_t[0] = block1[12];
1143       w3_t[1] = block1[13];
1144       w3_t[2] = pw_salt_len * 8;
1145       w3_t[3] = 0;
1146
1147       MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1148       MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1149       MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1150       MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1151       MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1152       MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1153       MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1154       MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1155       MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1156       MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1157       MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1158       MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1159       MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1160       MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1161       MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1162       MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1163
1164       MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1165       MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1166       MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1167       MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1168       MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1169       MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1170       MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1171       MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1172       MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1173       MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1174       MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1175       MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1176       MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1177       MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1178       MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1179       MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1180
1181       MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1182       MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1183       MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1184       MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1185       MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1186       MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1187       MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1188       MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1189       MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1190       MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1191       MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1192       MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1193       MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1194       MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1195       MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1196       MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1197
1198       MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1199       MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1200       MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1201       MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1202       MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1203       MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1204       MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1205       MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1206       MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1207       MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1208       MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1209       MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1210       MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1211       MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1212       MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1213       MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1214
1215       a += r_a;
1216       b += r_b;
1217       c += r_c;
1218       d += r_d;
1219     }
1220
1221     /*
1222      * final = md5 ($HA1 . $esalt)
1223      * we have at least 2 MD5 blocks/transformations, but we might need 3
1224      */
1225
1226     w0_t[0] = uint_to_hex_lower8 ((a >>  0) & 255) <<  0
1227             | uint_to_hex_lower8 ((a >>  8) & 255) << 16;
1228     w0_t[1] = uint_to_hex_lower8 ((a >> 16) & 255) <<  0
1229             | uint_to_hex_lower8 ((a >> 24) & 255) << 16;
1230     w0_t[2] = uint_to_hex_lower8 ((b >>  0) & 255) <<  0
1231             | uint_to_hex_lower8 ((b >>  8) & 255) << 16;
1232     w0_t[3] = uint_to_hex_lower8 ((b >> 16) & 255) <<  0
1233             | uint_to_hex_lower8 ((b >> 24) & 255) << 16;
1234     w1_t[0] = uint_to_hex_lower8 ((c >>  0) & 255) <<  0
1235             | uint_to_hex_lower8 ((c >>  8) & 255) << 16;
1236     w1_t[1] = uint_to_hex_lower8 ((c >> 16) & 255) <<  0
1237             | uint_to_hex_lower8 ((c >> 24) & 255) << 16;
1238     w1_t[2] = uint_to_hex_lower8 ((d >>  0) & 255) <<  0
1239             | uint_to_hex_lower8 ((d >>  8) & 255) << 16;
1240     w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) <<  0
1241             | uint_to_hex_lower8 ((d >> 24) & 255) << 16;
1242
1243     w2_t[0] = esalt_buf0[0];
1244     w2_t[1] = esalt_buf0[1];
1245     w2_t[2] = esalt_buf0[2];
1246     w2_t[3] = esalt_buf0[3];
1247
1248     w3_t[0] = esalt_buf0[4];
1249     w3_t[1] = esalt_buf0[5];
1250     w3_t[2] = esalt_buf0[6];
1251     w3_t[3] = esalt_buf0[7];
1252
1253     // md5
1254     // 1st transform
1255
1256     a = MD5M_A;
1257     b = MD5M_B;
1258     c = MD5M_C;
1259     d = MD5M_D;
1260
1261     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1262     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1263     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1264     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1265     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1266     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1267     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1268     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1269     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1270     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1271     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1272     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1273     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1274     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1275     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1276     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1277
1278     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1279     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1280     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1281     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1282     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1283     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1284     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1285     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1286     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1287     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1288     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1289     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1290     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1291     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1292     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1293     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1294
1295     MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1296     MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1297     MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1298     MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1299     MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1300     MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1301     MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1302     MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1303     MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1304     MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1305     MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1306     MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1307     MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1308     MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1309     MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1310     MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1311
1312     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1313     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1314     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1315     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1316     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1317     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1318     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1319     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1320     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1321     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1322     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1323     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1324     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1325     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1326     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1327     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1328
1329     a += MD5M_A;
1330     b += MD5M_B;
1331     c += MD5M_C;
1332     d += MD5M_D;
1333
1334     u32x r_a = a;
1335     u32x r_b = b;
1336     u32x r_c = c;
1337     u32x r_d = d;
1338
1339     // 2nd transform
1340
1341     w0_t[0] = esalt_buf0[ 8];
1342     w0_t[1] = esalt_buf0[ 9];
1343     w0_t[2] = esalt_buf0[10];
1344     w0_t[3] = esalt_buf0[11];
1345
1346     w1_t[0] = esalt_buf0[12];
1347     w1_t[1] = esalt_buf0[13];
1348     w1_t[2] = esalt_buf0[14];
1349     w1_t[3] = esalt_buf0[15];
1350
1351     w2_t[0] = esalt_buf1[ 0];
1352     w2_t[1] = esalt_buf1[ 1];
1353     w2_t[2] = esalt_buf1[ 2];
1354     w2_t[3] = esalt_buf1[ 3];
1355
1356     w3_t[0] = esalt_buf1[ 4];
1357     w3_t[1] = esalt_buf1[ 5];
1358     w3_t[2] = esalt_buf1[ 6];
1359     w3_t[3] = esalt_buf1[ 7];
1360
1361     // it is the final block when no more than 55 bytes left
1362
1363     if (remaining_bytes < 56)
1364     {
1365       // it is the last block !
1366
1367       w3_t[2] = digest_esalt_len * 8;
1368     }
1369
1370     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1371     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1372     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1373     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1374     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1375     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1376     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1377     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1378     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1379     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1380     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1381     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1382     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1383     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1384     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1385     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1386
1387     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1388     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1389     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1390     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1391     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1392     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1393     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1394     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1395     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1396     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1397     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1398     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1399     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1400     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1401     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1402     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1403
1404     MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1405     MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1406     MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1407     MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1408     MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1409     MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1410     MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1411     MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1412     MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1413     MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1414     MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1415     MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1416     MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1417     MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1418     MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1419     MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1420
1421     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1422     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1423     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1424     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1425     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1426     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1427     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1428     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1429     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1430     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1431     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1432     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1433     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1434     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1435     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1436     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1437
1438     // sometimes (not rare at all) we need a third block :(
1439
1440     if (remaining_bytes > 55)
1441     {
1442       // this is for sure the final block
1443
1444       a += r_a;
1445       b += r_b;
1446       c += r_c;
1447       d += r_d;
1448
1449       r_a = a;
1450       r_b = b;
1451       r_c = c;
1452       r_d = d;
1453
1454       w0_t[0] = esalt_buf1[ 8];
1455       w0_t[1] = esalt_buf1[ 9];
1456       w0_t[2] = esalt_buf1[10];
1457       w0_t[3] = esalt_buf1[11];
1458
1459       w1_t[0] = esalt_buf1[12];
1460       w1_t[1] = esalt_buf1[13];
1461       w1_t[2] = esalt_buf1[14];
1462       w1_t[3] = esalt_buf1[15];
1463
1464       w2_t[0] = esalt_buf2[ 0];
1465       w2_t[1] = esalt_buf2[ 1];
1466       w2_t[2] = esalt_buf2[ 2];
1467       w2_t[3] = esalt_buf2[ 3];
1468
1469       w3_t[0] = esalt_buf2[ 4];
1470       w3_t[1] = esalt_buf2[ 5];
1471       w3_t[2] = digest_esalt_len * 8;
1472       w3_t[3] = 0;
1473
1474       MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1475       MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1476       MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1477       MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1478       MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1479       MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1480       MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1481       MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1482       MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1483       MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1484       MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1485       MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1486       MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1487       MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1488       MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1489       MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1490
1491       MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1492       MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1493       MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1494       MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1495       MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1496       MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1497       MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1498       MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1499       MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1500       MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1501       MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1502       MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1503       MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1504       MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1505       MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1506       MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1507
1508       MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1509       MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1510       MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1511       MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1512       MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1513       MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1514       MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1515       MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1516       MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1517       MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1518       MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1519       MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1520       MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1521       MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1522       MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1523       MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1524
1525       MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1526       MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1527       MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1528       MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1529       MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1530       MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1531       MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1532       MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1533       MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1534       MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1535       MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1536       MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1537       MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1538       MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1539       MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1540       MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1541     }
1542
1543     a += r_a;
1544     b += r_b;
1545     c += r_c;
1546     d += r_d;
1547
1548     const u32x r0 = a;
1549     const u32x r1 = d;
1550     const u32x r2 = c;
1551     const u32x r3 = b;
1552
1553     #include VECT_COMPARE_M
1554   }
1555 }
1556
1557 extern "C" __global__ void __launch_bounds__ (256, 1) m11400_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 sip_t *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)
1558 {
1559 }
1560
1561 extern "C" __global__ void __launch_bounds__ (256, 1) m11400_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 sip_t *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)
1562 {
1563 }
1564
1565 extern "C" __global__ void __launch_bounds__ (256, 1) m11400_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 sip_t *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)
1566 {
1567   /**
1568    * modifier
1569    */
1570
1571   const u32 lid = threadIdx.x;
1572
1573   /**
1574    * base
1575    */
1576
1577   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1578
1579   u32x wordl0[4];
1580
1581   wordl0[0] = pws[gid].i[ 0];
1582   wordl0[1] = pws[gid].i[ 1];
1583   wordl0[2] = pws[gid].i[ 2];
1584   wordl0[3] = pws[gid].i[ 3];
1585
1586   u32x wordl1[4];
1587
1588   wordl1[0] = pws[gid].i[ 4];
1589   wordl1[1] = pws[gid].i[ 5];
1590   wordl1[2] = pws[gid].i[ 6];
1591   wordl1[3] = pws[gid].i[ 7];
1592
1593   u32x wordl2[4];
1594
1595   wordl2[0] = 0;
1596   wordl2[1] = 0;
1597   wordl2[2] = 0;
1598   wordl2[3] = 0;
1599
1600   u32x wordl3[4];
1601
1602   wordl3[0] = 0;
1603   wordl3[1] = 0;
1604   wordl3[2] = 0;
1605   wordl3[3] = 0;
1606
1607   /**
1608    * bin2asc table
1609    */
1610
1611   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
1612                  | c_bin2asc[(lid >> 4) & 15] << 0;
1613
1614   __syncthreads ();
1615
1616   if (gid >= gid_max) return;
1617
1618   const u32 pw_l_len = pws[gid].pw_len;
1619
1620   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
1621   {
1622     append_0x80_2 (wordl0, wordl1, pw_l_len);
1623
1624     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
1625   }
1626
1627   /**
1628    * salt
1629    */
1630
1631   const u32 salt_len = esalt_bufs[salt_pos].salt_len; // not a bug, we need to get it from the esalt
1632
1633   u32 salt_buf0[16];
1634
1635   salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0];
1636   salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1];
1637   salt_buf0[ 2] = esalt_bufs[salt_pos].salt_buf[ 2];
1638   salt_buf0[ 3] = esalt_bufs[salt_pos].salt_buf[ 3];
1639   salt_buf0[ 4] = esalt_bufs[salt_pos].salt_buf[ 4];
1640   salt_buf0[ 5] = esalt_bufs[salt_pos].salt_buf[ 5];
1641   salt_buf0[ 6] = esalt_bufs[salt_pos].salt_buf[ 6];
1642   salt_buf0[ 7] = esalt_bufs[salt_pos].salt_buf[ 7];
1643   salt_buf0[ 8] = esalt_bufs[salt_pos].salt_buf[ 8];
1644   salt_buf0[ 9] = esalt_bufs[salt_pos].salt_buf[ 9];
1645   salt_buf0[10] = esalt_bufs[salt_pos].salt_buf[10];
1646   salt_buf0[11] = esalt_bufs[salt_pos].salt_buf[11];
1647   salt_buf0[12] = esalt_bufs[salt_pos].salt_buf[12];
1648   salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13];
1649   salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14];
1650   salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15];
1651
1652   u32 salt_buf1[16];
1653
1654   salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16];
1655   salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17];
1656   salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18];
1657   salt_buf1[ 3] = esalt_bufs[salt_pos].salt_buf[19];
1658   salt_buf1[ 4] = esalt_bufs[salt_pos].salt_buf[20];
1659   salt_buf1[ 5] = esalt_bufs[salt_pos].salt_buf[21];
1660   salt_buf1[ 6] = esalt_bufs[salt_pos].salt_buf[22];
1661   salt_buf1[ 7] = esalt_bufs[salt_pos].salt_buf[23];
1662   salt_buf1[ 8] = esalt_bufs[salt_pos].salt_buf[24];
1663   salt_buf1[ 9] = esalt_bufs[salt_pos].salt_buf[25];
1664   salt_buf1[10] = esalt_bufs[salt_pos].salt_buf[26];
1665   salt_buf1[11] = esalt_bufs[salt_pos].salt_buf[27];
1666   salt_buf1[12] = esalt_bufs[salt_pos].salt_buf[28];
1667   salt_buf1[13] = esalt_bufs[salt_pos].salt_buf[29];
1668   salt_buf1[14] = 0;
1669   salt_buf1[15] = 0;
1670
1671   /**
1672    * esalt
1673    */
1674
1675   const u32 esalt_len = esalt_bufs[salt_pos].esalt_len;
1676
1677   u32 esalt_buf0[16];
1678
1679   esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0];
1680   esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1];
1681   esalt_buf0[ 2] = esalt_bufs[salt_pos].esalt_buf[ 2];
1682   esalt_buf0[ 3] = esalt_bufs[salt_pos].esalt_buf[ 3];
1683   esalt_buf0[ 4] = esalt_bufs[salt_pos].esalt_buf[ 4];
1684   esalt_buf0[ 5] = esalt_bufs[salt_pos].esalt_buf[ 5];
1685   esalt_buf0[ 6] = esalt_bufs[salt_pos].esalt_buf[ 6];
1686   esalt_buf0[ 7] = esalt_bufs[salt_pos].esalt_buf[ 7];
1687   esalt_buf0[ 8] = esalt_bufs[salt_pos].esalt_buf[ 8];
1688   esalt_buf0[ 9] = esalt_bufs[salt_pos].esalt_buf[ 9];
1689   esalt_buf0[10] = esalt_bufs[salt_pos].esalt_buf[10];
1690   esalt_buf0[11] = esalt_bufs[salt_pos].esalt_buf[11];
1691   esalt_buf0[12] = esalt_bufs[salt_pos].esalt_buf[12];
1692   esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13];
1693   esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14];
1694   esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15];
1695
1696   u32 esalt_buf1[16];
1697
1698   esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16];
1699   esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17];
1700   esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18];
1701   esalt_buf1[ 3] = esalt_bufs[salt_pos].esalt_buf[19];
1702   esalt_buf1[ 4] = esalt_bufs[salt_pos].esalt_buf[20];
1703   esalt_buf1[ 5] = esalt_bufs[salt_pos].esalt_buf[21];
1704   esalt_buf1[ 6] = esalt_bufs[salt_pos].esalt_buf[22];
1705   esalt_buf1[ 7] = esalt_bufs[salt_pos].esalt_buf[23];
1706   esalt_buf1[ 8] = esalt_bufs[salt_pos].esalt_buf[24];
1707   esalt_buf1[ 9] = esalt_bufs[salt_pos].esalt_buf[25];
1708   esalt_buf1[10] = esalt_bufs[salt_pos].esalt_buf[26];
1709   esalt_buf1[11] = esalt_bufs[salt_pos].esalt_buf[27];
1710   esalt_buf1[12] = esalt_bufs[salt_pos].esalt_buf[28];
1711   esalt_buf1[13] = esalt_bufs[salt_pos].esalt_buf[29];
1712   esalt_buf1[14] = esalt_bufs[salt_pos].esalt_buf[30];
1713   esalt_buf1[15] = esalt_bufs[salt_pos].esalt_buf[31];
1714
1715   u32 esalt_buf2[16];
1716
1717   esalt_buf2[ 0] = esalt_bufs[salt_pos].esalt_buf[32];
1718   esalt_buf2[ 1] = esalt_bufs[salt_pos].esalt_buf[33];
1719   esalt_buf2[ 2] = esalt_bufs[salt_pos].esalt_buf[34];
1720   esalt_buf2[ 3] = esalt_bufs[salt_pos].esalt_buf[35];
1721   esalt_buf2[ 4] = esalt_bufs[salt_pos].esalt_buf[36];
1722   esalt_buf2[ 5] = esalt_bufs[salt_pos].esalt_buf[37];
1723   esalt_buf2[ 6] = 0;
1724   esalt_buf2[ 7] = 0;
1725   esalt_buf2[ 8] = 0;
1726   esalt_buf2[ 9] = 0;
1727   esalt_buf2[10] = 0;
1728   esalt_buf2[11] = 0;
1729   esalt_buf2[12] = 0;
1730   esalt_buf2[13] = 0;
1731   esalt_buf2[14] = 0;
1732   esalt_buf2[15] = 0;
1733
1734   const u32 digest_esalt_len = 32 + esalt_len;
1735   const u32 remaining_bytes  = digest_esalt_len + 1 - 64; // substract previous block
1736
1737   /**
1738    * digest
1739    */
1740
1741   const u32 search[4] =
1742   {
1743     digests_buf[digests_offset].digest_buf[DGST_R0],
1744     digests_buf[digests_offset].digest_buf[DGST_R1],
1745     digests_buf[digests_offset].digest_buf[DGST_R2],
1746     digests_buf[digests_offset].digest_buf[DGST_R3]
1747   };
1748
1749   /**
1750    * loop
1751    */
1752
1753   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
1754   {
1755     const u32 pw_r_len = c_combs[il_pos].pw_len;
1756
1757     const u32 pw_len = pw_l_len + pw_r_len;
1758
1759     u32 wordr0[4];
1760
1761     wordr0[0] = c_combs[il_pos].i[0];
1762     wordr0[1] = c_combs[il_pos].i[1];
1763     wordr0[2] = c_combs[il_pos].i[2];
1764     wordr0[3] = c_combs[il_pos].i[3];
1765
1766     u32 wordr1[4];
1767
1768     wordr1[0] = c_combs[il_pos].i[4];
1769     wordr1[1] = c_combs[il_pos].i[5];
1770     wordr1[2] = c_combs[il_pos].i[6];
1771     wordr1[3] = c_combs[il_pos].i[7];
1772
1773     u32 wordr2[4];
1774
1775     wordr2[0] = 0;
1776     wordr2[1] = 0;
1777     wordr2[2] = 0;
1778     wordr2[3] = 0;
1779
1780     u32 wordr3[4];
1781
1782     wordr3[0] = 0;
1783     wordr3[1] = 0;
1784     wordr3[2] = 0;
1785     wordr3[3] = 0;
1786
1787     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
1788     {
1789       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
1790     }
1791
1792     u32x w0[4];
1793
1794     w0[0] = wordl0[0] | wordr0[0];
1795     w0[1] = wordl0[1] | wordr0[1];
1796     w0[2] = wordl0[2] | wordr0[2];
1797     w0[3] = wordl0[3] | wordr0[3];
1798
1799     u32x w1[4];
1800
1801     w1[0] = wordl1[0] | wordr1[0];
1802     w1[1] = wordl1[1] | wordr1[1];
1803     w1[2] = wordl1[2] | wordr1[2];
1804     w1[3] = wordl1[3] | wordr1[3];
1805
1806     u32x w2[4];
1807
1808     w2[0] = wordl2[0] | wordr2[0];
1809     w2[1] = wordl2[1] | wordr2[1];
1810     w2[2] = wordl2[2] | wordr2[2];
1811     w2[3] = wordl2[3] | wordr2[3];
1812
1813     u32x w3[4];
1814
1815     w3[0] = wordl3[0] | wordr3[0];
1816     w3[1] = wordl3[1] | wordr3[1];
1817     w3[2] = wordl3[2] | wordr3[2];
1818     w3[3] = wordl3[3] | wordr3[3];
1819
1820     const u32 pw_salt_len = salt_len + pw_len;
1821
1822     /*
1823      * HA1 = md5 ($salt . $pass)
1824      */
1825
1826     // append the pass to the salt
1827
1828     u32x block0[16];
1829
1830     block0[ 0] = salt_buf0[ 0];
1831     block0[ 1] = salt_buf0[ 1];
1832     block0[ 2] = salt_buf0[ 2];
1833     block0[ 3] = salt_buf0[ 3];
1834     block0[ 4] = salt_buf0[ 4];
1835     block0[ 5] = salt_buf0[ 5];
1836     block0[ 6] = salt_buf0[ 6];
1837     block0[ 7] = salt_buf0[ 7];
1838     block0[ 8] = salt_buf0[ 8];
1839     block0[ 9] = salt_buf0[ 9];
1840     block0[10] = salt_buf0[10];
1841     block0[11] = salt_buf0[11];
1842     block0[12] = salt_buf0[12];
1843     block0[13] = salt_buf0[13];
1844     block0[14] = salt_buf0[14];
1845     block0[15] = salt_buf0[15];
1846
1847     u32x block1[16];
1848
1849     block1[ 0] = salt_buf1[ 0];
1850     block1[ 1] = salt_buf1[ 1];
1851     block1[ 2] = salt_buf1[ 2];
1852     block1[ 3] = salt_buf1[ 3];
1853     block1[ 4] = salt_buf1[ 4];
1854     block1[ 5] = salt_buf1[ 5];
1855     block1[ 6] = salt_buf1[ 6];
1856     block1[ 7] = salt_buf1[ 7];
1857     block1[ 8] = salt_buf1[ 8];
1858     block1[ 9] = salt_buf1[ 9];
1859     block1[10] = salt_buf1[10];
1860     block1[11] = salt_buf1[11];
1861     block1[12] = salt_buf1[12];
1862     block1[13] = salt_buf1[13];
1863     block1[14] = salt_buf1[14];
1864     block1[15] = salt_buf1[15];
1865
1866     u32 block_len = 0;
1867
1868     block_len = memcat32 (block0, block1, salt_len, w0, w1, w2, w3, pw_len);
1869
1870     u32x w0_t[4];
1871
1872     w0_t[0] = block0[ 0];
1873     w0_t[1] = block0[ 1];
1874     w0_t[2] = block0[ 2];
1875     w0_t[3] = block0[ 3];
1876
1877     u32x w1_t[4];
1878
1879     w1_t[0] = block0[ 4];
1880     w1_t[1] = block0[ 5];
1881     w1_t[2] = block0[ 6];
1882     w1_t[3] = block0[ 7];
1883
1884     u32x w2_t[4];
1885
1886     w2_t[0] = block0[ 8];
1887     w2_t[1] = block0[ 9];
1888     w2_t[2] = block0[10];
1889     w2_t[3] = block0[11];
1890
1891     u32x w3_t[4];
1892
1893     w3_t[0] = block0[12];
1894     w3_t[1] = block0[13];
1895     w3_t[2] = block0[14];
1896     w3_t[3] = block0[15];
1897
1898     if (block_len < 56)
1899     {
1900       w3_t[2] = pw_salt_len * 8;
1901     }
1902
1903     // md5
1904
1905     u32x tmp2;
1906
1907     u32x a = MD5M_A;
1908     u32x b = MD5M_B;
1909     u32x c = MD5M_C;
1910     u32x d = MD5M_D;
1911
1912     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1913     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1914     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1915     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1916     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1917     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1918     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1919     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1920     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1921     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1922     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1923     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1924     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1925     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1926     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1927     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1928
1929     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1930     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1931     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1932     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1933     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1934     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1935     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1936     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1937     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1938     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1939     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1940     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1941     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1942     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1943     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1944     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1945
1946     MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1947     MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1948     MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1949     MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1950     MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1951     MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1952     MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1953     MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1954     MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1955     MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1956     MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1957     MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1958     MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1959     MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1960     MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1961     MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1962
1963     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1964     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1965     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1966     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1967     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1968     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1969     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1970     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1971     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1972     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1973     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1974     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1975     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1976     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1977     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1978     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1979
1980     a += MD5M_A;
1981     b += MD5M_B;
1982     c += MD5M_C;
1983     d += MD5M_D;
1984
1985     if (block_len > 55)
1986     {
1987       u32x r_a = a;
1988       u32x r_b = b;
1989       u32x r_c = c;
1990       u32x r_d = d;
1991
1992       w0_t[0] = block1[ 0];
1993       w0_t[1] = block1[ 1];
1994       w0_t[2] = block1[ 2];
1995       w0_t[3] = block1[ 3];
1996
1997       w1_t[0] = block1[ 4];
1998       w1_t[1] = block1[ 5];
1999       w1_t[2] = block1[ 6];
2000       w1_t[3] = block1[ 7];
2001
2002       w2_t[0] = block1[ 8];
2003       w2_t[1] = block1[ 9];
2004       w2_t[2] = block1[10];
2005       w2_t[3] = block1[11];
2006
2007       w3_t[0] = block1[12];
2008       w3_t[1] = block1[13];
2009       w3_t[2] = pw_salt_len * 8;
2010       w3_t[3] = 0;
2011
2012       MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
2013       MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
2014       MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
2015       MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
2016       MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
2017       MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
2018       MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
2019       MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
2020       MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
2021       MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
2022       MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
2023       MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
2024       MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
2025       MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
2026       MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
2027       MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
2028
2029       MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
2030       MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
2031       MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
2032       MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
2033       MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
2034       MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
2035       MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
2036       MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
2037       MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
2038       MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
2039       MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
2040       MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
2041       MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
2042       MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
2043       MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
2044       MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
2045
2046       MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
2047       MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
2048       MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
2049       MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
2050       MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
2051       MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
2052       MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
2053       MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
2054       MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
2055       MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
2056       MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
2057       MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
2058       MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
2059       MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
2060       MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
2061       MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
2062
2063       MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
2064       MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
2065       MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
2066       MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
2067       MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
2068       MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
2069       MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
2070       MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
2071       MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
2072       MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
2073       MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
2074       MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
2075       MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
2076       MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
2077       MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
2078       MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
2079
2080       a += r_a;
2081       b += r_b;
2082       c += r_c;
2083       d += r_d;
2084     }
2085
2086     /*
2087      * final = md5 ($HA1 . $esalt)
2088      * we have at least 2 MD5 blocks/transformations, but we might need 3
2089      */
2090
2091     w0_t[0] = uint_to_hex_lower8 ((a >>  0) & 255) <<  0
2092             | uint_to_hex_lower8 ((a >>  8) & 255) << 16;
2093     w0_t[1] = uint_to_hex_lower8 ((a >> 16) & 255) <<  0
2094             | uint_to_hex_lower8 ((a >> 24) & 255) << 16;
2095     w0_t[2] = uint_to_hex_lower8 ((b >>  0) & 255) <<  0
2096             | uint_to_hex_lower8 ((b >>  8) & 255) << 16;
2097     w0_t[3] = uint_to_hex_lower8 ((b >> 16) & 255) <<  0
2098             | uint_to_hex_lower8 ((b >> 24) & 255) << 16;
2099     w1_t[0] = uint_to_hex_lower8 ((c >>  0) & 255) <<  0
2100             | uint_to_hex_lower8 ((c >>  8) & 255) << 16;
2101     w1_t[1] = uint_to_hex_lower8 ((c >> 16) & 255) <<  0
2102             | uint_to_hex_lower8 ((c >> 24) & 255) << 16;
2103     w1_t[2] = uint_to_hex_lower8 ((d >>  0) & 255) <<  0
2104             | uint_to_hex_lower8 ((d >>  8) & 255) << 16;
2105     w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) <<  0
2106             | uint_to_hex_lower8 ((d >> 24) & 255) << 16;
2107
2108     w2_t[0] = esalt_buf0[0];
2109     w2_t[1] = esalt_buf0[1];
2110     w2_t[2] = esalt_buf0[2];
2111     w2_t[3] = esalt_buf0[3];
2112
2113     w3_t[0] = esalt_buf0[4];
2114     w3_t[1] = esalt_buf0[5];
2115     w3_t[2] = esalt_buf0[6];
2116     w3_t[3] = esalt_buf0[7];
2117
2118     // md5
2119     // 1st transform
2120
2121     a = MD5M_A;
2122     b = MD5M_B;
2123     c = MD5M_C;
2124     d = MD5M_D;
2125
2126     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
2127     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
2128     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
2129     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
2130     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
2131     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
2132     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
2133     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
2134     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
2135     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
2136     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
2137     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
2138     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
2139     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
2140     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
2141     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
2142
2143     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
2144     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
2145     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
2146     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
2147     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
2148     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
2149     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
2150     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
2151     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
2152     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
2153     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
2154     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
2155     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
2156     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
2157     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
2158     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
2159
2160     MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
2161     MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
2162     MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
2163     MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
2164     MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
2165     MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
2166     MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
2167     MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
2168     MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
2169     MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
2170     MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
2171     MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
2172     MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
2173     MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
2174     MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
2175     MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
2176
2177     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
2178     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
2179     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
2180     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
2181     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
2182     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
2183     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
2184     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
2185     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
2186     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
2187     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
2188     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
2189     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
2190     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
2191     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
2192     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
2193
2194     a += MD5M_A;
2195     b += MD5M_B;
2196     c += MD5M_C;
2197     d += MD5M_D;
2198
2199     u32x r_a = a;
2200     u32x r_b = b;
2201     u32x r_c = c;
2202     u32x r_d = d;
2203
2204     // 2nd transform
2205
2206     w0_t[0] = esalt_buf0[ 8];
2207     w0_t[1] = esalt_buf0[ 9];
2208     w0_t[2] = esalt_buf0[10];
2209     w0_t[3] = esalt_buf0[11];
2210
2211     w1_t[0] = esalt_buf0[12];
2212     w1_t[1] = esalt_buf0[13];
2213     w1_t[2] = esalt_buf0[14];
2214     w1_t[3] = esalt_buf0[15];
2215
2216     w2_t[0] = esalt_buf1[ 0];
2217     w2_t[1] = esalt_buf1[ 1];
2218     w2_t[2] = esalt_buf1[ 2];
2219     w2_t[3] = esalt_buf1[ 3];
2220
2221     w3_t[0] = esalt_buf1[ 4];
2222     w3_t[1] = esalt_buf1[ 5];
2223     w3_t[2] = esalt_buf1[ 6];
2224     w3_t[3] = esalt_buf1[ 7];
2225
2226     // it is the final block when no more than 55 bytes left
2227
2228     if (remaining_bytes < 56)
2229     {
2230       // it is the last block !
2231
2232       w3_t[2] = digest_esalt_len * 8;
2233     }
2234
2235     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
2236     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
2237     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
2238     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
2239     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
2240     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
2241     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
2242     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
2243     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
2244     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
2245     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
2246     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
2247     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
2248     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
2249     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
2250     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
2251
2252     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
2253     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
2254     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
2255     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
2256     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
2257     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
2258     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
2259     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
2260     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
2261     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
2262     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
2263     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
2264     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
2265     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
2266     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
2267     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
2268
2269     MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
2270     MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
2271     MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
2272     MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
2273     MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
2274     MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
2275     MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
2276     MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
2277     MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
2278     MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
2279     MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
2280     MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
2281     MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
2282     MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
2283     MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
2284     MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
2285
2286     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
2287     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
2288     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
2289     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
2290     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
2291     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
2292     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
2293     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
2294     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
2295     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
2296     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
2297     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
2298     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
2299     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
2300     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
2301     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
2302
2303     // sometimes (not rare at all) we need a third block :(
2304
2305     if (remaining_bytes > 55)
2306     {
2307       // this is for sure the final block
2308
2309       a += r_a;
2310       b += r_b;
2311       c += r_c;
2312       d += r_d;
2313
2314       r_a = a;
2315       r_b = b;
2316       r_c = c;
2317       r_d = d;
2318
2319       w0_t[0] = esalt_buf1[ 8];
2320       w0_t[1] = esalt_buf1[ 9];
2321       w0_t[2] = esalt_buf1[10];
2322       w0_t[3] = esalt_buf1[11];
2323
2324       w1_t[0] = esalt_buf1[12];
2325       w1_t[1] = esalt_buf1[13];
2326       w1_t[2] = esalt_buf1[14];
2327       w1_t[3] = esalt_buf1[15];
2328
2329       w2_t[0] = esalt_buf2[ 0];
2330       w2_t[1] = esalt_buf2[ 1];
2331       w2_t[2] = esalt_buf2[ 2];
2332       w2_t[3] = esalt_buf2[ 3];
2333
2334       w3_t[0] = esalt_buf2[ 4];
2335       w3_t[1] = esalt_buf2[ 5];
2336       w3_t[2] = digest_esalt_len * 8;
2337       w3_t[3] = 0;
2338
2339       MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
2340       MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
2341       MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
2342       MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
2343       MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
2344       MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
2345       MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
2346       MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
2347       MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
2348       MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
2349       MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
2350       MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
2351       MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
2352       MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
2353       MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
2354       MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
2355
2356       MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
2357       MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
2358       MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
2359       MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
2360       MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
2361       MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
2362       MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
2363       MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
2364       MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
2365       MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
2366       MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
2367       MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
2368       MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
2369       MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
2370       MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
2371       MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
2372
2373       MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
2374       MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
2375       MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
2376       MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
2377       MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
2378       MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
2379       MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
2380       MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
2381       MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
2382       MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
2383       MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
2384       MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
2385       MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
2386       MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
2387       MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
2388       MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
2389
2390       MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
2391       MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
2392       MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
2393       MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
2394       MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
2395       MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
2396       MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
2397       MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
2398       MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
2399       MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
2400       MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
2401       MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
2402       MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
2403       MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
2404       MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
2405       MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
2406     }
2407
2408     a += r_a;
2409     b += r_b;
2410     c += r_c;
2411     d += r_d;
2412
2413     const u32x r0 = a;
2414     const u32x r1 = d;
2415     const u32x r2 = c;
2416     const u32x r3 = b;
2417
2418     #include VECT_COMPARE_S
2419   }
2420 }
2421
2422 extern "C" __global__ void __launch_bounds__ (256, 1) m11400_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 sip_t *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)
2423 {
2424 }
2425
2426 extern "C" __global__ void __launch_bounds__ (256, 1) m11400_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 sip_t *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)
2427 {
2428 }