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