nv/m06223.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SHA512_
   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 #include "gpu_aes256_nv.c"
  29 #include "gpu_twofish256_nv.c"
  30 #include "gpu_serpent256_nv.c"
  31
  32 __device__ __constant__ u64 k_sha512[80] =
  33 {
  34   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
  35   SHA512C04, SHA512C05, SHA512C06, SHA512C07,
  36   SHA512C08, SHA512C09, SHA512C0a, SHA512C0b,
  37   SHA512C0c, SHA512C0d, SHA512C0e, SHA512C0f,
  38   SHA512C10, SHA512C11, SHA512C12, SHA512C13,
  39   SHA512C14, SHA512C15, SHA512C16, SHA512C17,
  40   SHA512C18, SHA512C19, SHA512C1a, SHA512C1b,
  41   SHA512C1c, SHA512C1d, SHA512C1e, SHA512C1f,
  42   SHA512C20, SHA512C21, SHA512C22, SHA512C23,
  43   SHA512C24, SHA512C25, SHA512C26, SHA512C27,
  44   SHA512C28, SHA512C29, SHA512C2a, SHA512C2b,
  45   SHA512C2c, SHA512C2d, SHA512C2e, SHA512C2f,
  46   SHA512C30, SHA512C31, SHA512C32, SHA512C33,
  47   SHA512C34, SHA512C35, SHA512C36, SHA512C37,
  48   SHA512C38, SHA512C39, SHA512C3a, SHA512C3b,
  49   SHA512C3c, SHA512C3d, SHA512C3e, SHA512C3f,
  50   SHA512C40, SHA512C41, SHA512C42, SHA512C43,
  51   SHA512C44, SHA512C45, SHA512C46, SHA512C47,
  52   SHA512C48, SHA512C49, SHA512C4a, SHA512C4b,
  53   SHA512C4c, SHA512C4d, SHA512C4e, SHA512C4f,
  54 };
  55
  56 __device__ static void sha512_transform (const u64 w[16], u64 dgst[8])
  57 {
  58   u64 a = dgst[0];
  59   u64 b = dgst[1];
  60   u64 c = dgst[2];
  61   u64 d = dgst[3];
  62   u64 e = dgst[4];
  63   u64 f = dgst[5];
  64   u64 g = dgst[6];
  65   u64 h = dgst[7];
  66
  67   u64 w0_t = w[ 0];
  68   u64 w1_t = w[ 1];
  69   u64 w2_t = w[ 2];
  70   u64 w3_t = w[ 3];
  71   u64 w4_t = w[ 4];
  72   u64 w5_t = w[ 5];
  73   u64 w6_t = w[ 6];
  74   u64 w7_t = w[ 7];
  75   u64 w8_t = w[ 8];
  76   u64 w9_t = w[ 9];
  77   u64 wa_t = w[10];
  78   u64 wb_t = w[11];
  79   u64 wc_t = w[12];
  80   u64 wd_t = w[13];
  81   u64 we_t = w[14];
  82   u64 wf_t = w[15];
  83
  84   #define ROUND_EXPAND()                            \
  85   {                                                 \
  86     w0_t = SHA512_EXPAND (we_t, w9_t, w1_t, w0_t);  \
  87     w1_t = SHA512_EXPAND (wf_t, wa_t, w2_t, w1_t);  \
  88     w2_t = SHA512_EXPAND (w0_t, wb_t, w3_t, w2_t);  \
  89     w3_t = SHA512_EXPAND (w1_t, wc_t, w4_t, w3_t);  \
  90     w4_t = SHA512_EXPAND (w2_t, wd_t, w5_t, w4_t);  \
  91     w5_t = SHA512_EXPAND (w3_t, we_t, w6_t, w5_t);  \
  92     w6_t = SHA512_EXPAND (w4_t, wf_t, w7_t, w6_t);  \
  93     w7_t = SHA512_EXPAND (w5_t, w0_t, w8_t, w7_t);  \
  94     w8_t = SHA512_EXPAND (w6_t, w1_t, w9_t, w8_t);  \
  95     w9_t = SHA512_EXPAND (w7_t, w2_t, wa_t, w9_t);  \
  96     wa_t = SHA512_EXPAND (w8_t, w3_t, wb_t, wa_t);  \
  97     wb_t = SHA512_EXPAND (w9_t, w4_t, wc_t, wb_t);  \
  98     wc_t = SHA512_EXPAND (wa_t, w5_t, wd_t, wc_t);  \
  99     wd_t = SHA512_EXPAND (wb_t, w6_t, we_t, wd_t);  \
 100     we_t = SHA512_EXPAND (wc_t, w7_t, wf_t, we_t);  \
 101     wf_t = SHA512_EXPAND (wd_t, w8_t, w0_t, wf_t);  \
 102   }
 103
 104   #define ROUND_STEP(i)                                                                   \
 105   {                                                                                       \
 106     SHA512_STEP (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha512[i +  0]); \
 107     SHA512_STEP (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha512[i +  1]); \
 108     SHA512_STEP (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha512[i +  2]); \
 109     SHA512_STEP (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha512[i +  3]); \
 110     SHA512_STEP (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha512[i +  4]); \
 111     SHA512_STEP (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha512[i +  5]); \
 112     SHA512_STEP (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha512[i +  6]); \
 113     SHA512_STEP (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha512[i +  7]); \
 114     SHA512_STEP (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha512[i +  8]); \
 115     SHA512_STEP (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha512[i +  9]); \
 116     SHA512_STEP (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha512[i + 10]); \
 117     SHA512_STEP (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha512[i + 11]); \
 118     SHA512_STEP (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha512[i + 12]); \
 119     SHA512_STEP (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha512[i + 13]); \
 120     SHA512_STEP (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, we_t, k_sha512[i + 14]); \
 121     SHA512_STEP (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha512[i + 15]); \
 122   }
 123
 124   ROUND_STEP (0);
 125
 126   for (int i = 16; i < 80; i += 16)
 127   {
 128     ROUND_EXPAND (); ROUND_STEP (i);
 129   }
 130
 131   dgst[0] += a;
 132   dgst[1] += b;
 133   dgst[2] += c;
 134   dgst[3] += d;
 135   dgst[4] += e;
 136   dgst[5] += f;
 137   dgst[6] += g;
 138   dgst[7] += h;
 139 }
 140
 141 __device__ static void hmac_run (const u64 w1[16], const u64 ipad[8], const u64 opad[8], u64 dgst[8])
 142 {
 143   dgst[0] = ipad[0];
 144   dgst[1] = ipad[1];
 145   dgst[2] = ipad[2];
 146   dgst[3] = ipad[3];
 147   dgst[4] = ipad[4];
 148   dgst[5] = ipad[5];
 149   dgst[6] = ipad[6];
 150   dgst[7] = ipad[7];
 151
 152   sha512_transform (w1, dgst);
 153
 154   u64 w[16];
 155
 156   w[ 0] = dgst[0];
 157   w[ 1] = dgst[1];
 158   w[ 2] = dgst[2];
 159   w[ 3] = dgst[3];
 160   w[ 4] = dgst[4];
 161   w[ 5] = dgst[5];
 162   w[ 6] = dgst[6];
 163   w[ 7] = dgst[7];
 164   w[ 8] = 0x8000000000000000;
 165   w[ 9] = 0;
 166   w[10] = 0;
 167   w[11] = 0;
 168   w[12] = 0;
 169   w[13] = 0;
 170   w[14] = 0;
 171   w[15] = (128 + 64) * 8;
 172
 173   dgst[0] = opad[0];
 174   dgst[1] = opad[1];
 175   dgst[2] = opad[2];
 176   dgst[3] = opad[3];
 177   dgst[4] = opad[4];
 178   dgst[5] = opad[5];
 179   dgst[6] = opad[6];
 180   dgst[7] = opad[7];
 181
 182   sha512_transform (w, dgst);
 183 }
 184
 185 __device__ static void hmac_init (u64 w[16], u64 ipad[8], u64 opad[8])
 186 {
 187   w[ 0] ^= 0x3636363636363636;
 188   w[ 1] ^= 0x3636363636363636;
 189   w[ 2] ^= 0x3636363636363636;
 190   w[ 3] ^= 0x3636363636363636;
 191   w[ 4] ^= 0x3636363636363636;
 192   w[ 5] ^= 0x3636363636363636;
 193   w[ 6] ^= 0x3636363636363636;
 194   w[ 7] ^= 0x3636363636363636;
 195   w[ 8] ^= 0x3636363636363636;
 196   w[ 9] ^= 0x3636363636363636;
 197   w[10] ^= 0x3636363636363636;
 198   w[11] ^= 0x3636363636363636;
 199   w[12] ^= 0x3636363636363636;
 200   w[13] ^= 0x3636363636363636;
 201   w[14] ^= 0x3636363636363636;
 202   w[15] ^= 0x3636363636363636;
 203
 204   ipad[0] = SHA512M_A;
 205   ipad[1] = SHA512M_B;
 206   ipad[2] = SHA512M_C;
 207   ipad[3] = SHA512M_D;
 208   ipad[4] = SHA512M_E;
 209   ipad[5] = SHA512M_F;
 210   ipad[6] = SHA512M_G;
 211   ipad[7] = SHA512M_H;
 212
 213   sha512_transform (w, ipad);
 214
 215   w[ 0] ^= 0x6a6a6a6a6a6a6a6a;
 216   w[ 1] ^= 0x6a6a6a6a6a6a6a6a;
 217   w[ 2] ^= 0x6a6a6a6a6a6a6a6a;
 218   w[ 3] ^= 0x6a6a6a6a6a6a6a6a;
 219   w[ 4] ^= 0x6a6a6a6a6a6a6a6a;
 220   w[ 5] ^= 0x6a6a6a6a6a6a6a6a;
 221   w[ 6] ^= 0x6a6a6a6a6a6a6a6a;
 222   w[ 7] ^= 0x6a6a6a6a6a6a6a6a;
 223   w[ 8] ^= 0x6a6a6a6a6a6a6a6a;
 224   w[ 9] ^= 0x6a6a6a6a6a6a6a6a;
 225   w[10] ^= 0x6a6a6a6a6a6a6a6a;
 226   w[11] ^= 0x6a6a6a6a6a6a6a6a;
 227   w[12] ^= 0x6a6a6a6a6a6a6a6a;
 228   w[13] ^= 0x6a6a6a6a6a6a6a6a;
 229   w[14] ^= 0x6a6a6a6a6a6a6a6a;
 230   w[15] ^= 0x6a6a6a6a6a6a6a6a;
 231
 232   opad[0] = SHA512M_A;
 233   opad[1] = SHA512M_B;
 234   opad[2] = SHA512M_C;
 235   opad[3] = SHA512M_D;
 236   opad[4] = SHA512M_E;
 237   opad[5] = SHA512M_F;
 238   opad[6] = SHA512M_G;
 239   opad[7] = SHA512M_H;
 240
 241   sha512_transform (w, opad);
 242 }
 243
 244 __device__ static u32 u8add (const u32 a, const u32 b)
 245 {
 246   const u32 a1 = (a >>  0) & 0xff;
 247   const u32 a2 = (a >>  8) & 0xff;
 248   const u32 a3 = (a >> 16) & 0xff;
 249   const u32 a4 = (a >> 24) & 0xff;
 250
 251   const u32 b1 = (b >>  0) & 0xff;
 252   const u32 b2 = (b >>  8) & 0xff;
 253   const u32 b3 = (b >> 16) & 0xff;
 254   const u32 b4 = (b >> 24) & 0xff;
 255
 256   const u32 r1 = (a1 + b1) & 0xff;
 257   const u32 r2 = (a2 + b2) & 0xff;
 258   const u32 r3 = (a3 + b3) & 0xff;
 259   const u32 r4 = (a4 + b4) & 0xff;
 260
 261   const u32 r = r1 <<  0
 262                | r2 <<  8
 263                | r3 << 16
 264                | r4 << 24;
 265
 266   return r;
 267 }
 268
 269 extern "C" __global__ void __launch_bounds__ (256, 1) m06223_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc64_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 tc_t *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 270 {
 271   /**
 272    * base
 273    */
 274
 275   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 276
 277   if (gid >= gid_max) return;
 278
 279   u32x w0[4];
 280
 281   w0[0] = pws[gid].i[ 0];
 282   w0[1] = pws[gid].i[ 1];
 283   w0[2] = pws[gid].i[ 2];
 284   w0[3] = pws[gid].i[ 3];
 285
 286   u32x w1[4];
 287
 288   w1[0] = pws[gid].i[ 4];
 289   w1[1] = pws[gid].i[ 5];
 290   w1[2] = pws[gid].i[ 6];
 291   w1[3] = pws[gid].i[ 7];
 292
 293   u32x w2[4];
 294
 295   w2[0] = pws[gid].i[ 8];
 296   w2[1] = pws[gid].i[ 9];
 297   w2[2] = pws[gid].i[10];
 298   w2[3] = pws[gid].i[11];
 299
 300   u32x w3[4];
 301
 302   w3[0] = pws[gid].i[12];
 303   w3[1] = pws[gid].i[13];
 304   w3[2] = pws[gid].i[14];
 305   w3[3] = pws[gid].i[15];
 306
 307   /**
 308    * keyfile
 309    */
 310
 311   w0[0] = u8add (w0[0], esalt_bufs[salt_pos].keyfile_buf[ 0]);
 312   w0[1] = u8add (w0[1], esalt_bufs[salt_pos].keyfile_buf[ 1]);
 313   w0[2] = u8add (w0[2], esalt_bufs[salt_pos].keyfile_buf[ 2]);
 314   w0[3] = u8add (w0[3], esalt_bufs[salt_pos].keyfile_buf[ 3]);
 315   w1[0] = u8add (w1[0], esalt_bufs[salt_pos].keyfile_buf[ 4]);
 316   w1[1] = u8add (w1[1], esalt_bufs[salt_pos].keyfile_buf[ 5]);
 317   w1[2] = u8add (w1[2], esalt_bufs[salt_pos].keyfile_buf[ 6]);
 318   w1[3] = u8add (w1[3], esalt_bufs[salt_pos].keyfile_buf[ 7]);
 319   w2[0] = u8add (w2[0], esalt_bufs[salt_pos].keyfile_buf[ 8]);
 320   w2[1] = u8add (w2[1], esalt_bufs[salt_pos].keyfile_buf[ 9]);
 321   w2[2] = u8add (w2[2], esalt_bufs[salt_pos].keyfile_buf[10]);
 322   w2[3] = u8add (w2[3], esalt_bufs[salt_pos].keyfile_buf[11]);
 323   w3[0] = u8add (w3[0], esalt_bufs[salt_pos].keyfile_buf[12]);
 324   w3[1] = u8add (w3[1], esalt_bufs[salt_pos].keyfile_buf[13]);
 325   w3[2] = u8add (w3[2], esalt_bufs[salt_pos].keyfile_buf[14]);
 326   w3[3] = u8add (w3[3], esalt_bufs[salt_pos].keyfile_buf[15]);
 327
 328   /**
 329    * salt
 330    */
 331
 332   u64 salt_buf[16];
 333
 334   // swap fehlt
 335
 336   salt_buf[ 0] = ((u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 0])) << 32 | (u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 1]);
 337   salt_buf[ 1] = ((u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 2])) << 32 | (u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 3]);
 338   salt_buf[ 2] = ((u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 4])) << 32 | (u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 5]);
 339   salt_buf[ 3] = ((u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 6])) << 32 | (u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 7]);
 340   salt_buf[ 4] = ((u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 8])) << 32 | (u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[ 9]);
 341   salt_buf[ 5] = ((u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[10])) << 32 | (u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[11]);
 342   salt_buf[ 6] = ((u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[12])) << 32 | (u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[13]);
 343   salt_buf[ 7] = ((u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[14])) << 32 | (u64) swap_workaround (esalt_bufs[salt_pos].salt_buf[15]);
 344   salt_buf[ 8] = 0;
 345   salt_buf[ 9] = 0;
 346   salt_buf[10] = 0;
 347   salt_buf[11] = 0;
 348   salt_buf[12] = 0;
 349   salt_buf[13] = 0;
 350   salt_buf[14] = 0;
 351   salt_buf[15] = (128 + 64 + 4) * 8;
 352
 353   const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen;
 354
 355   u64 w[16];
 356
 357   w[ 0] = ((u64) swap_workaround (w0[0])) << 32 | (u64) swap_workaround (w0[1]);
 358   w[ 1] = ((u64) swap_workaround (w0[2])) << 32 | (u64) swap_workaround (w0[3]);
 359   w[ 2] = ((u64) swap_workaround (w1[0])) << 32 | (u64) swap_workaround (w1[1]);
 360   w[ 3] = ((u64) swap_workaround (w1[2])) << 32 | (u64) swap_workaround (w1[3]);
 361   w[ 4] = ((u64) swap_workaround (w2[0])) << 32 | (u64) swap_workaround (w2[1]);
 362   w[ 5] = ((u64) swap_workaround (w2[2])) << 32 | (u64) swap_workaround (w2[3]);
 363   w[ 6] = ((u64) swap_workaround (w3[0])) << 32 | (u64) swap_workaround (w3[1]);
 364   w[ 7] = ((u64) swap_workaround (w3[2])) << 32 | (u64) swap_workaround (w3[3]);
 365   w[ 8] = 0;
 366   w[ 9] = 0;
 367   w[10] = 0;
 368   w[11] = 0;
 369   w[12] = 0;
 370   w[13] = 0;
 371   w[14] = 0;
 372   w[15] = 0;
 373
 374   u64 ipad[8];
 375   u64 opad[8];
 376
 377   hmac_init (w, ipad, opad);
 378
 379   tmps[gid].ipad[0] = ipad[0];
 380   tmps[gid].ipad[1] = ipad[1];
 381   tmps[gid].ipad[2] = ipad[2];
 382   tmps[gid].ipad[3] = ipad[3];
 383   tmps[gid].ipad[4] = ipad[4];
 384   tmps[gid].ipad[5] = ipad[5];
 385   tmps[gid].ipad[6] = ipad[6];
 386   tmps[gid].ipad[7] = ipad[7];
 387
 388   tmps[gid].opad[0] = opad[0];
 389   tmps[gid].opad[1] = opad[1];
 390   tmps[gid].opad[2] = opad[2];
 391   tmps[gid].opad[3] = opad[3];
 392   tmps[gid].opad[4] = opad[4];
 393   tmps[gid].opad[5] = opad[5];
 394   tmps[gid].opad[6] = opad[6];
 395   tmps[gid].opad[7] = opad[7];
 396
 397   for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 8); i += 8, j += 1)
 398   {
 399     salt_buf[8] = (u64) j << 32 | (u64) 0x80000000;
 400
 401     u64 dgst[8];
 402
 403     hmac_run (salt_buf, ipad, opad, dgst);
 404
 405     tmps[gid].dgst[i + 0] = dgst[0];
 406     tmps[gid].dgst[i + 1] = dgst[1];
 407     tmps[gid].dgst[i + 2] = dgst[2];
 408     tmps[gid].dgst[i + 3] = dgst[3];
 409     tmps[gid].dgst[i + 4] = dgst[4];
 410     tmps[gid].dgst[i + 5] = dgst[5];
 411     tmps[gid].dgst[i + 6] = dgst[6];
 412     tmps[gid].dgst[i + 7] = dgst[7];
 413
 414     tmps[gid].out[i + 0] = dgst[0];
 415     tmps[gid].out[i + 1] = dgst[1];
 416     tmps[gid].out[i + 2] = dgst[2];
 417     tmps[gid].out[i + 3] = dgst[3];
 418     tmps[gid].out[i + 4] = dgst[4];
 419     tmps[gid].out[i + 5] = dgst[5];
 420     tmps[gid].out[i + 6] = dgst[6];
 421     tmps[gid].out[i + 7] = dgst[7];
 422   }
 423 }
 424
 425 extern "C" __global__ void __launch_bounds__ (256, 1) m06223_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc64_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 tc_t *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 426 {
 427   const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen;
 428
 429   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 430
 431   if (gid >= gid_max) return;
 432
 433   u64 ipad[8];
 434
 435   ipad[0] = tmps[gid].ipad[0];
 436   ipad[1] = tmps[gid].ipad[1];
 437   ipad[2] = tmps[gid].ipad[2];
 438   ipad[3] = tmps[gid].ipad[3];
 439   ipad[4] = tmps[gid].ipad[4];
 440   ipad[5] = tmps[gid].ipad[5];
 441   ipad[6] = tmps[gid].ipad[6];
 442   ipad[7] = tmps[gid].ipad[7];
 443
 444   u64 opad[8];
 445
 446   opad[0] = tmps[gid].opad[0];
 447   opad[1] = tmps[gid].opad[1];
 448   opad[2] = tmps[gid].opad[2];
 449   opad[3] = tmps[gid].opad[3];
 450   opad[4] = tmps[gid].opad[4];
 451   opad[5] = tmps[gid].opad[5];
 452   opad[6] = tmps[gid].opad[6];
 453   opad[7] = tmps[gid].opad[7];
 454
 455   for (u32 i = 0; i < (truecrypt_mdlen / 8 / 8); i += 8)
 456   {
 457     u64 dgst[8];
 458
 459     dgst[0] = tmps[gid].dgst[i + 0];
 460     dgst[1] = tmps[gid].dgst[i + 1];
 461     dgst[2] = tmps[gid].dgst[i + 2];
 462     dgst[3] = tmps[gid].dgst[i + 3];
 463     dgst[4] = tmps[gid].dgst[i + 4];
 464     dgst[5] = tmps[gid].dgst[i + 5];
 465     dgst[6] = tmps[gid].dgst[i + 6];
 466     dgst[7] = tmps[gid].dgst[i + 7];
 467
 468     u64 out[8];
 469
 470     out[0] = tmps[gid].out[i + 0];
 471     out[1] = tmps[gid].out[i + 1];
 472     out[2] = tmps[gid].out[i + 2];
 473     out[3] = tmps[gid].out[i + 3];
 474     out[4] = tmps[gid].out[i + 4];
 475     out[5] = tmps[gid].out[i + 5];
 476     out[6] = tmps[gid].out[i + 6];
 477     out[7] = tmps[gid].out[i + 7];
 478
 479     for (u32 j = 0; j < loop_cnt; j++)
 480     {
 481       u64 w[16];
 482
 483       w[ 0] = dgst[0];
 484       w[ 1] = dgst[1];
 485       w[ 2] = dgst[2];
 486       w[ 3] = dgst[3];
 487       w[ 4] = dgst[4];
 488       w[ 5] = dgst[5];
 489       w[ 6] = dgst[6];
 490       w[ 7] = dgst[7];
 491       w[ 8] = 0x8000000000000000;
 492       w[ 9] = 0;
 493       w[10] = 0;
 494       w[11] = 0;
 495       w[12] = 0;
 496       w[13] = 0;
 497       w[14] = 0;
 498       w[15] = (128 + 64) * 8;
 499
 500       hmac_run (w, ipad, opad, dgst);
 501
 502       out[0] ^= dgst[0];
 503       out[1] ^= dgst[1];
 504       out[2] ^= dgst[2];
 505       out[3] ^= dgst[3];
 506       out[4] ^= dgst[4];
 507       out[5] ^= dgst[5];
 508       out[6] ^= dgst[6];
 509       out[7] ^= dgst[7];
 510     }
 511
 512     tmps[gid].dgst[i + 0] = dgst[0];
 513     tmps[gid].dgst[i + 1] = dgst[1];
 514     tmps[gid].dgst[i + 2] = dgst[2];
 515     tmps[gid].dgst[i + 3] = dgst[3];
 516     tmps[gid].dgst[i + 4] = dgst[4];
 517     tmps[gid].dgst[i + 5] = dgst[5];
 518     tmps[gid].dgst[i + 6] = dgst[6];
 519     tmps[gid].dgst[i + 7] = dgst[7];
 520
 521     tmps[gid].out[i + 0] = out[0];
 522     tmps[gid].out[i + 1] = out[1];
 523     tmps[gid].out[i + 2] = out[2];
 524     tmps[gid].out[i + 3] = out[3];
 525     tmps[gid].out[i + 4] = out[4];
 526     tmps[gid].out[i + 5] = out[5];
 527     tmps[gid].out[i + 6] = out[6];
 528     tmps[gid].out[i + 7] = out[7];
 529   }
 530 }
 531
 532 extern "C" __global__ void __launch_bounds__ (256, 1) m06223_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc64_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 tc_t *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 533 {
 534   /**
 535    * base
 536    */
 537
 538   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 539   const u32 lid = threadIdx.x;
 540
 541   if (gid >= gid_max) return;
 542
 543   u32 ukey1[8];
 544
 545   ukey1[0] = swap_workaround (h32_from_64 (tmps[gid].out[ 0]));
 546   ukey1[1] = swap_workaround (l32_from_64 (tmps[gid].out[ 0]));
 547   ukey1[2] = swap_workaround (h32_from_64 (tmps[gid].out[ 1]));
 548   ukey1[3] = swap_workaround (l32_from_64 (tmps[gid].out[ 1]));
 549   ukey1[4] = swap_workaround (h32_from_64 (tmps[gid].out[ 2]));
 550   ukey1[5] = swap_workaround (l32_from_64 (tmps[gid].out[ 2]));
 551   ukey1[6] = swap_workaround (h32_from_64 (tmps[gid].out[ 3]));
 552   ukey1[7] = swap_workaround (l32_from_64 (tmps[gid].out[ 3]));
 553
 554   u32 ukey2[8];
 555
 556   ukey2[0] = swap_workaround (h32_from_64 (tmps[gid].out[ 4]));
 557   ukey2[1] = swap_workaround (l32_from_64 (tmps[gid].out[ 4]));
 558   ukey2[2] = swap_workaround (h32_from_64 (tmps[gid].out[ 5]));
 559   ukey2[3] = swap_workaround (l32_from_64 (tmps[gid].out[ 5]));
 560   ukey2[4] = swap_workaround (h32_from_64 (tmps[gid].out[ 6]));
 561   ukey2[5] = swap_workaround (l32_from_64 (tmps[gid].out[ 6]));
 562   ukey2[6] = swap_workaround (h32_from_64 (tmps[gid].out[ 7]));
 563   ukey2[7] = swap_workaround (l32_from_64 (tmps[gid].out[ 7]));
 564
 565   u32 data[4];
 566
 567   data[0] = esalt_bufs[0].data_buf[0];
 568   data[1] = esalt_bufs[0].data_buf[1];
 569   data[2] = esalt_bufs[0].data_buf[2];
 570   data[3] = esalt_bufs[0].data_buf[3];
 571
 572   u32 tmp[4];
 573
 574   {
 575     tmp[0] = data[0];
 576     tmp[1] = data[1];
 577     tmp[2] = data[2];
 578     tmp[3] = data[3];
 579
 580     aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
 581
 582     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 583     {
 584       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 585
 586       d_return_buf[lid] = 1;
 587     }
 588   }
 589
 590   {
 591     tmp[0] = data[0];
 592     tmp[1] = data[1];
 593     tmp[2] = data[2];
 594     tmp[3] = data[3];
 595
 596     serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
 597
 598     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 599     {
 600       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 601
 602       d_return_buf[lid] = 1;
 603     }
 604   }
 605
 606   {
 607     tmp[0] = data[0];
 608     tmp[1] = data[1];
 609     tmp[2] = data[2];
 610     tmp[3] = data[3];
 611
 612     twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
 613
 614     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 615     {
 616       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 617
 618       d_return_buf[lid] = 1;
 619     }
 620   }
 621
 622   u32 ukey3[8];
 623
 624   ukey3[0] = swap_workaround (h32_from_64 (tmps[gid].out[ 8]));
 625   ukey3[1] = swap_workaround (l32_from_64 (tmps[gid].out[ 8]));
 626   ukey3[2] = swap_workaround (h32_from_64 (tmps[gid].out[ 9]));
 627   ukey3[3] = swap_workaround (l32_from_64 (tmps[gid].out[ 9]));
 628   ukey3[4] = swap_workaround (h32_from_64 (tmps[gid].out[10]));
 629   ukey3[5] = swap_workaround (l32_from_64 (tmps[gid].out[10]));
 630   ukey3[6] = swap_workaround (h32_from_64 (tmps[gid].out[11]));
 631   ukey3[7] = swap_workaround (l32_from_64 (tmps[gid].out[11]));
 632
 633   u32 ukey4[8];
 634
 635   ukey4[0] = swap_workaround (h32_from_64 (tmps[gid].out[12]));
 636   ukey4[1] = swap_workaround (l32_from_64 (tmps[gid].out[12]));
 637   ukey4[2] = swap_workaround (h32_from_64 (tmps[gid].out[13]));
 638   ukey4[3] = swap_workaround (l32_from_64 (tmps[gid].out[13]));
 639   ukey4[4] = swap_workaround (h32_from_64 (tmps[gid].out[14]));
 640   ukey4[5] = swap_workaround (l32_from_64 (tmps[gid].out[14]));
 641   ukey4[6] = swap_workaround (h32_from_64 (tmps[gid].out[15]));
 642   ukey4[7] = swap_workaround (l32_from_64 (tmps[gid].out[15]));
 643
 644   {
 645     tmp[0] = data[0];
 646     tmp[1] = data[1];
 647     tmp[2] = data[2];
 648     tmp[3] = data[3];
 649
 650     aes256_decrypt_xts     (ukey2, ukey4, tmp, tmp);
 651     twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
 652
 653     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 654     {
 655       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 656
 657       d_return_buf[lid] = 1;
 658     }
 659   }
 660
 661   {
 662     tmp[0] = data[0];
 663     tmp[1] = data[1];
 664     tmp[2] = data[2];
 665     tmp[3] = data[3];
 666
 667     serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
 668     aes256_decrypt_xts     (ukey1, ukey3, tmp, tmp);
 669
 670     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 671     {
 672       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 673
 674       d_return_buf[lid] = 1;
 675     }
 676   }
 677
 678   {
 679     tmp[0] = data[0];
 680     tmp[1] = data[1];
 681     tmp[2] = data[2];
 682     tmp[3] = data[3];
 683
 684     twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
 685     serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
 686
 687     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 688     {
 689       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 690
 691       d_return_buf[lid] = 1;
 692     }
 693   }
 694
 695   u32 ukey5[8];
 696
 697   ukey5[0] = swap_workaround (h32_from_64 (tmps[gid].out[16]));
 698   ukey5[1] = swap_workaround (l32_from_64 (tmps[gid].out[16]));
 699   ukey5[2] = swap_workaround (h32_from_64 (tmps[gid].out[17]));
 700   ukey5[3] = swap_workaround (l32_from_64 (tmps[gid].out[17]));
 701   ukey5[4] = swap_workaround (h32_from_64 (tmps[gid].out[18]));
 702   ukey5[5] = swap_workaround (l32_from_64 (tmps[gid].out[18]));
 703   ukey5[6] = swap_workaround (h32_from_64 (tmps[gid].out[19]));
 704   ukey5[7] = swap_workaround (l32_from_64 (tmps[gid].out[19]));
 705
 706   u32 ukey6[8];
 707
 708   ukey6[0] = swap_workaround (h32_from_64 (tmps[gid].out[20]));
 709   ukey6[1] = swap_workaround (l32_from_64 (tmps[gid].out[20]));
 710   ukey6[2] = swap_workaround (h32_from_64 (tmps[gid].out[21]));
 711   ukey6[3] = swap_workaround (l32_from_64 (tmps[gid].out[21]));
 712   ukey6[4] = swap_workaround (h32_from_64 (tmps[gid].out[22]));
 713   ukey6[5] = swap_workaround (l32_from_64 (tmps[gid].out[22]));
 714   ukey6[6] = swap_workaround (h32_from_64 (tmps[gid].out[23]));
 715   ukey6[7] = swap_workaround (l32_from_64 (tmps[gid].out[23]));
 716
 717   {
 718     tmp[0] = data[0];
 719     tmp[1] = data[1];
 720     tmp[2] = data[2];
 721     tmp[3] = data[3];
 722
 723     aes256_decrypt_xts     (ukey3, ukey6, tmp, tmp);
 724     twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
 725     serpent256_decrypt_xts (ukey1, ukey4, tmp, tmp);
 726
 727     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 728     {
 729       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 730
 731       d_return_buf[lid] = 1;
 732     }
 733   }
 734
 735   {
 736     tmp[0] = data[0];
 737     tmp[1] = data[1];
 738     tmp[2] = data[2];
 739     tmp[3] = data[3];
 740
 741     serpent256_decrypt_xts (ukey3, ukey6, tmp, tmp);
 742     twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
 743     aes256_decrypt_xts     (ukey1, ukey4, tmp, tmp);
 744
 745     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 746     {
 747       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 748
 749       d_return_buf[lid] = 1;
 750     }
 751   }
 752 }