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