nv/m07400.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SHA256_
   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 __device__ __constant__ u32 k_sha256[64] =
  33 {
  34   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
  35   SHA256C04, SHA256C05, SHA256C06, SHA256C07,
  36   SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
  37   SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
  38   SHA256C10, SHA256C11, SHA256C12, SHA256C13,
  39   SHA256C14, SHA256C15, SHA256C16, SHA256C17,
  40   SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
  41   SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
  42   SHA256C20, SHA256C21, SHA256C22, SHA256C23,
  43   SHA256C24, SHA256C25, SHA256C26, SHA256C27,
  44   SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
  45   SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
  46   SHA256C30, SHA256C31, SHA256C32, SHA256C33,
  47   SHA256C34, SHA256C35, SHA256C36, SHA256C37,
  48   SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
  49   SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
  50 };
  51
  52 __device__ static void sha256_transform (const u32x w[16], u32x digest[8])
  53 {
  54   u32x a = digest[0];
  55   u32x b = digest[1];
  56   u32x c = digest[2];
  57   u32x d = digest[3];
  58   u32x e = digest[4];
  59   u32x f = digest[5];
  60   u32x g = digest[6];
  61   u32x h = digest[7];
  62
  63   u32x w0_t = swap_workaround (w[ 0]);
  64   u32x w1_t = swap_workaround (w[ 1]);
  65   u32x w2_t = swap_workaround (w[ 2]);
  66   u32x w3_t = swap_workaround (w[ 3]);
  67   u32x w4_t = swap_workaround (w[ 4]);
  68   u32x w5_t = swap_workaround (w[ 5]);
  69   u32x w6_t = swap_workaround (w[ 6]);
  70   u32x w7_t = swap_workaround (w[ 7]);
  71   u32x w8_t = swap_workaround (w[ 8]);
  72   u32x w9_t = swap_workaround (w[ 9]);
  73   u32x wa_t = swap_workaround (w[10]);
  74   u32x wb_t = swap_workaround (w[11]);
  75   u32x wc_t = swap_workaround (w[12]);
  76   u32x wd_t = swap_workaround (w[13]);
  77   u32x we_t = swap_workaround (w[14]);
  78   u32x wf_t = swap_workaround (w[15]);
  79
  80   #define ROUND_EXPAND()                            \
  81   {                                                 \
  82     w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t);  \
  83     w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t);  \
  84     w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t);  \
  85     w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t);  \
  86     w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t);  \
  87     w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t);  \
  88     w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t);  \
  89     w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t);  \
  90     w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t);  \
  91     w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t);  \
  92     wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t);  \
  93     wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t);  \
  94     wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t);  \
  95     wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t);  \
  96     we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t);  \
  97     wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t);  \
  98   }
  99
 100   #define ROUND_STEP(i)                                                                   \
 101   {                                                                                       \
 102     SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i +  0]); \
 103     SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i +  1]); \
 104     SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i +  2]); \
 105     SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i +  3]); \
 106     SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i +  4]); \
 107     SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i +  5]); \
 108     SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i +  6]); \
 109     SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i +  7]); \
 110     SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i +  8]); \
 111     SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i +  9]); \
 112     SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
 113     SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
 114     SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
 115     SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
 116     SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
 117     SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
 118   }
 119
 120   ROUND_STEP (0);
 121
 122   for (int i = 16; i < 64; i += 16)
 123   {
 124     ROUND_EXPAND (); ROUND_STEP (i);
 125   }
 126
 127   digest[0] += a;
 128   digest[1] += b;
 129   digest[2] += c;
 130   digest[3] += d;
 131   digest[4] += e;
 132   digest[5] += f;
 133   digest[6] += g;
 134   digest[7] += h;
 135 }
 136
 137 __device__ static void sha256_transform_no14 (const u32x w[16], u32x digest[8])
 138 {
 139   u32x w_t[16];
 140
 141   w_t[ 0] = w[ 0];
 142   w_t[ 1] = w[ 1];
 143   w_t[ 2] = w[ 2];
 144   w_t[ 3] = w[ 3];
 145   w_t[ 4] = w[ 4];
 146   w_t[ 5] = w[ 5];
 147   w_t[ 6] = w[ 6];
 148   w_t[ 7] = w[ 7];
 149   w_t[ 8] = w[ 8];
 150   w_t[ 9] = w[ 9];
 151   w_t[10] = w[10];
 152   w_t[11] = w[11];
 153   w_t[12] = w[12];
 154   w_t[13] = w[13];
 155   w_t[14] = 0;
 156   w_t[15] = w[15];
 157
 158   sha256_transform (w_t, digest);
 159 }
 160
 161 __device__ static void init_ctx (u32x digest[8])
 162 {
 163   digest[0] = SHA256M_A;
 164   digest[1] = SHA256M_B;
 165   digest[2] = SHA256M_C;
 166   digest[3] = SHA256M_D;
 167   digest[4] = SHA256M_E;
 168   digest[5] = SHA256M_F;
 169   digest[6] = SHA256M_G;
 170   digest[7] = SHA256M_H;
 171 }
 172
 173 __device__ static void bzero16 (u32x block[16])
 174 {
 175   block[ 0] = 0;
 176   block[ 1] = 0;
 177   block[ 2] = 0;
 178   block[ 3] = 0;
 179   block[ 4] = 0;
 180   block[ 5] = 0;
 181   block[ 6] = 0;
 182   block[ 7] = 0;
 183   block[ 8] = 0;
 184   block[ 9] = 0;
 185   block[10] = 0;
 186   block[11] = 0;
 187   block[12] = 0;
 188   block[13] = 0;
 189   block[14] = 0;
 190   block[15] = 0;
 191 }
 192
 193 __device__ static void bswap8 (u32x block[16])
 194 {
 195   block[ 0] = swap_workaround (block[ 0]);
 196   block[ 1] = swap_workaround (block[ 1]);
 197   block[ 2] = swap_workaround (block[ 2]);
 198   block[ 3] = swap_workaround (block[ 3]);
 199   block[ 4] = swap_workaround (block[ 4]);
 200   block[ 5] = swap_workaround (block[ 5]);
 201   block[ 6] = swap_workaround (block[ 6]);
 202   block[ 7] = swap_workaround (block[ 7]);
 203 }
 204
 205 __device__ static u32 memcat16 (u32x block[16], const u32 block_len, const u32x append[4], const u32 append_len)
 206 {
 207   const u32 div = block_len / 4;
 208
 209   u32x tmp0;
 210   u32x tmp1;
 211   u32x tmp2;
 212   u32x tmp3;
 213   u32x tmp4;
 214
 215   #if __CUDA_ARCH__ >= 200
 216
 217   const int offset_minus_4 = 4 - (block_len & 3);
 218
 219   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
 220
 221   tmp0 = __byte_perm (        0, append[0], selector);
 222   tmp1 = __byte_perm (append[0], append[1], selector);
 223   tmp2 = __byte_perm (append[1], append[2], selector);
 224   tmp3 = __byte_perm (append[2], append[3], selector);
 225   tmp4 = __byte_perm (append[3],         0, selector);
 226
 227   #else
 228
 229   const u32 mod = block_len & 3;
 230
 231   switch (mod)
 232   {
 233     case 0: tmp0 = append[0];
 234             tmp1 = append[1];
 235             tmp2 = append[2];
 236             tmp3 = append[3];
 237             tmp4 = 0;
 238             break;
 239     case 1: tmp0 =                   append[0] <<  8;
 240             tmp1 = append[0] >> 24 | append[1] <<  8;
 241             tmp2 = append[1] >> 24 | append[2] <<  8;
 242             tmp3 = append[2] >> 24 | append[3] <<  8;
 243             tmp4 = append[3] >> 24;
 244             break;
 245     case 2: tmp0 =                   append[0] << 16;
 246             tmp1 = append[0] >> 16 | append[1] << 16;
 247             tmp2 = append[1] >> 16 | append[2] << 16;
 248             tmp3 = append[2] >> 16 | append[3] << 16;
 249             tmp4 = append[3] >> 16;
 250             break;
 251     case 3: tmp0 =                   append[0] << 24;
 252             tmp1 = append[0] >>  8 | append[1] << 24;
 253             tmp2 = append[1] >>  8 | append[2] << 24;
 254             tmp3 = append[2] >>  8 | append[3] << 24;
 255             tmp4 = append[3] >>  8;
 256             break;
 257   }
 258
 259   #endif
 260
 261   switch (div)
 262   {
 263     case  0:  block[ 0] |= tmp0;
 264               block[ 1]  = tmp1;
 265               block[ 2]  = tmp2;
 266               block[ 3]  = tmp3;
 267               block[ 4]  = tmp4;
 268               break;
 269     case  1:  block[ 1] |= tmp0;
 270               block[ 2]  = tmp1;
 271               block[ 3]  = tmp2;
 272               block[ 4]  = tmp3;
 273               block[ 5]  = tmp4;
 274               break;
 275     case  2:  block[ 2] |= tmp0;
 276               block[ 3]  = tmp1;
 277               block[ 4]  = tmp2;
 278               block[ 5]  = tmp3;
 279               block[ 6]  = tmp4;
 280               break;
 281     case  3:  block[ 3] |= tmp0;
 282               block[ 4]  = tmp1;
 283               block[ 5]  = tmp2;
 284               block[ 6]  = tmp3;
 285               block[ 7]  = tmp4;
 286               break;
 287     case  4:  block[ 4] |= tmp0;
 288               block[ 5]  = tmp1;
 289               block[ 6]  = tmp2;
 290               block[ 7]  = tmp3;
 291               block[ 8]  = tmp4;
 292               break;
 293     case  5:  block[ 5] |= tmp0;
 294               block[ 6]  = tmp1;
 295               block[ 7]  = tmp2;
 296               block[ 8]  = tmp3;
 297               block[ 9]  = tmp4;
 298               break;
 299     case  6:  block[ 6] |= tmp0;
 300               block[ 7]  = tmp1;
 301               block[ 8]  = tmp2;
 302               block[ 9]  = tmp3;
 303               block[10]  = tmp4;
 304               break;
 305     case  7:  block[ 7] |= tmp0;
 306               block[ 8]  = tmp1;
 307               block[ 9]  = tmp2;
 308               block[10]  = tmp3;
 309               block[11]  = tmp4;
 310               break;
 311     case  8:  block[ 8] |= tmp0;
 312               block[ 9]  = tmp1;
 313               block[10]  = tmp2;
 314               block[11]  = tmp3;
 315               block[12]  = tmp4;
 316               break;
 317     case  9:  block[ 9] |= tmp0;
 318               block[10]  = tmp1;
 319               block[11]  = tmp2;
 320               block[12]  = tmp3;
 321               block[13]  = tmp4;
 322               break;
 323     case 10:  block[10] |= tmp0;
 324               block[11]  = tmp1;
 325               block[12]  = tmp2;
 326               block[13]  = tmp3;
 327               block[14]  = tmp4;
 328               break;
 329     case 11:  block[11] |= tmp0;
 330               block[12]  = tmp1;
 331               block[13]  = tmp2;
 332               block[14]  = tmp3;
 333               block[15]  = tmp4;
 334               break;
 335     case 12:  block[12] |= tmp0;
 336               block[13]  = tmp1;
 337               block[14]  = tmp2;
 338               block[15]  = tmp3;
 339               break;
 340     case 13:  block[13] |= tmp0;
 341               block[14]  = tmp1;
 342               block[15]  = tmp2;
 343               break;
 344     case 14:  block[14] |= tmp0;
 345               block[15]  = tmp1;
 346               break;
 347     case 15:  block[15] |= tmp0;
 348               break;
 349   }
 350
 351   u32 new_len = block_len + append_len;
 352
 353   return new_len;
 354 }
 355
 356 __device__ static u32 memcat16c (u32x block[16], const u32 block_len, const u32x append[4], const u32 append_len, u32x digest[8])
 357 {
 358   const u32 div = block_len / 4;
 359
 360   u32x tmp0;
 361   u32x tmp1;
 362   u32x tmp2;
 363   u32x tmp3;
 364   u32x tmp4;
 365
 366   #if __CUDA_ARCH__ >= 200
 367
 368   const int offset_minus_4 = 4 - (block_len & 3);
 369
 370   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
 371
 372   tmp0 = __byte_perm (        0, append[0], selector);
 373   tmp1 = __byte_perm (append[0], append[1], selector);
 374   tmp2 = __byte_perm (append[1], append[2], selector);
 375   tmp3 = __byte_perm (append[2], append[3], selector);
 376   tmp4 = __byte_perm (append[3],         0, selector);
 377
 378   #else
 379
 380   const u32 mod = block_len & 3;
 381
 382   switch (mod)
 383   {
 384     case 0: tmp0 = append[0];
 385             tmp1 = append[1];
 386             tmp2 = append[2];
 387             tmp3 = append[3];
 388             tmp4 = 0;
 389             break;
 390     case 1: tmp0 =                   append[0] <<  8;
 391             tmp1 = append[0] >> 24 | append[1] <<  8;
 392             tmp2 = append[1] >> 24 | append[2] <<  8;
 393             tmp3 = append[2] >> 24 | append[3] <<  8;
 394             tmp4 = append[3] >> 24;
 395             break;
 396     case 2: tmp0 =                   append[0] << 16;
 397             tmp1 = append[0] >> 16 | append[1] << 16;
 398             tmp2 = append[1] >> 16 | append[2] << 16;
 399             tmp3 = append[2] >> 16 | append[3] << 16;
 400             tmp4 = append[3] >> 16;
 401             break;
 402     case 3: tmp0 =                   append[0] << 24;
 403             tmp1 = append[0] >>  8 | append[1] << 24;
 404             tmp2 = append[1] >>  8 | append[2] << 24;
 405             tmp3 = append[2] >>  8 | append[3] << 24;
 406             tmp4 = append[3] >>  8;
 407             break;
 408   }
 409
 410   #endif
 411
 412   u32x carry[4] = { 0, 0, 0, 0 };
 413
 414   switch (div)
 415   {
 416     case  0:  block[ 0] |= tmp0;
 417               block[ 1]  = tmp1;
 418               block[ 2]  = tmp2;
 419               block[ 3]  = tmp3;
 420               block[ 4]  = tmp4;
 421               break;
 422     case  1:  block[ 1] |= tmp0;
 423               block[ 2]  = tmp1;
 424               block[ 3]  = tmp2;
 425               block[ 4]  = tmp3;
 426               block[ 5]  = tmp4;
 427               break;
 428     case  2:  block[ 2] |= tmp0;
 429               block[ 3]  = tmp1;
 430               block[ 4]  = tmp2;
 431               block[ 5]  = tmp3;
 432               block[ 6]  = tmp4;
 433               break;
 434     case  3:  block[ 3] |= tmp0;
 435               block[ 4]  = tmp1;
 436               block[ 5]  = tmp2;
 437               block[ 6]  = tmp3;
 438               block[ 7]  = tmp4;
 439               break;
 440     case  4:  block[ 4] |= tmp0;
 441               block[ 5]  = tmp1;
 442               block[ 6]  = tmp2;
 443               block[ 7]  = tmp3;
 444               block[ 8]  = tmp4;
 445               break;
 446     case  5:  block[ 5] |= tmp0;
 447               block[ 6]  = tmp1;
 448               block[ 7]  = tmp2;
 449               block[ 8]  = tmp3;
 450               block[ 9]  = tmp4;
 451               break;
 452     case  6:  block[ 6] |= tmp0;
 453               block[ 7]  = tmp1;
 454               block[ 8]  = tmp2;
 455               block[ 9]  = tmp3;
 456               block[10]  = tmp4;
 457               break;
 458     case  7:  block[ 7] |= tmp0;
 459               block[ 8]  = tmp1;
 460               block[ 9]  = tmp2;
 461               block[10]  = tmp3;
 462               block[11]  = tmp4;
 463               break;
 464     case  8:  block[ 8] |= tmp0;
 465               block[ 9]  = tmp1;
 466               block[10]  = tmp2;
 467               block[11]  = tmp3;
 468               block[12]  = tmp4;
 469               break;
 470     case  9:  block[ 9] |= tmp0;
 471               block[10]  = tmp1;
 472               block[11]  = tmp2;
 473               block[12]  = tmp3;
 474               block[13]  = tmp4;
 475               break;
 476     case 10:  block[10] |= tmp0;
 477               block[11]  = tmp1;
 478               block[12]  = tmp2;
 479               block[13]  = tmp3;
 480               block[14]  = tmp4;
 481               break;
 482     case 11:  block[11] |= tmp0;
 483               block[12]  = tmp1;
 484               block[13]  = tmp2;
 485               block[14]  = tmp3;
 486               block[15]  = tmp4;
 487               break;
 488     case 12:  block[12] |= tmp0;
 489               block[13]  = tmp1;
 490               block[14]  = tmp2;
 491               block[15]  = tmp3;
 492               carry[ 0]  = tmp4;
 493               break;
 494     case 13:  block[13] |= tmp0;
 495               block[14]  = tmp1;
 496               block[15]  = tmp2;
 497               carry[ 0]  = tmp3;
 498               carry[ 1]  = tmp4;
 499               break;
 500     case 14:  block[14] |= tmp0;
 501               block[15]  = tmp1;
 502               carry[ 0]  = tmp2;
 503               carry[ 1]  = tmp3;
 504               carry[ 2]  = tmp4;
 505               break;
 506     case 15:  block[15] |= tmp0;
 507               carry[ 0]  = tmp1;
 508               carry[ 1]  = tmp2;
 509               carry[ 2]  = tmp3;
 510               carry[ 3]  = tmp4;
 511               break;
 512   }
 513
 514   u32 new_len = block_len + append_len;
 515
 516   if (new_len >= 64)
 517   {
 518     new_len -= 64;
 519
 520     sha256_transform (block, digest);
 521
 522     bzero16 (block);
 523
 524     block[0] = carry[0];
 525     block[1] = carry[1];
 526     block[2] = carry[2];
 527     block[3] = carry[3];
 528   }
 529
 530   return new_len;
 531 }
 532
 533 __device__ static u32 memcat20 (u32x block[20], const u32 block_len, const u32x append[4], const u32 append_len)
 534 {
 535   const u32 div = block_len / 4;
 536
 537   u32x tmp0;
 538   u32x tmp1;
 539   u32x tmp2;
 540   u32x tmp3;
 541   u32x tmp4;
 542
 543   #if __CUDA_ARCH__ >= 200
 544
 545   const int offset_minus_4 = 4 - (block_len & 3);
 546
 547   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
 548
 549   tmp0 = __byte_perm (        0, append[0], selector);
 550   tmp1 = __byte_perm (append[0], append[1], selector);
 551   tmp2 = __byte_perm (append[1], append[2], selector);
 552   tmp3 = __byte_perm (append[2], append[3], selector);
 553   tmp4 = __byte_perm (append[3],         0, selector);
 554
 555   #else
 556
 557   const u32 mod = block_len & 3;
 558
 559   switch (mod)
 560   {
 561     case 0: tmp0 = append[0];
 562             tmp1 = append[1];
 563             tmp2 = append[2];
 564             tmp3 = append[3];
 565             tmp4 = 0;
 566             break;
 567     case 1: tmp0 =                   append[0] <<  8;
 568             tmp1 = append[0] >> 24 | append[1] <<  8;
 569             tmp2 = append[1] >> 24 | append[2] <<  8;
 570             tmp3 = append[2] >> 24 | append[3] <<  8;
 571             tmp4 = append[3] >> 24;
 572             break;
 573     case 2: tmp0 =                   append[0] << 16;
 574             tmp1 = append[0] >> 16 | append[1] << 16;
 575             tmp2 = append[1] >> 16 | append[2] << 16;
 576             tmp3 = append[2] >> 16 | append[3] << 16;
 577             tmp4 = append[3] >> 16;
 578             break;
 579     case 3: tmp0 =                   append[0] << 24;
 580             tmp1 = append[0] >>  8 | append[1] << 24;
 581             tmp2 = append[1] >>  8 | append[2] << 24;
 582             tmp3 = append[2] >>  8 | append[3] << 24;
 583             tmp4 = append[3] >>  8;
 584             break;
 585   }
 586
 587   #endif
 588
 589   switch (div)
 590   {
 591     case  0:  block[ 0] |= tmp0;
 592               block[ 1]  = tmp1;
 593               block[ 2]  = tmp2;
 594               block[ 3]  = tmp3;
 595               block[ 4]  = tmp4;
 596               break;
 597     case  1:  block[ 1] |= tmp0;
 598               block[ 2]  = tmp1;
 599               block[ 3]  = tmp2;
 600               block[ 4]  = tmp3;
 601               block[ 5]  = tmp4;
 602               break;
 603     case  2:  block[ 2] |= tmp0;
 604               block[ 3]  = tmp1;
 605               block[ 4]  = tmp2;
 606               block[ 5]  = tmp3;
 607               block[ 6]  = tmp4;
 608               break;
 609     case  3:  block[ 3] |= tmp0;
 610               block[ 4]  = tmp1;
 611               block[ 5]  = tmp2;
 612               block[ 6]  = tmp3;
 613               block[ 7]  = tmp4;
 614               break;
 615     case  4:  block[ 4] |= tmp0;
 616               block[ 5]  = tmp1;
 617               block[ 6]  = tmp2;
 618               block[ 7]  = tmp3;
 619               block[ 8]  = tmp4;
 620               break;
 621     case  5:  block[ 5] |= tmp0;
 622               block[ 6]  = tmp1;
 623               block[ 7]  = tmp2;
 624               block[ 8]  = tmp3;
 625               block[ 9]  = tmp4;
 626               break;
 627     case  6:  block[ 6] |= tmp0;
 628               block[ 7]  = tmp1;
 629               block[ 8]  = tmp2;
 630               block[ 9]  = tmp3;
 631               block[10]  = tmp4;
 632               break;
 633     case  7:  block[ 7] |= tmp0;
 634               block[ 8]  = tmp1;
 635               block[ 9]  = tmp2;
 636               block[10]  = tmp3;
 637               block[11]  = tmp4;
 638               break;
 639     case  8:  block[ 8] |= tmp0;
 640               block[ 9]  = tmp1;
 641               block[10]  = tmp2;
 642               block[11]  = tmp3;
 643               block[12]  = tmp4;
 644               break;
 645     case  9:  block[ 9] |= tmp0;
 646               block[10]  = tmp1;
 647               block[11]  = tmp2;
 648               block[12]  = tmp3;
 649               block[13]  = tmp4;
 650               break;
 651     case 10:  block[10] |= tmp0;
 652               block[11]  = tmp1;
 653               block[12]  = tmp2;
 654               block[13]  = tmp3;
 655               block[14]  = tmp4;
 656               break;
 657     case 11:  block[11] |= tmp0;
 658               block[12]  = tmp1;
 659               block[13]  = tmp2;
 660               block[14]  = tmp3;
 661               block[15]  = tmp4;
 662               break;
 663     case 12:  block[12] |= tmp0;
 664               block[13]  = tmp1;
 665               block[14]  = tmp2;
 666               block[15]  = tmp3;
 667               block[16]  = tmp4;
 668               break;
 669     case 13:  block[13] |= tmp0;
 670               block[14]  = tmp1;
 671               block[15]  = tmp2;
 672               block[16]  = tmp3;
 673               block[17]  = tmp4;
 674               break;
 675     case 14:  block[14] |= tmp0;
 676               block[15]  = tmp1;
 677               block[16]  = tmp2;
 678               block[17]  = tmp3;
 679               block[18]  = tmp4;
 680               break;
 681     case 15:  block[15] |= tmp0;
 682               block[16]  = tmp1;
 683               block[17]  = tmp2;
 684               block[18]  = tmp3;
 685               block[19]  = tmp4;
 686               break;
 687   }
 688
 689   return block_len + append_len;
 690 }
 691
 692 __device__ static u32 memcat20_x80 (u32x block[20], const u32 block_len, const u32x append[4], const u32 append_len)
 693 {
 694   const u32 div = block_len / 4;
 695
 696   u32x tmp0;
 697   u32x tmp1;
 698   u32x tmp2;
 699   u32x tmp3;
 700   u32x tmp4;
 701
 702   #if __CUDA_ARCH__ >= 200
 703
 704   const int offset_minus_4 = 4 - (block_len & 3);
 705
 706   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
 707
 708   tmp0 = __byte_perm (        0, append[0], selector);
 709   tmp1 = __byte_perm (append[0], append[1], selector);
 710   tmp2 = __byte_perm (append[1], append[2], selector);
 711   tmp3 = __byte_perm (append[2], append[3], selector);
 712   tmp4 = __byte_perm (append[3],      0x80, selector);
 713
 714   #else
 715
 716   const u32 mod = block_len & 3;
 717
 718   switch (mod)
 719   {
 720     case 0: tmp0 = append[0];
 721             tmp1 = append[1];
 722             tmp2 = append[2];
 723             tmp3 = append[3];
 724             tmp4 = 0;
 725             break;
 726     case 1: tmp0 =                   append[0] <<  8;
 727             tmp1 = append[0] >> 24 | append[1] <<  8;
 728             tmp2 = append[1] >> 24 | append[2] <<  8;
 729             tmp3 = append[2] >> 24 | append[3] <<  8;
 730             tmp4 = append[3] >> 24;
 731             break;
 732     case 2: tmp0 =                   append[0] << 16;
 733             tmp1 = append[0] >> 16 | append[1] << 16;
 734             tmp2 = append[1] >> 16 | append[2] << 16;
 735             tmp3 = append[2] >> 16 | append[3] << 16;
 736             tmp4 = append[3] >> 16;
 737             break;
 738     case 3: tmp0 =                   append[0] << 24;
 739             tmp1 = append[0] >>  8 | append[1] << 24;
 740             tmp2 = append[1] >>  8 | append[2] << 24;
 741             tmp3 = append[2] >>  8 | append[3] << 24;
 742             tmp4 = append[3] >>  8;
 743             break;
 744   }
 745
 746   #endif
 747
 748   switch (div)
 749   {
 750     case  0:  block[ 0] |= tmp0;
 751               block[ 1]  = tmp1;
 752               block[ 2]  = tmp2;
 753               block[ 3]  = tmp3;
 754               block[ 4]  = tmp4;
 755               break;
 756     case  1:  block[ 1] |= tmp0;
 757               block[ 2]  = tmp1;
 758               block[ 3]  = tmp2;
 759               block[ 4]  = tmp3;
 760               block[ 5]  = tmp4;
 761               break;
 762     case  2:  block[ 2] |= tmp0;
 763               block[ 3]  = tmp1;
 764               block[ 4]  = tmp2;
 765               block[ 5]  = tmp3;
 766               block[ 6]  = tmp4;
 767               break;
 768     case  3:  block[ 3] |= tmp0;
 769               block[ 4]  = tmp1;
 770               block[ 5]  = tmp2;
 771               block[ 6]  = tmp3;
 772               block[ 7]  = tmp4;
 773               break;
 774     case  4:  block[ 4] |= tmp0;
 775               block[ 5]  = tmp1;
 776               block[ 6]  = tmp2;
 777               block[ 7]  = tmp3;
 778               block[ 8]  = tmp4;
 779               break;
 780     case  5:  block[ 5] |= tmp0;
 781               block[ 6]  = tmp1;
 782               block[ 7]  = tmp2;
 783               block[ 8]  = tmp3;
 784               block[ 9]  = tmp4;
 785               break;
 786     case  6:  block[ 6] |= tmp0;
 787               block[ 7]  = tmp1;
 788               block[ 8]  = tmp2;
 789               block[ 9]  = tmp3;
 790               block[10]  = tmp4;
 791               break;
 792     case  7:  block[ 7] |= tmp0;
 793               block[ 8]  = tmp1;
 794               block[ 9]  = tmp2;
 795               block[10]  = tmp3;
 796               block[11]  = tmp4;
 797               break;
 798     case  8:  block[ 8] |= tmp0;
 799               block[ 9]  = tmp1;
 800               block[10]  = tmp2;
 801               block[11]  = tmp3;
 802               block[12]  = tmp4;
 803               break;
 804     case  9:  block[ 9] |= tmp0;
 805               block[10]  = tmp1;
 806               block[11]  = tmp2;
 807               block[12]  = tmp3;
 808               block[13]  = tmp4;
 809               break;
 810     case 10:  block[10] |= tmp0;
 811               block[11]  = tmp1;
 812               block[12]  = tmp2;
 813               block[13]  = tmp3;
 814               block[14]  = tmp4;
 815               break;
 816     case 11:  block[11] |= tmp0;
 817               block[12]  = tmp1;
 818               block[13]  = tmp2;
 819               block[14]  = tmp3;
 820               block[15]  = tmp4;
 821               break;
 822     case 12:  block[12] |= tmp0;
 823               block[13]  = tmp1;
 824               block[14]  = tmp2;
 825               block[15]  = tmp3;
 826               block[16]  = tmp4;
 827               break;
 828     case 13:  block[13] |= tmp0;
 829               block[14]  = tmp1;
 830               block[15]  = tmp2;
 831               block[16]  = tmp3;
 832               block[17]  = tmp4;
 833               break;
 834     case 14:  block[14] |= tmp0;
 835               block[15]  = tmp1;
 836               block[16]  = tmp2;
 837               block[17]  = tmp3;
 838               block[18]  = tmp4;
 839               break;
 840     case 15:  block[15] |= tmp0;
 841               block[16]  = tmp1;
 842               block[17]  = tmp2;
 843               block[18]  = tmp3;
 844               block[19]  = tmp4;
 845               break;
 846   }
 847
 848   return block_len + append_len;
 849 }
 850
 851 extern "C" __global__ void __launch_bounds__ (256, 1) m07400_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, sha256crypt_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)
 852 {
 853   /**
 854    * base
 855    */
 856
 857   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 858
 859   if (gid >= gid_max) return;
 860
 861   u32x w0[4];
 862
 863   w0[0] = pws[gid].i[0];
 864   w0[1] = pws[gid].i[1];
 865   w0[2] = pws[gid].i[2];
 866   w0[3] = pws[gid].i[3];
 867
 868   const u32 pw_len = pws[gid].pw_len;
 869
 870   /**
 871    * salt
 872    */
 873
 874   u32 salt_buf[4];
 875
 876   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 877   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 878   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 879   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 880
 881   u32 salt_len = salt_bufs[salt_pos].salt_len;
 882
 883   /**
 884    * buffers
 885    */
 886
 887   u32 block_len;     // never reaches > 64
 888   u32 transform_len; // required for w[15] = len * 8
 889
 890   u32x block[16];
 891
 892   u32x alt_result[8];
 893   u32x p_bytes[8];
 894   u32x s_bytes[8];
 895
 896   /* Prepare for the real work.  */
 897
 898   block_len = 0;
 899
 900   bzero16 (block);
 901
 902   /* Add key.  */
 903
 904   block_len = memcat16 (block, block_len, w0, pw_len);
 905
 906   /* Add salt.  */
 907
 908   block_len = memcat16 (block, block_len, salt_buf, salt_len);
 909
 910   /* Add key again.  */
 911
 912   block_len = memcat16 (block, block_len, w0, pw_len);
 913
 914   append_0x80_4 (block, block_len);
 915
 916   block[15] = swap_workaround (block_len * 8);
 917
 918   init_ctx (alt_result);
 919
 920   sha256_transform (block, alt_result);
 921
 922   bswap8 (alt_result);
 923
 924   block_len = 0;
 925
 926   bzero16 (block);
 927
 928   u32x alt_result_tmp[8];
 929
 930   alt_result_tmp[0] = alt_result[0];
 931   alt_result_tmp[1] = alt_result[1];
 932   alt_result_tmp[2] = alt_result[2];
 933   alt_result_tmp[3] = alt_result[3];
 934   alt_result_tmp[4] = 0;
 935   alt_result_tmp[5] = 0;
 936   alt_result_tmp[6] = 0;
 937   alt_result_tmp[7] = 0;
 938
 939   truncate_block (alt_result_tmp, pw_len);
 940
 941   /* Add the key string.  */
 942
 943   block_len = memcat16 (block, block_len, w0, pw_len);
 944
 945   /* The last part is the salt string.  This must be at most 8
 946      characters and it ends at the first `$' character (for
 947      compatibility with existing implementations).  */
 948
 949   block_len = memcat16 (block, block_len, salt_buf, salt_len);
 950
 951   /* Now get result of this (32 bytes) and add it to the other
 952      context.  */
 953
 954   block_len = memcat16 (block, block_len, alt_result_tmp, pw_len);
 955
 956   transform_len = block_len;
 957
 958   /* Take the binary representation of the length of the key and for every
 959      1 add the alternate sum, for every 0 the key.  */
 960
 961   alt_result_tmp[0] = alt_result[0];
 962   alt_result_tmp[1] = alt_result[1];
 963   alt_result_tmp[2] = alt_result[2];
 964   alt_result_tmp[3] = alt_result[3];
 965   alt_result_tmp[4] = alt_result[4];
 966   alt_result_tmp[5] = alt_result[5];
 967   alt_result_tmp[6] = alt_result[6];
 968   alt_result_tmp[7] = alt_result[7];
 969
 970   init_ctx (alt_result);
 971
 972   for (u32 j = pw_len; j; j >>= 1)
 973   {
 974     if (j & 1)
 975     {
 976       block_len = memcat16c (block, block_len, &alt_result_tmp[0], 16, alt_result);
 977       block_len = memcat16c (block, block_len, &alt_result_tmp[4], 16, alt_result);
 978
 979       transform_len += 32;
 980     }
 981     else
 982     {
 983       block_len = memcat16c (block, block_len, w0, pw_len, alt_result);
 984
 985       transform_len += pw_len;
 986     }
 987   }
 988
 989   append_0x80_4 (block, block_len);
 990
 991   if (block_len >= 56)
 992   {
 993     sha256_transform (block, alt_result);
 994
 995     bzero16 (block);
 996   }
 997
 998   block[15] = swap_workaround (transform_len * 8);
 999
1000   sha256_transform (block, alt_result);
1001
1002   bswap8 (alt_result);
1003
1004   tmps[gid].alt_result[0] = alt_result[0];
1005   tmps[gid].alt_result[1] = alt_result[1];
1006   tmps[gid].alt_result[2] = alt_result[2];
1007   tmps[gid].alt_result[3] = alt_result[3];
1008   tmps[gid].alt_result[4] = alt_result[4];
1009   tmps[gid].alt_result[5] = alt_result[5];
1010   tmps[gid].alt_result[6] = alt_result[6];
1011   tmps[gid].alt_result[7] = alt_result[7];
1012
1013   /* Start computation of P byte sequence.  */
1014
1015   block_len = 0;
1016
1017   transform_len = 0;
1018
1019   bzero16 (block);
1020
1021   /* For every character in the password add the entire password.  */
1022
1023   init_ctx (p_bytes);
1024
1025   for (u32 j = 0; j < pw_len; j++)
1026   {
1027     block_len = memcat16c (block, block_len, w0, pw_len, p_bytes);
1028
1029     transform_len += pw_len;
1030   }
1031
1032   /* Finish the digest.  */
1033
1034   append_0x80_4 (block, block_len);
1035
1036   if (block_len >= 56)
1037   {
1038     sha256_transform (block, p_bytes);
1039
1040     bzero16 (block);
1041   }
1042
1043   block[15] = swap_workaround (transform_len * 8);
1044
1045   sha256_transform (block, p_bytes);
1046
1047   bswap8 (p_bytes);
1048
1049   truncate_block (p_bytes, pw_len);
1050
1051   tmps[gid].p_bytes[0] = p_bytes[0];
1052   tmps[gid].p_bytes[1] = p_bytes[1];
1053   tmps[gid].p_bytes[2] = p_bytes[2];
1054   tmps[gid].p_bytes[3] = p_bytes[3];
1055
1056   /* Start computation of S byte sequence.  */
1057
1058   block_len = 0;
1059
1060   transform_len = 0;
1061
1062   bzero16 (block);
1063
1064   /* For every character in the password add the entire password.  */
1065
1066   init_ctx (s_bytes);
1067
1068   for (u32 j = 0; j < 16 + (alt_result[0] & 0xff); j++)
1069   {
1070     block_len = memcat16c (block, block_len, salt_buf, salt_len, s_bytes);
1071
1072     transform_len += salt_len;
1073   }
1074
1075   /* Finish the digest.  */
1076
1077   append_0x80_4 (block, block_len);
1078
1079   if (block_len >= 56)
1080   {
1081     sha256_transform (block, s_bytes);
1082
1083     bzero16 (block);
1084   }
1085
1086   block[15] = swap_workaround (transform_len * 8);
1087
1088   sha256_transform (block, s_bytes);
1089
1090   bswap8 (s_bytes);
1091
1092   truncate_block (s_bytes, salt_len);
1093
1094   tmps[gid].s_bytes[0] = s_bytes[0];
1095   tmps[gid].s_bytes[1] = s_bytes[1];
1096   tmps[gid].s_bytes[2] = s_bytes[2];
1097   tmps[gid].s_bytes[3] = s_bytes[3];
1098 }
1099
1100 extern "C" __global__ void __launch_bounds__ (256, 1) m07400_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, sha256crypt_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)
1101 {
1102   /**
1103    * base
1104    */
1105
1106   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1107
1108   if (gid >= gid_max) return;
1109
1110   const u32 pw_len = pws[gid].pw_len;
1111
1112   /**
1113    * base
1114    */
1115
1116   u32x p_bytes[4];
1117
1118   p_bytes[0] = tmps[gid].p_bytes[0];
1119   p_bytes[1] = tmps[gid].p_bytes[1];
1120   p_bytes[2] = tmps[gid].p_bytes[2];
1121   p_bytes[3] = tmps[gid].p_bytes[3];
1122
1123   u32x p_bytes_x80[4];
1124
1125   p_bytes_x80[0] = tmps[gid].p_bytes[0];
1126   p_bytes_x80[1] = tmps[gid].p_bytes[1];
1127   p_bytes_x80[2] = tmps[gid].p_bytes[2];
1128   p_bytes_x80[3] = tmps[gid].p_bytes[3];
1129
1130   append_0x80_1 (p_bytes_x80, pw_len);
1131
1132   u32x s_bytes[4];
1133
1134   s_bytes[0] = tmps[gid].s_bytes[0];
1135   s_bytes[1] = tmps[gid].s_bytes[1];
1136   s_bytes[2] = tmps[gid].s_bytes[2];
1137   s_bytes[3] = tmps[gid].s_bytes[3];
1138
1139   u32x alt_result[8];
1140
1141   alt_result[0] = tmps[gid].alt_result[0];
1142   alt_result[1] = tmps[gid].alt_result[1];
1143   alt_result[2] = tmps[gid].alt_result[2];
1144   alt_result[3] = tmps[gid].alt_result[3];
1145   alt_result[4] = tmps[gid].alt_result[4];
1146   alt_result[5] = tmps[gid].alt_result[5];
1147   alt_result[6] = tmps[gid].alt_result[6];
1148   alt_result[7] = tmps[gid].alt_result[7];
1149
1150   u32 salt_len = salt_bufs[salt_pos].salt_len;
1151
1152   /* Repeatedly run the collected hash value through SHA256 to burn
1153      CPU cycles.  */
1154
1155   for (u32 i = 0, j = loop_pos; i < loop_cnt; i++, j++)
1156   {
1157     u32x tmp[8];
1158
1159     init_ctx (tmp);
1160
1161     u32x block[32];
1162
1163     bzero16 (&block[ 0]);
1164     bzero16 (&block[16]);
1165
1166     u32 block_len = 0;
1167
1168     const u32 j1 = (j & 1) ? 1 : 0;
1169     const u32 j3 = (j % 3) ? 1 : 0;
1170     const u32 j7 = (j % 7) ? 1 : 0;
1171
1172     if (j1)
1173     {
1174       block[0] = p_bytes[0];
1175       block[1] = p_bytes[1];
1176       block[2] = p_bytes[2];
1177       block[3] = p_bytes[3];
1178
1179       block_len = pw_len;
1180     }
1181     else
1182     {
1183       block[0] = alt_result[0];
1184       block[1] = alt_result[1];
1185       block[2] = alt_result[2];
1186       block[3] = alt_result[3];
1187       block[4] = alt_result[4];
1188       block[5] = alt_result[5];
1189       block[6] = alt_result[6];
1190       block[7] = alt_result[7];
1191
1192       block_len = 32;
1193     }
1194
1195     if (j3)
1196     {
1197       block_len = memcat20 (block, block_len, s_bytes, salt_len);
1198     }
1199
1200     if (j7)
1201     {
1202       block_len = memcat20 (block, block_len, p_bytes, pw_len);
1203     }
1204
1205     if (j1)
1206     {
1207       block_len = memcat20     (block, block_len, &alt_result[0], 16);
1208       block_len = memcat20_x80 (block, block_len, &alt_result[4], 16);
1209     }
1210     else
1211     {
1212       block_len = memcat20 (block, block_len, p_bytes_x80, pw_len);
1213     }
1214
1215     if (block_len >= 56)
1216     {
1217       sha256_transform (block, tmp);
1218
1219       block[ 0] = block[16];
1220       block[ 1] = block[17];
1221       block[ 2] = block[18];
1222       block[ 3] = block[19];
1223       block[ 4] = 0;
1224       block[ 5] = 0;
1225       block[ 6] = 0;
1226       block[ 7] = 0;
1227       block[ 8] = 0;
1228       block[ 9] = 0;
1229       block[10] = 0;
1230       block[11] = 0;
1231       block[12] = 0;
1232       block[13] = 0;
1233       block[14] = 0;
1234       block[15] = 0;
1235     }
1236
1237     block[15] = swap_workaround (block_len * 8);
1238
1239     sha256_transform_no14 (block, tmp);
1240
1241     bswap8 (tmp);
1242
1243     alt_result[0] = tmp[0];
1244     alt_result[1] = tmp[1];
1245     alt_result[2] = tmp[2];
1246     alt_result[3] = tmp[3];
1247     alt_result[4] = tmp[4];
1248     alt_result[5] = tmp[5];
1249     alt_result[6] = tmp[6];
1250     alt_result[7] = tmp[7];
1251   }
1252
1253   tmps[gid].alt_result[0] = alt_result[0];
1254   tmps[gid].alt_result[1] = alt_result[1];
1255   tmps[gid].alt_result[2] = alt_result[2];
1256   tmps[gid].alt_result[3] = alt_result[3];
1257   tmps[gid].alt_result[4] = alt_result[4];
1258   tmps[gid].alt_result[5] = alt_result[5];
1259   tmps[gid].alt_result[6] = alt_result[6];
1260   tmps[gid].alt_result[7] = alt_result[7];
1261 }
1262
1263 extern "C" __global__ void __launch_bounds__ (256, 1) m07400_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, sha256crypt_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)
1264 {
1265   /**
1266    * base
1267    */
1268
1269   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1270
1271   if (gid >= gid_max) return;
1272
1273   const u32 lid = threadIdx.x;
1274
1275   const u32x r0 = tmps[gid].alt_result[0];
1276   const u32x r1 = tmps[gid].alt_result[1];
1277   const u32x r2 = tmps[gid].alt_result[2];
1278   const u32x r3 = tmps[gid].alt_result[3];
1279
1280   #define il_pos 0
1281
1282   #include VECT_COMPARE_M
1283 }