nv/m00500.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_SIZE2
  17 #endif
  18
  19 #define DGST_R0 0
  20 #define DGST_R1 1
  21 #define DGST_R2 2
  22 #define DGST_R3 3
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27
  28 #ifdef  VECT_SIZE1
  29 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  30 #endif
  31
  32 #ifdef  VECT_SIZE2
  33 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  34 #endif
  35
  36 #ifdef  VECT_SIZE4
  37 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  38 #endif
  39
  40 #define md5crypt_magic 0x00243124
  41
  42 __device__ static void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
  43 {
  44   u32x a = digest[0];
  45   u32x b = digest[1];
  46   u32x c = digest[2];
  47   u32x d = digest[3];
  48
  49   u32x w0_t = w0[0];
  50   u32x w1_t = w0[1];
  51   u32x w2_t = w0[2];
  52   u32x w3_t = w0[3];
  53   u32x w4_t = w1[0];
  54   u32x w5_t = w1[1];
  55   u32x w6_t = w1[2];
  56   u32x w7_t = w1[3];
  57   u32x w8_t = w2[0];
  58   u32x w9_t = w2[1];
  59   u32x wa_t = w2[2];
  60   u32x wb_t = w2[3];
  61   u32x wc_t = w3[0];
  62   u32x wd_t = w3[1];
  63   u32x we_t = w3[2];
  64   u32x wf_t = 0;
  65
  66   u32x tmp2;
  67
  68   MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
  69   MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
  70   MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
  71   MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
  72   MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
  73   MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
  74   MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
  75   MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
  76   MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
  77   MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
  78   MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
  79   MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
  80   MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
  81   MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
  82   MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
  83   MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
  84
  85   MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
  86   MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
  87   MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
  88   MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
  89   MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
  90   MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
  91   MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
  92   MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
  93   MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
  94   MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
  95   MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
  96   MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
  97   MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
  98   MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
  99   MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
 100   MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
 101
 102   MD5_STEP (MD5_H1, a, b, c, d, w5_t, MD5C20, MD5S20);
 103   MD5_STEP (MD5_H2, d, a, b, c, w8_t, MD5C21, MD5S21);
 104   MD5_STEP (MD5_H1, c, d, a, b, wb_t, MD5C22, MD5S22);
 105   MD5_STEP (MD5_H2, b, c, d, a, we_t, MD5C23, MD5S23);
 106   MD5_STEP (MD5_H1, a, b, c, d, w1_t, MD5C24, MD5S20);
 107   MD5_STEP (MD5_H2, d, a, b, c, w4_t, MD5C25, MD5S21);
 108   MD5_STEP (MD5_H1, c, d, a, b, w7_t, MD5C26, MD5S22);
 109   MD5_STEP (MD5_H2, b, c, d, a, wa_t, MD5C27, MD5S23);
 110   MD5_STEP (MD5_H1, a, b, c, d, wd_t, MD5C28, MD5S20);
 111   MD5_STEP (MD5_H2, d, a, b, c, w0_t, MD5C29, MD5S21);
 112   MD5_STEP (MD5_H1, c, d, a, b, w3_t, MD5C2a, MD5S22);
 113   MD5_STEP (MD5_H2, b, c, d, a, w6_t, MD5C2b, MD5S23);
 114   MD5_STEP (MD5_H1, a, b, c, d, w9_t, MD5C2c, MD5S20);
 115   MD5_STEP (MD5_H2, d, a, b, c, wc_t, MD5C2d, MD5S21);
 116   MD5_STEP (MD5_H1, c, d, a, b, wf_t, MD5C2e, MD5S22);
 117   MD5_STEP (MD5_H2, b, c, d, a, w2_t, MD5C2f, MD5S23);
 118
 119   MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
 120   MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
 121   MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
 122   MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
 123   MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
 124   MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
 125   MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
 126   MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
 127   MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
 128   MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
 129   MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
 130   MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
 131   MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
 132   MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
 133   MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
 134   MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
 135
 136   digest[0] += a;
 137   digest[1] += b;
 138   digest[2] += c;
 139   digest[3] += d;
 140 }
 141
 142 __device__ static void memcat16 (u32x block0[4], u32x block1[4], u32x block2[4], u32x block3[4], const u32 block_len, const u32x append[4])
 143 {
 144   u32x tmp0;
 145   u32x tmp1;
 146   u32x tmp2;
 147   u32x tmp3;
 148   u32x tmp4;
 149
 150   #if __CUDA_ARCH__ >= 200
 151
 152   const int offset_minus_4 = 4 - (block_len & 3);
 153
 154   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
 155
 156   tmp0 = __byte_perm (        0, append[0], selector);
 157   tmp1 = __byte_perm (append[0], append[1], selector);
 158   tmp2 = __byte_perm (append[1], append[2], selector);
 159   tmp3 = __byte_perm (append[2], append[3], selector);
 160   tmp4 = __byte_perm (append[3],         0, selector);
 161
 162   #else
 163
 164   const u32 mod = block_len & 3;
 165
 166   switch (mod)
 167   {
 168     case 0: tmp0 = append[0];
 169             tmp1 = append[1];
 170             tmp2 = append[2];
 171             tmp3 = append[3];
 172             tmp4 = 0;
 173             break;
 174     case 1: tmp0 =                   append[0] <<  8;
 175             tmp1 = append[0] >> 24 | append[1] <<  8;
 176             tmp2 = append[1] >> 24 | append[2] <<  8;
 177             tmp3 = append[2] >> 24 | append[3] <<  8;
 178             tmp4 = append[3] >> 24;
 179             break;
 180     case 2: tmp0 =                   append[0] << 16;
 181             tmp1 = append[0] >> 16 | append[1] << 16;
 182             tmp2 = append[1] >> 16 | append[2] << 16;
 183             tmp3 = append[2] >> 16 | append[3] << 16;
 184             tmp4 = append[3] >> 16;
 185             break;
 186     case 3: tmp0 =                   append[0] << 24;
 187             tmp1 = append[0] >>  8 | append[1] << 24;
 188             tmp2 = append[1] >>  8 | append[2] << 24;
 189             tmp3 = append[2] >>  8 | append[3] << 24;
 190             tmp4 = append[3] >>  8;
 191             break;
 192   }
 193
 194   #endif
 195
 196   const u32 div = block_len / 4;
 197
 198   switch (div)
 199   {
 200     case  0:  block0[0] |= tmp0;
 201               block0[1]  = tmp1;
 202               block0[2]  = tmp2;
 203               block0[3]  = tmp3;
 204               block1[0]  = tmp4;
 205               break;
 206     case  1:  block0[1] |= tmp0;
 207               block0[2]  = tmp1;
 208               block0[3]  = tmp2;
 209               block1[0]  = tmp3;
 210               block1[1]  = tmp4;
 211               break;
 212     case  2:  block0[2] |= tmp0;
 213               block0[3]  = tmp1;
 214               block1[0]  = tmp2;
 215               block1[1]  = tmp3;
 216               block1[2]  = tmp4;
 217               break;
 218     case  3:  block0[3] |= tmp0;
 219               block1[0]  = tmp1;
 220               block1[1]  = tmp2;
 221               block1[2]  = tmp3;
 222               block1[3]  = tmp4;
 223               break;
 224     case  4:  block1[0] |= tmp0;
 225               block1[1]  = tmp1;
 226               block1[2]  = tmp2;
 227               block1[3]  = tmp3;
 228               block2[0]  = tmp4;
 229               break;
 230     case  5:  block1[1] |= tmp0;
 231               block1[2]  = tmp1;
 232               block1[3]  = tmp2;
 233               block2[0]  = tmp3;
 234               block2[1]  = tmp4;
 235               break;
 236     case  6:  block1[2] |= tmp0;
 237               block1[3]  = tmp1;
 238               block2[0]  = tmp2;
 239               block2[1]  = tmp3;
 240               block2[2]  = tmp4;
 241               break;
 242     case  7:  block1[3] |= tmp0;
 243               block2[0]  = tmp1;
 244               block2[1]  = tmp2;
 245               block2[2]  = tmp3;
 246               block2[3]  = tmp4;
 247               break;
 248     case  8:  block2[0] |= tmp0;
 249               block2[1]  = tmp1;
 250               block2[2]  = tmp2;
 251               block2[3]  = tmp3;
 252               block3[0]  = tmp4;
 253               break;
 254     case  9:  block2[1] |= tmp0;
 255               block2[2]  = tmp1;
 256               block2[3]  = tmp2;
 257               block3[0]  = tmp3;
 258               block3[1]  = tmp4;
 259               break;
 260   }
 261
 262   return;
 263 }
 264
 265 __device__ static void memcat16_x80 (u32x block0[4], u32x block1[4], u32x block2[4], u32x block3[4], const u32 block_len, const u32x append[4])
 266 {
 267   u32x tmp0;
 268   u32x tmp1;
 269   u32x tmp2;
 270   u32x tmp3;
 271   u32x tmp4;
 272
 273   #if __CUDA_ARCH__ >= 200
 274
 275   const int offset_minus_4 = 4 - (block_len & 3);
 276
 277   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
 278
 279   tmp0 = __byte_perm (        0, append[0], selector);
 280   tmp1 = __byte_perm (append[0], append[1], selector);
 281   tmp2 = __byte_perm (append[1], append[2], selector);
 282   tmp3 = __byte_perm (append[2], append[3], selector);
 283   tmp4 = __byte_perm (append[3],      0x80, selector);
 284
 285   #else
 286
 287   const u32 mod = block_len & 3;
 288
 289   switch (mod)
 290   {
 291     case 0: tmp0 = append[0];
 292             tmp1 = append[1];
 293             tmp2 = append[2];
 294             tmp3 = append[3];
 295             tmp4 = 0x80;
 296             break;
 297     case 1: tmp0 =                   append[0] <<  8;
 298             tmp1 = append[0] >> 24 | append[1] <<  8;
 299             tmp2 = append[1] >> 24 | append[2] <<  8;
 300             tmp3 = append[2] >> 24 | append[3] <<  8;
 301             tmp4 = append[3] >> 24;
 302             break;
 303     case 2: tmp0 =                   append[0] << 16;
 304             tmp1 = append[0] >> 16 | append[1] << 16;
 305             tmp2 = append[1] >> 16 | append[2] << 16;
 306             tmp3 = append[2] >> 16 | append[3] << 16;
 307             tmp4 = append[3] >> 16;
 308             break;
 309     case 3: tmp0 =                   append[0] << 24;
 310             tmp1 = append[0] >>  8 | append[1] << 24;
 311             tmp2 = append[1] >>  8 | append[2] << 24;
 312             tmp3 = append[2] >>  8 | append[3] << 24;
 313             tmp4 = append[3] >>  8;
 314             break;
 315   }
 316
 317   #endif
 318
 319   const u32 div = block_len / 4;
 320
 321   switch (div)
 322   {
 323     case  0:  block0[0] |= tmp0;
 324               block0[1]  = tmp1;
 325               block0[2]  = tmp2;
 326               block0[3]  = tmp3;
 327               block1[0]  = tmp4;
 328               break;
 329     case  1:  block0[1] |= tmp0;
 330               block0[2]  = tmp1;
 331               block0[3]  = tmp2;
 332               block1[0]  = tmp3;
 333               block1[1]  = tmp4;
 334               break;
 335     case  2:  block0[2] |= tmp0;
 336               block0[3]  = tmp1;
 337               block1[0]  = tmp2;
 338               block1[1]  = tmp3;
 339               block1[2]  = tmp4;
 340               break;
 341     case  3:  block0[3] |= tmp0;
 342               block1[0]  = tmp1;
 343               block1[1]  = tmp2;
 344               block1[2]  = tmp3;
 345               block1[3]  = tmp4;
 346               break;
 347     case  4:  block1[0] |= tmp0;
 348               block1[1]  = tmp1;
 349               block1[2]  = tmp2;
 350               block1[3]  = tmp3;
 351               block2[0]  = tmp4;
 352               break;
 353     case  5:  block1[1] |= tmp0;
 354               block1[2]  = tmp1;
 355               block1[3]  = tmp2;
 356               block2[0]  = tmp3;
 357               block2[1]  = tmp4;
 358               break;
 359     case  6:  block1[2] |= tmp0;
 360               block1[3]  = tmp1;
 361               block2[0]  = tmp2;
 362               block2[1]  = tmp3;
 363               block2[2]  = tmp4;
 364               break;
 365     case  7:  block1[3] |= tmp0;
 366               block2[0]  = tmp1;
 367               block2[1]  = tmp2;
 368               block2[2]  = tmp3;
 369               block2[3]  = tmp4;
 370               break;
 371     case  8:  block2[0] |= tmp0;
 372               block2[1]  = tmp1;
 373               block2[2]  = tmp2;
 374               block2[3]  = tmp3;
 375               block3[0]  = tmp4;
 376               break;
 377     case  9:  block2[1] |= tmp0;
 378               block2[2]  = tmp1;
 379               block2[3]  = tmp2;
 380               block3[0]  = tmp3;
 381               block3[1]  = tmp4;
 382               break;
 383   }
 384
 385   return;
 386 }
 387
 388 __device__ static void memcat8 (u32x block0[4], u32x block1[4], u32x block2[4], u32x block3[4], const u32 block_len, const u32 append[2])
 389 {
 390   u32x tmp0;
 391   u32x tmp1;
 392   u32x tmp2;
 393
 394   #if __CUDA_ARCH__ >= 200
 395
 396   const int offset_minus_4 = 4 - (block_len & 3);
 397
 398   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
 399
 400   tmp0 = __byte_perm (        0, append[0], selector);
 401   tmp1 = __byte_perm (append[0], append[1], selector);
 402   tmp2 = __byte_perm (append[1],         0, selector);
 403
 404   #else
 405
 406   const u32 mod = block_len & 3;
 407
 408   switch (mod)
 409   {
 410     case 0: tmp0 = append[0];
 411             tmp1 = append[1];
 412             tmp2 = 0;
 413             break;
 414     case 1: tmp0 =                   append[0] <<  8;
 415             tmp1 = append[0] >> 24 | append[1] <<  8;
 416             tmp2 = append[1] >> 24;
 417             break;
 418     case 2: tmp0 =                   append[0] << 16;
 419             tmp1 = append[0] >> 16 | append[1] << 16;
 420             tmp2 = append[1] >> 16;
 421             break;
 422     case 3: tmp0 =                   append[0] << 24;
 423             tmp1 = append[0] >>  8 | append[1] << 24;
 424             tmp2 = append[1] >>  8;
 425             break;
 426   }
 427
 428   #endif
 429
 430   const u32 div = block_len / 4;
 431
 432   switch (div)
 433   {
 434     case  0:  block0[0] |= tmp0;
 435               block0[1]  = tmp1;
 436               block0[2]  = tmp2;
 437               break;
 438     case  1:  block0[1] |= tmp0;
 439               block0[2]  = tmp1;
 440               block0[3]  = tmp2;
 441               break;
 442     case  2:  block0[2] |= tmp0;
 443               block0[3]  = tmp1;
 444               block1[0]  = tmp2;
 445               break;
 446     case  3:  block0[3] |= tmp0;
 447               block1[0]  = tmp1;
 448               block1[1]  = tmp2;
 449               break;
 450     case  4:  block1[0] |= tmp0;
 451               block1[1]  = tmp1;
 452               block1[2]  = tmp2;
 453               break;
 454     case  5:  block1[1] |= tmp0;
 455               block1[2]  = tmp1;
 456               block1[3]  = tmp2;
 457               break;
 458     case  6:  block1[2] |= tmp0;
 459               block1[3]  = tmp1;
 460               block2[0]  = tmp2;
 461               break;
 462     case  7:  block1[3] |= tmp0;
 463               block2[0]  = tmp1;
 464               block2[1]  = tmp2;
 465               break;
 466     case  8:  block2[0] |= tmp0;
 467               block2[1]  = tmp1;
 468               block2[2]  = tmp2;
 469               break;
 470     case  9:  block2[1] |= tmp0;
 471               block2[2]  = tmp1;
 472               block2[3]  = tmp2;
 473               break;
 474     case 10:  block2[2] |= tmp0;
 475               block2[3]  = tmp1;
 476               block3[0]  = tmp2;
 477               break;
 478     case 11:  block2[3] |= tmp0;
 479               block3[0]  = tmp1;
 480               block3[1]  = tmp2;
 481               break;
 482   }
 483
 484   return;
 485 }
 486
 487 __device__ static void append_sign (u32x block0[4], u32x block1[4], const u32 block_len)
 488 {
 489   switch (block_len)
 490   {
 491     case 0:
 492       block0[0] = md5crypt_magic;
 493       break;
 494
 495     case 1:
 496       block0[0] = block0[0]            | md5crypt_magic <<  8;
 497       block0[1] = md5crypt_magic >> 24;
 498       break;
 499
 500     case 2:
 501       block0[0] = block0[0]            | md5crypt_magic << 16;
 502       block0[1] = md5crypt_magic >> 16;
 503       break;
 504
 505     case 3:
 506       block0[0] = block0[0]            | md5crypt_magic << 24;
 507       block0[1] = md5crypt_magic >>  8;
 508       break;
 509
 510     case 4:
 511       block0[1] = md5crypt_magic;
 512       break;
 513
 514     case 5:
 515       block0[1] = block0[1]            | md5crypt_magic <<  8;
 516       block0[2] = md5crypt_magic >> 24;
 517       break;
 518
 519     case 6:
 520       block0[1] = block0[1]            | md5crypt_magic << 16;
 521       block0[2] = md5crypt_magic >> 16;
 522       break;
 523
 524     case 7:
 525       block0[1] = block0[1]            | md5crypt_magic << 24;
 526       block0[2] = md5crypt_magic >>  8;
 527       break;
 528
 529     case 8:
 530       block0[2] = md5crypt_magic;
 531       break;
 532
 533     case 9:
 534       block0[2] = block0[2]            | md5crypt_magic <<  8;
 535       block0[3] = md5crypt_magic >> 24;
 536       break;
 537
 538     case 10:
 539       block0[2] = block0[2]            | md5crypt_magic << 16;
 540       block0[3] = md5crypt_magic >> 16;
 541       break;
 542
 543     case 11:
 544       block0[2] = block0[2]            | md5crypt_magic << 24;
 545       block0[3] = md5crypt_magic >>  8;
 546       break;
 547
 548     case 12:
 549       block0[3] = md5crypt_magic;
 550       break;
 551
 552     case 13:
 553       block0[3] = block0[3]            | md5crypt_magic <<  8;
 554       block1[0] = md5crypt_magic >> 24;
 555       break;
 556
 557     case 14:
 558       block0[3] = block0[3]            | md5crypt_magic << 16;
 559       block1[0] = md5crypt_magic >> 16;
 560       break;
 561
 562     case 15:
 563       block0[3] = block0[3]            | md5crypt_magic << 24;
 564       block1[0] = md5crypt_magic >>  8;
 565       break;
 566   }
 567 }
 568
 569 __device__ static void append_1st (u32x block0[4], u32x block1[4], u32x block2[4], u32x block3[4], const u32 block_len, const u32x append)
 570 {
 571   switch (block_len)
 572   {
 573     case 0:
 574       block0[0] = append;
 575       break;
 576
 577     case 1:
 578       block0[0] = block0[0] | append <<  8;
 579       break;
 580
 581     case 2:
 582       block0[0] = block0[0] | append << 16;
 583       break;
 584
 585     case 3:
 586       block0[0] = block0[0] | append << 24;
 587       break;
 588
 589     case 4:
 590       block0[1] = append;
 591       break;
 592
 593     case 5:
 594       block0[1] = block0[1] | append <<  8;
 595       break;
 596
 597     case 6:
 598       block0[1] = block0[1] | append << 16;
 599       break;
 600
 601     case 7:
 602       block0[1] = block0[1] | append << 24;
 603       break;
 604
 605     case 8:
 606       block0[2] = append;
 607       break;
 608
 609     case 9:
 610       block0[2] = block0[2] | append <<  8;
 611       break;
 612
 613     case 10:
 614       block0[2] = block0[2] | append << 16;
 615       break;
 616
 617     case 11:
 618       block0[2] = block0[2] | append << 24;
 619       break;
 620
 621     case 12:
 622       block0[3] = append;
 623       break;
 624
 625     case 13:
 626       block0[3] = block0[3] | append <<  8;
 627       break;
 628
 629     case 14:
 630       block0[3] = block0[3] | append << 16;
 631       break;
 632
 633     case 15:
 634       block0[3] = block0[3] | append << 24;
 635       break;
 636
 637     case 16:
 638       block1[0] = append;
 639       break;
 640
 641     case 17:
 642       block1[0] = block1[0] | append <<  8;
 643       break;
 644
 645     case 18:
 646       block1[0] = block1[0] | append << 16;
 647       break;
 648
 649     case 19:
 650       block1[0] = block1[0] | append << 24;
 651       break;
 652
 653     case 20:
 654       block1[1] = append;
 655       break;
 656
 657     case 21:
 658       block1[1] = block1[1] | append <<  8;
 659       break;
 660
 661     case 22:
 662       block1[1] = block1[1] | append << 16;
 663       break;
 664
 665     case 23:
 666       block1[1] = block1[1] | append << 24;
 667       break;
 668
 669     case 24:
 670       block1[2] = append;
 671       break;
 672
 673     case 25:
 674       block1[2] = block1[2] | append <<  8;
 675       break;
 676
 677     case 26:
 678       block1[2] = block1[2] | append << 16;
 679       break;
 680
 681     case 27:
 682       block1[2] = block1[2] | append << 24;
 683       break;
 684
 685     case 28:
 686       block1[3] = append;
 687       break;
 688
 689     case 29:
 690       block1[3] = block1[3] | append <<  8;
 691       break;
 692
 693     case 30:
 694       block1[3] = block1[3] | append << 16;
 695       break;
 696
 697     case 31:
 698       block1[3] = block1[3] | append << 24;
 699       break;
 700
 701     case 32:
 702       block2[0] = append;
 703       break;
 704
 705     case 33:
 706       block2[0] = block2[0] | append <<  8;
 707       break;
 708
 709     case 34:
 710       block2[0] = block2[0] | append << 16;
 711       break;
 712
 713     case 35:
 714       block2[0] = block2[0] | append << 24;
 715       break;
 716
 717     case 36:
 718       block2[1] = append;
 719       break;
 720
 721     case 37:
 722       block2[1] = block2[1] | append <<  8;
 723       break;
 724
 725     case 38:
 726       block2[1] = block2[1] | append << 16;
 727       break;
 728
 729     case 39:
 730       block2[1] = block2[1] | append << 24;
 731       break;
 732
 733     case 40:
 734       block2[2] = append;
 735       break;
 736
 737     case 41:
 738       block2[2] = block2[2] | append <<  8;
 739       break;
 740
 741     case 42:
 742       block2[2] = block2[2] | append << 16;
 743       break;
 744
 745     case 43:
 746       block2[2] = block2[2] | append << 24;
 747       break;
 748
 749     case 44:
 750       block2[3] = append;
 751       break;
 752
 753     case 45:
 754       block2[3] = block2[3] | append <<  8;
 755       break;
 756
 757     case 46:
 758       block2[3] = block2[3] | append << 16;
 759       break;
 760
 761     case 47:
 762       block2[3] = block2[3] | append << 24;
 763       break;
 764
 765     case 48:
 766       block3[0] = append;
 767       break;
 768
 769     case 49:
 770       block3[0] = block3[0] | append <<  8;
 771       break;
 772
 773     case 50:
 774       block3[0] = block3[0] | append << 16;
 775       break;
 776
 777     case 51:
 778       block3[0] = block3[0] | append << 24;
 779       break;
 780
 781     case 52:
 782       block3[1] = append;
 783       break;
 784
 785     case 53:
 786       block3[1] = block3[1] | append <<  8;
 787       break;
 788
 789     case 54:
 790       block3[1] = block3[1] | append << 16;
 791       break;
 792
 793     case 55:
 794       block3[1] = block3[1] | append << 24;
 795       break;
 796
 797     case 56:
 798       block3[2] = append;
 799       break;
 800   }
 801 }
 802
 803 extern "C" __global__ void __launch_bounds__ (256, 1) m00500_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, md5crypt_tmp_t *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 804 {
 805   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 806
 807   if (gid >= gid_max) return;
 808
 809   u32x w0[4];
 810
 811   w0[0] = pws[gid].i[0];
 812   w0[1] = pws[gid].i[1];
 813   w0[2] = pws[gid].i[2];
 814   w0[3] = pws[gid].i[3];
 815
 816   const u32 pw_len = pws[gid].pw_len;
 817
 818   /**
 819    * salt
 820    */
 821
 822   u32 salt_buf[2];
 823
 824   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 825   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 826
 827   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 828
 829   /**
 830    * init
 831    */
 832
 833   //memcat16 (block0, block1, block2, block3, block_len, w0);
 834   //block_len += pw_len;
 835
 836   u32 block_len = pw_len;
 837
 838   u32x block0[4];
 839
 840   block0[0] = w0[0];
 841   block0[1] = w0[1];
 842   block0[2] = w0[2];
 843   block0[3] = w0[3];
 844
 845   u32x block1[4];
 846
 847   block1[0] = 0;
 848   block1[1] = 0;
 849   block1[2] = 0;
 850   block1[3] = 0;
 851
 852   u32x block2[4];
 853
 854   block2[0] = 0;
 855   block2[1] = 0;
 856   block2[2] = 0;
 857   block2[3] = 0;
 858
 859   u32x block3[4];
 860
 861   block3[0] = 0;
 862   block3[1] = 0;
 863   block3[2] = 0;
 864   block3[3] = 0;
 865
 866   memcat8 (block0, block1, block2, block3, block_len, salt_buf);
 867
 868   block_len += salt_len;
 869
 870   memcat16 (block0, block1, block2, block3, block_len, w0);
 871
 872   block_len += pw_len;
 873
 874   append_0x80_4 (block0, block1, block2, block3, block_len);
 875
 876   block3[2] = block_len * 8;
 877
 878   u32x digest[4];
 879
 880   digest[0] = MD5M_A;
 881   digest[1] = MD5M_B;
 882   digest[2] = MD5M_C;
 883   digest[3] = MD5M_D;
 884
 885   md5_transform (block0, block1, block2, block3, digest);
 886
 887   /* The password first, since that is what is most unknown */
 888   /* Then our magic string */
 889   /* Then the raw salt */
 890   /* Then just as many characters of the MD5(pw,salt,pw) */
 891
 892   //memcat16 (block0, block1, block2, block3, block_len, w);
 893   //block_len += pw_len;
 894
 895   block_len = pw_len;
 896
 897   block0[0] = w0[0];
 898   block0[1] = w0[1];
 899   block0[2] = w0[2];
 900   block0[3] = w0[3];
 901
 902   block1[0] = 0;
 903   block1[1] = 0;
 904   block1[2] = 0;
 905   block1[3] = 0;
 906
 907   block2[0] = 0;
 908   block2[1] = 0;
 909   block2[2] = 0;
 910   block2[3] = 0;
 911
 912   block3[0] = 0;
 913   block3[1] = 0;
 914   block3[2] = 0;
 915   block3[3] = 0;
 916
 917   append_sign (block0, block1, block_len);
 918
 919   block_len += 3;
 920
 921   memcat8 (block0, block1, block2, block3, block_len, salt_buf);
 922
 923   block_len += salt_len;
 924
 925   truncate_block (digest, pw_len);
 926
 927   memcat16 (block0, block1, block2, block3, block_len, digest);
 928
 929   block_len += pw_len;
 930
 931   /* Then something really weird... */
 932
 933   u32x append = block0[0] & 0xFF;
 934
 935   for (u32 j = pw_len; j; j >>= 1)
 936   {
 937     if ((j & 1) == 0)
 938     {
 939       append_1st (block0, block1, block2, block3, block_len, append);
 940     }
 941
 942     block_len++;
 943   }
 944
 945   append_0x80_4 (block0, block1, block2, block3, block_len);
 946
 947   block3[2] = block_len * 8;
 948
 949   digest[0] = MD5M_A;
 950   digest[1] = MD5M_B;
 951   digest[2] = MD5M_C;
 952   digest[3] = MD5M_D;
 953
 954   md5_transform (block0, block1, block2, block3, digest);
 955
 956   tmps[gid].digest_buf[0] = digest[0];
 957   tmps[gid].digest_buf[1] = digest[1];
 958   tmps[gid].digest_buf[2] = digest[2];
 959   tmps[gid].digest_buf[3] = digest[3];
 960 }
 961
 962 extern "C" __global__ void __launch_bounds__ (256, 1) m00500_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, md5crypt_tmp_t *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 963 {
 964   /**
 965    * base
 966    */
 967
 968   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 969
 970   if (gid >= gid_max) return;
 971
 972   u32x w0[4];
 973
 974   w0[0] = pws[gid].i[0];
 975   w0[1] = pws[gid].i[1];
 976   w0[2] = pws[gid].i[2];
 977   w0[3] = pws[gid].i[3];
 978
 979   const u32 pw_len = pws[gid].pw_len;
 980
 981   u32x w0_x80[4];
 982
 983   w0_x80[0] = w0[0];
 984   w0_x80[1] = w0[1];
 985   w0_x80[2] = w0[2];
 986   w0_x80[3] = w0[3];
 987
 988   append_0x80_1 (w0_x80, pw_len);
 989
 990   /**
 991    * salt
 992    */
 993
 994   u32 salt_buf[2];
 995
 996   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 997   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 998
 999   const u32 salt_len = salt_bufs[salt_pos].salt_len;
1000
1001   /**
1002    * digest
1003    */
1004
1005   u32x digest[4];
1006
1007   digest[0] = tmps[gid].digest_buf[0];
1008   digest[1] = tmps[gid].digest_buf[1];
1009   digest[2] = tmps[gid].digest_buf[2];
1010   digest[3] = tmps[gid].digest_buf[3];
1011
1012   /**
1013    * loop
1014    */
1015
1016   /* and now, just to make sure things don't run too fast */
1017
1018   u32 block_len;
1019
1020   u32x block0[4];
1021
1022   block0[0] = 0;
1023   block0[1] = 0;
1024   block0[2] = 0;
1025   block0[3] = 0;
1026
1027   u32x block1[4];
1028
1029   block1[0] = 0;
1030   block1[1] = 0;
1031   block1[2] = 0;
1032   block1[3] = 0;
1033
1034   u32x block2[4];
1035
1036   block2[0] = 0;
1037   block2[1] = 0;
1038   block2[2] = 0;
1039   block2[3] = 0;
1040
1041   u32x block3[4];
1042
1043   block3[0] = 0;
1044   block3[1] = 0;
1045   block3[2] = 0;
1046   block3[3] = 0;
1047
1048   for (u32 i = 0, j = loop_pos; i < loop_cnt; i++, j++)
1049   {
1050     block1[0] = 0;
1051     block1[1] = 0;
1052     block1[2] = 0;
1053     block1[3] = 0;
1054     block2[0] = 0;
1055     block2[1] = 0;
1056     block2[2] = 0;
1057     block2[3] = 0;
1058     block3[0] = 0;
1059     block3[1] = 0;
1060
1061     const u32 j1 = (j & 1) ? 1 : 0;
1062     const u32 j3 = (j % 3) ? 1 : 0;
1063     const u32 j7 = (j % 7) ? 1 : 0;
1064
1065     if (j1)
1066     {
1067       block0[0] = w0[0];
1068       block0[1] = w0[1];
1069       block0[2] = w0[2];
1070       block0[3] = w0[3];
1071
1072       block_len = pw_len;
1073
1074       if (j3)
1075       {
1076         memcat8 (block0, block1, block2, block3, block_len, salt_buf);
1077
1078         block_len += salt_len;
1079       }
1080
1081       if (j7)
1082       {
1083         memcat16 (block0, block1, block2, block3, block_len, w0);
1084
1085         block_len += pw_len;
1086       }
1087
1088       memcat16_x80 (block0, block1, block2, block3, block_len, digest);
1089
1090       block_len += 16;
1091     }
1092     else
1093     {
1094       block0[0] = digest[0];
1095       block0[1] = digest[1];
1096       block0[2] = digest[2];
1097       block0[3] = digest[3];
1098
1099       block_len = 16;
1100
1101       if (j3 && j7)
1102       {
1103         block1[0] = salt_buf[0];
1104         block1[1] = salt_buf[1];
1105
1106         block_len += salt_len;
1107
1108         memcat16 (block0, block1, block2, block3, block_len, w0);
1109
1110         block_len += pw_len;
1111       }
1112       else if (j3)
1113       {
1114         block1[0] = salt_buf[0];
1115         block1[1] = salt_buf[1];
1116
1117         block_len += salt_len;
1118       }
1119       else if (j7)
1120       {
1121         block1[0] = w0[0];
1122         block1[1] = w0[1];
1123         block1[2] = w0[2];
1124         block1[3] = w0[3];
1125
1126         block_len += pw_len;
1127       }
1128
1129       memcat16 (block0, block1, block2, block3, block_len, w0_x80);
1130
1131       block_len += pw_len;
1132     }
1133
1134     block3[2] = block_len * 8;
1135
1136     digest[0] = MD5M_A;
1137     digest[1] = MD5M_B;
1138     digest[2] = MD5M_C;
1139     digest[3] = MD5M_D;
1140
1141     md5_transform (block0, block1, block2, block3, digest);
1142   }
1143
1144   tmps[gid].digest_buf[0] = digest[0];
1145   tmps[gid].digest_buf[1] = digest[1];
1146   tmps[gid].digest_buf[2] = digest[2];
1147   tmps[gid].digest_buf[3] = digest[3];
1148 }
1149
1150 extern "C" __global__ void __launch_bounds__ (256, 1) m00500_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, md5crypt_tmp_t *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1151 {
1152   /**
1153    * modifier
1154    */
1155
1156   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1157
1158   if (gid >= gid_max) return;
1159
1160   const u32 lid = threadIdx.x;
1161
1162   /**
1163    * digest
1164    */
1165
1166   const u32x r0 = tmps[gid].digest_buf[DGST_R0];
1167   const u32x r1 = tmps[gid].digest_buf[DGST_R1];
1168   const u32x r2 = tmps[gid].digest_buf[DGST_R2];
1169   const u32x r3 = tmps[gid].digest_buf[DGST_R3];
1170
1171   #define il_pos 0
1172
1173   #include VECT_COMPARE_M
1174 }