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