nv/m10900.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _PBKDF2_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
  26 #include "types_nv.c"
  27 #include "common_nv.c"
  28
  29 #ifdef  VECT_SIZE1
  30 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  31 #endif
  32
  33 __device__ static void sha256_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[8])
  34 {
  35   u32x a = digest[0];
  36   u32x b = digest[1];
  37   u32x c = digest[2];
  38   u32x d = digest[3];
  39   u32x e = digest[4];
  40   u32x f = digest[5];
  41   u32x g = digest[6];
  42   u32x h = digest[7];
  43
  44   u32x w0_t = w0[0];
  45   u32x w1_t = w0[1];
  46   u32x w2_t = w0[2];
  47   u32x w3_t = w0[3];
  48   u32x w4_t = w1[0];
  49   u32x w5_t = w1[1];
  50   u32x w6_t = w1[2];
  51   u32x w7_t = w1[3];
  52   u32x w8_t = w2[0];
  53   u32x w9_t = w2[1];
  54   u32x wa_t = w2[2];
  55   u32x wb_t = w2[3];
  56   u32x wc_t = w3[0];
  57   u32x wd_t = w3[1];
  58   u32x we_t = w3[2];
  59   u32x wf_t = w3[3];
  60
  61   SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C00);
  62   SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C01);
  63   SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C02);
  64   SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C03);
  65   SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C04);
  66   SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C05);
  67   SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C06);
  68   SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C07);
  69   SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C08);
  70   SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C09);
  71   SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C0a);
  72   SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C0b);
  73   SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C0c);
  74   SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C0d);
  75   SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C0e);
  76   SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C0f);
  77   w0_t = SHA256_S1(we_t) + w9_t + SHA256_S0(w1_t) + w0_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C10);
  78   w1_t = SHA256_S1(wf_t) + wa_t + SHA256_S0(w2_t) + w1_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C11);
  79   w2_t = SHA256_S1(w0_t) + wb_t + SHA256_S0(w3_t) + w2_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C12);
  80   w3_t = SHA256_S1(w1_t) + wc_t + SHA256_S0(w4_t) + w3_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C13);
  81   w4_t = SHA256_S1(w2_t) + wd_t + SHA256_S0(w5_t) + w4_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C14);
  82   w5_t = SHA256_S1(w3_t) + we_t + SHA256_S0(w6_t) + w5_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C15);
  83   w6_t = SHA256_S1(w4_t) + wf_t + SHA256_S0(w7_t) + w6_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C16);
  84   w7_t = SHA256_S1(w5_t) + w0_t + SHA256_S0(w8_t) + w7_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C17);
  85   w8_t = SHA256_S1(w6_t) + w1_t + SHA256_S0(w9_t) + w8_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C18);
  86   w9_t = SHA256_S1(w7_t) + w2_t + SHA256_S0(wa_t) + w9_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C19);
  87   wa_t = SHA256_S1(w8_t) + w3_t + SHA256_S0(wb_t) + wa_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C1a);
  88   wb_t = SHA256_S1(w9_t) + w4_t + SHA256_S0(wc_t) + wb_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C1b);
  89   wc_t = SHA256_S1(wa_t) + w5_t + SHA256_S0(wd_t) + wc_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C1c);
  90   wd_t = SHA256_S1(wb_t) + w6_t + SHA256_S0(we_t) + wd_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C1d);
  91   we_t = SHA256_S1(wc_t) + w7_t + SHA256_S0(wf_t) + we_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C1e);
  92   wf_t = SHA256_S1(wd_t) + w8_t + SHA256_S0(w0_t) + wf_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C1f);
  93   w0_t = SHA256_S1(we_t) + w9_t + SHA256_S0(w1_t) + w0_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C20);
  94   w1_t = SHA256_S1(wf_t) + wa_t + SHA256_S0(w2_t) + w1_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C21);
  95   w2_t = SHA256_S1(w0_t) + wb_t + SHA256_S0(w3_t) + w2_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C22);
  96   w3_t = SHA256_S1(w1_t) + wc_t + SHA256_S0(w4_t) + w3_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C23);
  97   w4_t = SHA256_S1(w2_t) + wd_t + SHA256_S0(w5_t) + w4_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C24);
  98   w5_t = SHA256_S1(w3_t) + we_t + SHA256_S0(w6_t) + w5_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C25);
  99   w6_t = SHA256_S1(w4_t) + wf_t + SHA256_S0(w7_t) + w6_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C26);
 100   w7_t = SHA256_S1(w5_t) + w0_t + SHA256_S0(w8_t) + w7_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C27);
 101   w8_t = SHA256_S1(w6_t) + w1_t + SHA256_S0(w9_t) + w8_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C28);
 102   w9_t = SHA256_S1(w7_t) + w2_t + SHA256_S0(wa_t) + w9_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C29);
 103   wa_t = SHA256_S1(w8_t) + w3_t + SHA256_S0(wb_t) + wa_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C2a);
 104   wb_t = SHA256_S1(w9_t) + w4_t + SHA256_S0(wc_t) + wb_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C2b);
 105   wc_t = SHA256_S1(wa_t) + w5_t + SHA256_S0(wd_t) + wc_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C2c);
 106   wd_t = SHA256_S1(wb_t) + w6_t + SHA256_S0(we_t) + wd_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C2d);
 107   we_t = SHA256_S1(wc_t) + w7_t + SHA256_S0(wf_t) + we_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C2e);
 108   wf_t = SHA256_S1(wd_t) + w8_t + SHA256_S0(w0_t) + wf_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C2f);
 109   w0_t = SHA256_S1(we_t) + w9_t + SHA256_S0(w1_t) + w0_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C30);
 110   w1_t = SHA256_S1(wf_t) + wa_t + SHA256_S0(w2_t) + w1_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C31);
 111   w2_t = SHA256_S1(w0_t) + wb_t + SHA256_S0(w3_t) + w2_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C32);
 112   w3_t = SHA256_S1(w1_t) + wc_t + SHA256_S0(w4_t) + w3_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C33);
 113   w4_t = SHA256_S1(w2_t) + wd_t + SHA256_S0(w5_t) + w4_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C34);
 114   w5_t = SHA256_S1(w3_t) + we_t + SHA256_S0(w6_t) + w5_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C35);
 115   w6_t = SHA256_S1(w4_t) + wf_t + SHA256_S0(w7_t) + w6_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C36);
 116   w7_t = SHA256_S1(w5_t) + w0_t + SHA256_S0(w8_t) + w7_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C37);
 117   w8_t = SHA256_S1(w6_t) + w1_t + SHA256_S0(w9_t) + w8_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C38);
 118   w9_t = SHA256_S1(w7_t) + w2_t + SHA256_S0(wa_t) + w9_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C39);
 119   wa_t = SHA256_S1(w8_t) + w3_t + SHA256_S0(wb_t) + wa_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C3a);
 120   wb_t = SHA256_S1(w9_t) + w4_t + SHA256_S0(wc_t) + wb_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C3b);
 121   wc_t = SHA256_S1(wa_t) + w5_t + SHA256_S0(wd_t) + wc_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C3c);
 122   wd_t = SHA256_S1(wb_t) + w6_t + SHA256_S0(we_t) + wd_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C3d);
 123   we_t = SHA256_S1(wc_t) + w7_t + SHA256_S0(wf_t) + we_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C3e);
 124   wf_t = SHA256_S1(wd_t) + w8_t + SHA256_S0(w0_t) + wf_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C3f);
 125
 126   digest[0] += a;
 127   digest[1] += b;
 128   digest[2] += c;
 129   digest[3] += d;
 130   digest[4] += e;
 131   digest[5] += f;
 132   digest[6] += g;
 133   digest[7] += h;
 134 }
 135
 136 __device__ static void hmac_sha256_pad (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[8], u32x opad[8])
 137 {
 138   w0[0] = w0[0] ^ 0x36363636;
 139   w0[1] = w0[1] ^ 0x36363636;
 140   w0[2] = w0[2] ^ 0x36363636;
 141   w0[3] = w0[3] ^ 0x36363636;
 142   w1[0] = w1[0] ^ 0x36363636;
 143   w1[1] = w1[1] ^ 0x36363636;
 144   w1[2] = w1[2] ^ 0x36363636;
 145   w1[3] = w1[3] ^ 0x36363636;
 146   w2[0] = w2[0] ^ 0x36363636;
 147   w2[1] = w2[1] ^ 0x36363636;
 148   w2[2] = w2[2] ^ 0x36363636;
 149   w2[3] = w2[3] ^ 0x36363636;
 150   w3[0] = w3[0] ^ 0x36363636;
 151   w3[1] = w3[1] ^ 0x36363636;
 152   w3[2] = w3[2] ^ 0x36363636;
 153   w3[3] = w3[3] ^ 0x36363636;
 154
 155   ipad[0] = SHA256M_A;
 156   ipad[1] = SHA256M_B;
 157   ipad[2] = SHA256M_C;
 158   ipad[3] = SHA256M_D;
 159   ipad[4] = SHA256M_E;
 160   ipad[5] = SHA256M_F;
 161   ipad[6] = SHA256M_G;
 162   ipad[7] = SHA256M_H;
 163
 164   sha256_transform (w0, w1, w2, w3, ipad);
 165
 166   w0[0] = w0[0] ^ 0x6a6a6a6a;
 167   w0[1] = w0[1] ^ 0x6a6a6a6a;
 168   w0[2] = w0[2] ^ 0x6a6a6a6a;
 169   w0[3] = w0[3] ^ 0x6a6a6a6a;
 170   w1[0] = w1[0] ^ 0x6a6a6a6a;
 171   w1[1] = w1[1] ^ 0x6a6a6a6a;
 172   w1[2] = w1[2] ^ 0x6a6a6a6a;
 173   w1[3] = w1[3] ^ 0x6a6a6a6a;
 174   w2[0] = w2[0] ^ 0x6a6a6a6a;
 175   w2[1] = w2[1] ^ 0x6a6a6a6a;
 176   w2[2] = w2[2] ^ 0x6a6a6a6a;
 177   w2[3] = w2[3] ^ 0x6a6a6a6a;
 178   w3[0] = w3[0] ^ 0x6a6a6a6a;
 179   w3[1] = w3[1] ^ 0x6a6a6a6a;
 180   w3[2] = w3[2] ^ 0x6a6a6a6a;
 181   w3[3] = w3[3] ^ 0x6a6a6a6a;
 182
 183   opad[0] = SHA256M_A;
 184   opad[1] = SHA256M_B;
 185   opad[2] = SHA256M_C;
 186   opad[3] = SHA256M_D;
 187   opad[4] = SHA256M_E;
 188   opad[5] = SHA256M_F;
 189   opad[6] = SHA256M_G;
 190   opad[7] = SHA256M_H;
 191
 192   sha256_transform (w0, w1, w2, w3, opad);
 193 }
 194
 195 __device__ static void hmac_sha256_run (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[8], u32x opad[8], u32x digest[8])
 196 {
 197   digest[0] = ipad[0];
 198   digest[1] = ipad[1];
 199   digest[2] = ipad[2];
 200   digest[3] = ipad[3];
 201   digest[4] = ipad[4];
 202   digest[5] = ipad[5];
 203   digest[6] = ipad[6];
 204   digest[7] = ipad[7];
 205
 206   sha256_transform (w0, w1, w2, w3, digest);
 207
 208   w0[0] = digest[0];
 209   w0[1] = digest[1];
 210   w0[2] = digest[2];
 211   w0[3] = digest[3];
 212   w1[0] = digest[4];
 213   w1[1] = digest[5];
 214   w1[2] = digest[6];
 215   w1[3] = digest[7];
 216   w2[0] = 0x80000000;
 217   w2[1] = 0;
 218   w2[2] = 0;
 219   w2[3] = 0;
 220   w3[0] = 0;
 221   w3[1] = 0;
 222   w3[2] = 0;
 223   w3[3] = (64 + 32) * 8;
 224
 225   digest[0] = opad[0];
 226   digest[1] = opad[1];
 227   digest[2] = opad[2];
 228   digest[3] = opad[3];
 229   digest[4] = opad[4];
 230   digest[5] = opad[5];
 231   digest[6] = opad[6];
 232   digest[7] = opad[7];
 233
 234   sha256_transform (w0, w1, w2, w3, digest);
 235 }
 236
 237 extern "C" __global__ void __launch_bounds__ (256, 1) m10900_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, pbkdf2_sha256_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 pbkdf2_sha256_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)
 238 {
 239   /**
 240    * base
 241    */
 242
 243   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 244
 245   if (gid >= gid_max) return;
 246
 247   u32x w0[4];
 248
 249   w0[0] = swap_workaround (pws[gid].i[ 0]);
 250   w0[1] = swap_workaround (pws[gid].i[ 1]);
 251   w0[2] = swap_workaround (pws[gid].i[ 2]);
 252   w0[3] = swap_workaround (pws[gid].i[ 3]);
 253
 254   u32x w1[4];
 255
 256   w1[0] = swap_workaround (pws[gid].i[ 4]);
 257   w1[1] = swap_workaround (pws[gid].i[ 5]);
 258   w1[2] = swap_workaround (pws[gid].i[ 6]);
 259   w1[3] = swap_workaround (pws[gid].i[ 7]);
 260
 261   u32x w2[4];
 262
 263   w2[0] = swap_workaround (pws[gid].i[ 8]);
 264   w2[1] = swap_workaround (pws[gid].i[ 9]);
 265   w2[2] = swap_workaround (pws[gid].i[10]);
 266   w2[3] = swap_workaround (pws[gid].i[11]);
 267
 268   u32x w3[4];
 269
 270   w3[0] = swap_workaround (pws[gid].i[12]);
 271   w3[1] = swap_workaround (pws[gid].i[13]);
 272   w3[2] = swap_workaround (pws[gid].i[14]);
 273   w3[3] = swap_workaround (pws[gid].i[15]);
 274
 275   /**
 276    * salt
 277    */
 278
 279   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 280
 281   u32 esalt_buf0[4];
 282   u32 esalt_buf1[4];
 283   u32 esalt_buf2[4];
 284   u32 esalt_buf3[4];
 285
 286   esalt_buf0[0] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 0]);
 287   esalt_buf0[1] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 1]);
 288   esalt_buf0[2] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 2]);
 289   esalt_buf0[3] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 3]);
 290   esalt_buf1[0] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 4]);
 291   esalt_buf1[1] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 5]);
 292   esalt_buf1[2] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 6]);
 293   esalt_buf1[3] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 7]);
 294   esalt_buf2[0] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 8]);
 295   esalt_buf2[1] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 9]);
 296   esalt_buf2[2] = swap_workaround (esalt_bufs[salt_pos].salt_buf[10]);
 297   esalt_buf2[3] = swap_workaround (esalt_bufs[salt_pos].salt_buf[11]);
 298   esalt_buf3[0] = swap_workaround (esalt_bufs[salt_pos].salt_buf[12]);
 299   esalt_buf3[1] = swap_workaround (esalt_bufs[salt_pos].salt_buf[13]);
 300   esalt_buf3[2] = 0;
 301   esalt_buf3[3] = (64 + salt_len + 4) * 8;
 302
 303   u32 ipad[8];
 304   u32 opad[8];
 305
 306   hmac_sha256_pad (w0, w1, w2, w3, ipad, opad);
 307
 308   tmps[gid].ipad[0] = ipad[0];
 309   tmps[gid].ipad[1] = ipad[1];
 310   tmps[gid].ipad[2] = ipad[2];
 311   tmps[gid].ipad[3] = ipad[3];
 312   tmps[gid].ipad[4] = ipad[4];
 313   tmps[gid].ipad[5] = ipad[5];
 314   tmps[gid].ipad[6] = ipad[6];
 315   tmps[gid].ipad[7] = ipad[7];
 316
 317   tmps[gid].opad[0] = opad[0];
 318   tmps[gid].opad[1] = opad[1];
 319   tmps[gid].opad[2] = opad[2];
 320   tmps[gid].opad[3] = opad[3];
 321   tmps[gid].opad[4] = opad[4];
 322   tmps[gid].opad[5] = opad[5];
 323   tmps[gid].opad[6] = opad[6];
 324   tmps[gid].opad[7] = opad[7];
 325
 326   for (u32 i = 0, j = 1; i < 8; i += 8, j += 1)
 327   {
 328     u32 dgst[8];
 329
 330     hmac_sha256_run (esalt_buf0, esalt_buf1, esalt_buf2, esalt_buf3, ipad, opad, dgst);
 331
 332     tmps[gid].dgst[i + 0] = dgst[0];
 333     tmps[gid].dgst[i + 1] = dgst[1];
 334     tmps[gid].dgst[i + 2] = dgst[2];
 335     tmps[gid].dgst[i + 3] = dgst[3];
 336     tmps[gid].dgst[i + 4] = dgst[4];
 337     tmps[gid].dgst[i + 5] = dgst[5];
 338     tmps[gid].dgst[i + 6] = dgst[6];
 339     tmps[gid].dgst[i + 7] = dgst[7];
 340
 341     tmps[gid].out[i + 0] = dgst[0];
 342     tmps[gid].out[i + 1] = dgst[1];
 343     tmps[gid].out[i + 2] = dgst[2];
 344     tmps[gid].out[i + 3] = dgst[3];
 345     tmps[gid].out[i + 4] = dgst[4];
 346     tmps[gid].out[i + 5] = dgst[5];
 347     tmps[gid].out[i + 6] = dgst[6];
 348     tmps[gid].out[i + 7] = dgst[7];
 349   }
 350 }
 351
 352 extern "C" __global__ void __launch_bounds__ (256, 1) m10900_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, pbkdf2_sha256_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 pbkdf2_sha256_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)
 353 {
 354   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 355
 356   if (gid >= gid_max) return;
 357
 358   u32 ipad[8];
 359
 360   ipad[0] = tmps[gid].ipad[0];
 361   ipad[1] = tmps[gid].ipad[1];
 362   ipad[2] = tmps[gid].ipad[2];
 363   ipad[3] = tmps[gid].ipad[3];
 364   ipad[4] = tmps[gid].ipad[4];
 365   ipad[5] = tmps[gid].ipad[5];
 366   ipad[6] = tmps[gid].ipad[6];
 367   ipad[7] = tmps[gid].ipad[7];
 368
 369   u32 opad[8];
 370
 371   opad[0] = tmps[gid].opad[0];
 372   opad[1] = tmps[gid].opad[1];
 373   opad[2] = tmps[gid].opad[2];
 374   opad[3] = tmps[gid].opad[3];
 375   opad[4] = tmps[gid].opad[4];
 376   opad[5] = tmps[gid].opad[5];
 377   opad[6] = tmps[gid].opad[6];
 378   opad[7] = tmps[gid].opad[7];
 379
 380   for (u32 i = 0; i < 8; i += 8)
 381   {
 382     u32 dgst[8];
 383
 384     dgst[0] = tmps[gid].dgst[i + 0];
 385     dgst[1] = tmps[gid].dgst[i + 1];
 386     dgst[2] = tmps[gid].dgst[i + 2];
 387     dgst[3] = tmps[gid].dgst[i + 3];
 388     dgst[4] = tmps[gid].dgst[i + 4];
 389     dgst[5] = tmps[gid].dgst[i + 5];
 390     dgst[6] = tmps[gid].dgst[i + 6];
 391     dgst[7] = tmps[gid].dgst[i + 7];
 392
 393     u32 out[8];
 394
 395     out[0] = tmps[gid].out[i + 0];
 396     out[1] = tmps[gid].out[i + 1];
 397     out[2] = tmps[gid].out[i + 2];
 398     out[3] = tmps[gid].out[i + 3];
 399     out[4] = tmps[gid].out[i + 4];
 400     out[5] = tmps[gid].out[i + 5];
 401     out[6] = tmps[gid].out[i + 6];
 402     out[7] = tmps[gid].out[i + 7];
 403
 404     for (u32 j = 0; j < loop_cnt; j++)
 405     {
 406       u32 w0[4];
 407       u32 w1[4];
 408       u32 w2[4];
 409       u32 w3[4];
 410
 411       w0[0] = dgst[0];
 412       w0[1] = dgst[1];
 413       w0[2] = dgst[2];
 414       w0[3] = dgst[3];
 415       w1[0] = dgst[4];
 416       w1[1] = dgst[5];
 417       w1[2] = dgst[6];
 418       w1[3] = dgst[7];
 419       w2[0] = 0x80000000;
 420       w2[1] = 0;
 421       w2[2] = 0;
 422       w2[3] = 0;
 423       w3[0] = 0;
 424       w3[1] = 0;
 425       w3[2] = 0;
 426       w3[3] = (64 + 32) * 8;
 427
 428       hmac_sha256_run (w0, w1, w2, w3, ipad, opad, dgst);
 429
 430       out[0] ^= dgst[0];
 431       out[1] ^= dgst[1];
 432       out[2] ^= dgst[2];
 433       out[3] ^= dgst[3];
 434       out[4] ^= dgst[4];
 435       out[5] ^= dgst[5];
 436       out[6] ^= dgst[6];
 437       out[7] ^= dgst[7];
 438     }
 439
 440     tmps[gid].dgst[i + 0] = dgst[0];
 441     tmps[gid].dgst[i + 1] = dgst[1];
 442     tmps[gid].dgst[i + 2] = dgst[2];
 443     tmps[gid].dgst[i + 3] = dgst[3];
 444     tmps[gid].dgst[i + 4] = dgst[4];
 445     tmps[gid].dgst[i + 5] = dgst[5];
 446     tmps[gid].dgst[i + 6] = dgst[6];
 447     tmps[gid].dgst[i + 7] = dgst[7];
 448
 449     tmps[gid].out[i + 0] = out[0];
 450     tmps[gid].out[i + 1] = out[1];
 451     tmps[gid].out[i + 2] = out[2];
 452     tmps[gid].out[i + 3] = out[3];
 453     tmps[gid].out[i + 4] = out[4];
 454     tmps[gid].out[i + 5] = out[5];
 455     tmps[gid].out[i + 6] = out[6];
 456     tmps[gid].out[i + 7] = out[7];
 457   }
 458 }
 459
 460 extern "C" __global__ void __launch_bounds__ (256, 1) m10900_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, pbkdf2_sha256_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 pbkdf2_sha256_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)
 461 {
 462   /**
 463    * base
 464    */
 465
 466   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 467
 468   if (gid >= gid_max) return;
 469
 470   const u32 lid = threadIdx.x;
 471
 472   const u32x r0 = tmps[gid].out[DGST_R0];
 473   const u32x r1 = tmps[gid].out[DGST_R1];
 474   const u32x r2 = tmps[gid].out[DGST_R2];
 475   const u32x r3 = tmps[gid].out[DGST_R3];
 476
 477   #define il_pos 0
 478
 479   #include VECT_COMPARE_M
 480 }