nv/m12800.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _MS_DRSR_
   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 #ifdef VECT_SIZE1
  34 #define uint_to_hex_lower8(i) l_bin2asc[(i)]
  35 #endif
  36
  37 __device__ __constant__ char c_bin2asc[16] = { '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', 'A', 'B', 'C', 'D', 'E', 'F' };
  38
  39 __device__ __shared__ short l_bin2asc[256];
  40
  41 __device__ static void md4_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
  42 {
  43   u32x a = digest[0];
  44   u32x b = digest[1];
  45   u32x c = digest[2];
  46   u32x d = digest[3];
  47
  48   MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
  49   MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
  50   MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
  51   MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
  52   MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
  53   MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
  54   MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
  55   MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
  56   MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
  57   MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
  58   MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
  59   MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
  60   MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
  61   MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
  62   MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
  63   MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
  64
  65   MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
  66   MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
  67   MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
  68   MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
  69   MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
  70   MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
  71   MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
  72   MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
  73   MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
  74   MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
  75   MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
  76   MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
  77   MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
  78   MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
  79   MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
  80   MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
  81
  82   MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
  83   MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
  84   MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
  85   MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
  86   MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
  87   MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
  88   MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
  89   MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
  90   MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
  91   MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
  92   MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
  93   MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
  94   MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
  95   MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
  96   MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
  97   MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
  98
  99   digest[0] += a;
 100   digest[1] += b;
 101   digest[2] += c;
 102   digest[3] += d;
 103 }
 104
 105 __device__ static void sha256_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[8])
 106 {
 107   u32x a = digest[0];
 108   u32x b = digest[1];
 109   u32x c = digest[2];
 110   u32x d = digest[3];
 111   u32x e = digest[4];
 112   u32x f = digest[5];
 113   u32x g = digest[6];
 114   u32x h = digest[7];
 115
 116   u32x w0_t = w0[0];
 117   u32x w1_t = w0[1];
 118   u32x w2_t = w0[2];
 119   u32x w3_t = w0[3];
 120   u32x w4_t = w1[0];
 121   u32x w5_t = w1[1];
 122   u32x w6_t = w1[2];
 123   u32x w7_t = w1[3];
 124   u32x w8_t = w2[0];
 125   u32x w9_t = w2[1];
 126   u32x wa_t = w2[2];
 127   u32x wb_t = w2[3];
 128   u32x wc_t = w3[0];
 129   u32x wd_t = w3[1];
 130   u32x we_t = w3[2];
 131   u32x wf_t = w3[3];
 132
 133   SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C00);
 134   SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C01);
 135   SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C02);
 136   SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C03);
 137   SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C04);
 138   SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C05);
 139   SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C06);
 140   SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C07);
 141   SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C08);
 142   SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C09);
 143   SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C0a);
 144   SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C0b);
 145   SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C0c);
 146   SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C0d);
 147   SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C0e);
 148   SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C0f);
 149   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);
 150   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);
 151   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);
 152   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);
 153   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);
 154   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);
 155   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);
 156   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);
 157   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);
 158   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);
 159   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);
 160   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);
 161   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);
 162   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);
 163   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);
 164   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);
 165   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);
 166   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);
 167   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);
 168   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);
 169   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);
 170   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);
 171   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);
 172   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);
 173   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);
 174   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);
 175   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);
 176   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);
 177   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);
 178   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);
 179   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);
 180   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);
 181   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);
 182   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);
 183   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);
 184   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);
 185   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);
 186   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);
 187   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);
 188   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);
 189   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);
 190   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);
 191   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);
 192   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);
 193   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);
 194   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);
 195   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);
 196   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);
 197
 198   digest[0] += a;
 199   digest[1] += b;
 200   digest[2] += c;
 201   digest[3] += d;
 202   digest[4] += e;
 203   digest[5] += f;
 204   digest[6] += g;
 205   digest[7] += h;
 206 }
 207
 208 __device__ static void hmac_sha256_pad (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[8], u32x opad[8])
 209 {
 210   w0[0] = w0[0] ^ 0x36363636;
 211   w0[1] = w0[1] ^ 0x36363636;
 212   w0[2] = w0[2] ^ 0x36363636;
 213   w0[3] = w0[3] ^ 0x36363636;
 214   w1[0] = w1[0] ^ 0x36363636;
 215   w1[1] = w1[1] ^ 0x36363636;
 216   w1[2] = w1[2] ^ 0x36363636;
 217   w1[3] = w1[3] ^ 0x36363636;
 218   w2[0] = w2[0] ^ 0x36363636;
 219   w2[1] = w2[1] ^ 0x36363636;
 220   w2[2] = w2[2] ^ 0x36363636;
 221   w2[3] = w2[3] ^ 0x36363636;
 222   w3[0] = w3[0] ^ 0x36363636;
 223   w3[1] = w3[1] ^ 0x36363636;
 224   w3[2] = w3[2] ^ 0x36363636;
 225   w3[3] = w3[3] ^ 0x36363636;
 226
 227   ipad[0] = SHA256M_A;
 228   ipad[1] = SHA256M_B;
 229   ipad[2] = SHA256M_C;
 230   ipad[3] = SHA256M_D;
 231   ipad[4] = SHA256M_E;
 232   ipad[5] = SHA256M_F;
 233   ipad[6] = SHA256M_G;
 234   ipad[7] = SHA256M_H;
 235
 236   sha256_transform (w0, w1, w2, w3, ipad);
 237
 238   w0[0] = w0[0] ^ 0x6a6a6a6a;
 239   w0[1] = w0[1] ^ 0x6a6a6a6a;
 240   w0[2] = w0[2] ^ 0x6a6a6a6a;
 241   w0[3] = w0[3] ^ 0x6a6a6a6a;
 242   w1[0] = w1[0] ^ 0x6a6a6a6a;
 243   w1[1] = w1[1] ^ 0x6a6a6a6a;
 244   w1[2] = w1[2] ^ 0x6a6a6a6a;
 245   w1[3] = w1[3] ^ 0x6a6a6a6a;
 246   w2[0] = w2[0] ^ 0x6a6a6a6a;
 247   w2[1] = w2[1] ^ 0x6a6a6a6a;
 248   w2[2] = w2[2] ^ 0x6a6a6a6a;
 249   w2[3] = w2[3] ^ 0x6a6a6a6a;
 250   w3[0] = w3[0] ^ 0x6a6a6a6a;
 251   w3[1] = w3[1] ^ 0x6a6a6a6a;
 252   w3[2] = w3[2] ^ 0x6a6a6a6a;
 253   w3[3] = w3[3] ^ 0x6a6a6a6a;
 254
 255   opad[0] = SHA256M_A;
 256   opad[1] = SHA256M_B;
 257   opad[2] = SHA256M_C;
 258   opad[3] = SHA256M_D;
 259   opad[4] = SHA256M_E;
 260   opad[5] = SHA256M_F;
 261   opad[6] = SHA256M_G;
 262   opad[7] = SHA256M_H;
 263
 264   sha256_transform (w0, w1, w2, w3, opad);
 265 }
 266
 267 __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])
 268 {
 269   digest[0] = ipad[0];
 270   digest[1] = ipad[1];
 271   digest[2] = ipad[2];
 272   digest[3] = ipad[3];
 273   digest[4] = ipad[4];
 274   digest[5] = ipad[5];
 275   digest[6] = ipad[6];
 276   digest[7] = ipad[7];
 277
 278   sha256_transform (w0, w1, w2, w3, digest);
 279
 280   w0[0] = digest[0];
 281   w0[1] = digest[1];
 282   w0[2] = digest[2];
 283   w0[3] = digest[3];
 284   w1[0] = digest[4];
 285   w1[1] = digest[5];
 286   w1[2] = digest[6];
 287   w1[3] = digest[7];
 288   w2[0] = 0x80000000;
 289   w2[1] = 0;
 290   w2[2] = 0;
 291   w2[3] = 0;
 292   w3[0] = 0;
 293   w3[1] = 0;
 294   w3[2] = 0;
 295   w3[3] = (64 + 32) * 8;
 296
 297   digest[0] = opad[0];
 298   digest[1] = opad[1];
 299   digest[2] = opad[2];
 300   digest[3] = opad[3];
 301   digest[4] = opad[4];
 302   digest[5] = opad[5];
 303   digest[6] = opad[6];
 304   digest[7] = opad[7];
 305
 306   sha256_transform (w0, w1, w2, w3, digest);
 307 }
 308
 309 extern "C" __global__ void __launch_bounds__ (256, 1) m12800_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)
 310 {
 311   /**
 312    * base
 313    */
 314
 315   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 316   const u32 lid = threadIdx.x;
 317
 318   /**
 319    * lookup ascii table
 320    */
 321
 322   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
 323                  | c_bin2asc[(lid >> 4) & 15] << 0;
 324
 325   __syncthreads ();
 326
 327   if (gid >= gid_max) return;
 328
 329   /**
 330    * base
 331    */
 332
 333   u32x w0[4];
 334
 335   w0[0] = pws[gid].i[0];
 336   w0[1] = pws[gid].i[1];
 337   w0[2] = pws[gid].i[2];
 338   w0[3] = pws[gid].i[3];
 339
 340   u32x w1[4];
 341
 342   w1[0] = pws[gid].i[4];
 343   w1[1] = pws[gid].i[5];
 344   w1[2] = 0;
 345   w1[3] = 0;
 346
 347   u32x w2[4];
 348
 349   w2[0] = 0;
 350   w2[1] = 0;
 351   w2[2] = 0;
 352   w2[3] = 0;
 353
 354   u32x w3[4];
 355
 356   w3[0] = 0;
 357   w3[1] = 0;
 358   w3[2] = 0;
 359   w3[3] = 0;
 360
 361   const u32 pw_len = pws[gid].pw_len;
 362
 363   /**
 364    * salt
 365    */
 366
 367   u32 salt_len = salt_bufs[salt_pos].salt_len;
 368
 369   u32 salt_buf0[4];
 370
 371   salt_buf0[0] = swap_workaround (salt_bufs[salt_pos].salt_buf[0]);
 372   salt_buf0[1] = swap_workaround (salt_bufs[salt_pos].salt_buf[1]);
 373   salt_buf0[2] = swap_workaround (salt_bufs[salt_pos].salt_buf[2]);
 374   salt_buf0[3] = swap_workaround (salt_bufs[salt_pos].salt_buf[3]);
 375
 376   u32 salt_buf1[4];
 377
 378   salt_buf1[0] = 0;
 379   salt_buf1[1] = 0;
 380   salt_buf1[2] = 0;
 381   salt_buf1[3] = 0;
 382
 383   u32 salt_buf2[4];
 384
 385   salt_buf2[0] = 0;
 386   salt_buf2[1] = 0;
 387   salt_buf2[2] = 0;
 388   salt_buf2[3] = 0;
 389
 390   u32 salt_buf3[4];
 391
 392   salt_buf3[0] = 0;
 393   salt_buf3[1] = 0;
 394   salt_buf3[2] = 0;
 395   salt_buf3[3] = (64 + salt_len + 4) * 8;
 396
 397   /**
 398    * generate nthash
 399    */
 400
 401   append_0x80_2 (w0, w1, pw_len);
 402
 403   make_unicode (w1, w2, w3);
 404   make_unicode (w0, w0, w1);
 405
 406   w3[2] = pw_len * 2 * 8;
 407
 408   u32x digest_md4[4];
 409
 410   digest_md4[0] = MD4M_A;
 411   digest_md4[1] = MD4M_B;
 412   digest_md4[2] = MD4M_C;
 413   digest_md4[3] = MD4M_D;
 414
 415   md4_transform (w0, w1, w2, w3, digest_md4);
 416
 417   w0[0] = uint_to_hex_lower8 ((digest_md4[0] >>  0) & 255) <<  0
 418         | uint_to_hex_lower8 ((digest_md4[0] >>  8) & 255) << 16;
 419   w0[1] = uint_to_hex_lower8 ((digest_md4[0] >> 16) & 255) <<  0
 420         | uint_to_hex_lower8 ((digest_md4[0] >> 24) & 255) << 16;
 421   w0[2] = uint_to_hex_lower8 ((digest_md4[1] >>  0) & 255) <<  0
 422         | uint_to_hex_lower8 ((digest_md4[1] >>  8) & 255) << 16;
 423   w0[3] = uint_to_hex_lower8 ((digest_md4[1] >> 16) & 255) <<  0
 424         | uint_to_hex_lower8 ((digest_md4[1] >> 24) & 255) << 16;
 425   w1[0] = uint_to_hex_lower8 ((digest_md4[2] >>  0) & 255) <<  0
 426         | uint_to_hex_lower8 ((digest_md4[2] >>  8) & 255) << 16;
 427   w1[1] = uint_to_hex_lower8 ((digest_md4[2] >> 16) & 255) <<  0
 428         | uint_to_hex_lower8 ((digest_md4[2] >> 24) & 255) << 16;
 429   w1[2] = uint_to_hex_lower8 ((digest_md4[3] >>  0) & 255) <<  0
 430         | uint_to_hex_lower8 ((digest_md4[3] >>  8) & 255) << 16;
 431   w1[3] = uint_to_hex_lower8 ((digest_md4[3] >> 16) & 255) <<  0
 432         | uint_to_hex_lower8 ((digest_md4[3] >> 24) & 255) << 16;
 433
 434   make_unicode (w1, w2, w3);
 435   make_unicode (w0, w0, w1);
 436
 437   w0[0] = swap_workaround (w0[0]);
 438   w0[1] = swap_workaround (w0[1]);
 439   w0[2] = swap_workaround (w0[2]);
 440   w0[3] = swap_workaround (w0[3]);
 441   w1[0] = swap_workaround (w1[0]);
 442   w1[1] = swap_workaround (w1[1]);
 443   w1[2] = swap_workaround (w1[2]);
 444   w1[3] = swap_workaround (w1[3]);
 445   w2[0] = swap_workaround (w2[0]);
 446   w2[1] = swap_workaround (w2[1]);
 447   w2[2] = swap_workaround (w2[2]);
 448   w2[3] = swap_workaround (w2[3]);
 449   w3[0] = swap_workaround (w3[0]);
 450   w3[1] = swap_workaround (w3[1]);
 451   w3[2] = swap_workaround (w3[2]);
 452   w3[3] = swap_workaround (w3[3]);
 453
 454   u32 ipad[8];
 455   u32 opad[8];
 456
 457   hmac_sha256_pad (w0, w1, w2, w3, ipad, opad);
 458
 459   tmps[gid].ipad[0] = ipad[0];
 460   tmps[gid].ipad[1] = ipad[1];
 461   tmps[gid].ipad[2] = ipad[2];
 462   tmps[gid].ipad[3] = ipad[3];
 463   tmps[gid].ipad[4] = ipad[4];
 464   tmps[gid].ipad[5] = ipad[5];
 465   tmps[gid].ipad[6] = ipad[6];
 466   tmps[gid].ipad[7] = ipad[7];
 467
 468   tmps[gid].opad[0] = opad[0];
 469   tmps[gid].opad[1] = opad[1];
 470   tmps[gid].opad[2] = opad[2];
 471   tmps[gid].opad[3] = opad[3];
 472   tmps[gid].opad[4] = opad[4];
 473   tmps[gid].opad[5] = opad[5];
 474   tmps[gid].opad[6] = opad[6];
 475   tmps[gid].opad[7] = opad[7];
 476
 477   for (u32 i = 0, j = 1; i < 8; i += 8, j += 1)
 478   {
 479     u32 dgst[8];
 480
 481     hmac_sha256_run (salt_buf0, salt_buf1, salt_buf2, salt_buf3, ipad, opad, dgst);
 482
 483     tmps[gid].dgst[i + 0] = dgst[0];
 484     tmps[gid].dgst[i + 1] = dgst[1];
 485     tmps[gid].dgst[i + 2] = dgst[2];
 486     tmps[gid].dgst[i + 3] = dgst[3];
 487     tmps[gid].dgst[i + 4] = dgst[4];
 488     tmps[gid].dgst[i + 5] = dgst[5];
 489     tmps[gid].dgst[i + 6] = dgst[6];
 490     tmps[gid].dgst[i + 7] = dgst[7];
 491
 492     tmps[gid].out[i + 0] = dgst[0];
 493     tmps[gid].out[i + 1] = dgst[1];
 494     tmps[gid].out[i + 2] = dgst[2];
 495     tmps[gid].out[i + 3] = dgst[3];
 496     tmps[gid].out[i + 4] = dgst[4];
 497     tmps[gid].out[i + 5] = dgst[5];
 498     tmps[gid].out[i + 6] = dgst[6];
 499     tmps[gid].out[i + 7] = dgst[7];
 500   }
 501 }
 502
 503 extern "C" __global__ void __launch_bounds__ (256, 1) m12800_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)
 504 {
 505   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 506
 507   if (gid >= gid_max) return;
 508
 509   u32 ipad[8];
 510
 511   ipad[0] = tmps[gid].ipad[0];
 512   ipad[1] = tmps[gid].ipad[1];
 513   ipad[2] = tmps[gid].ipad[2];
 514   ipad[3] = tmps[gid].ipad[3];
 515   ipad[4] = tmps[gid].ipad[4];
 516   ipad[5] = tmps[gid].ipad[5];
 517   ipad[6] = tmps[gid].ipad[6];
 518   ipad[7] = tmps[gid].ipad[7];
 519
 520   u32 opad[8];
 521
 522   opad[0] = tmps[gid].opad[0];
 523   opad[1] = tmps[gid].opad[1];
 524   opad[2] = tmps[gid].opad[2];
 525   opad[3] = tmps[gid].opad[3];
 526   opad[4] = tmps[gid].opad[4];
 527   opad[5] = tmps[gid].opad[5];
 528   opad[6] = tmps[gid].opad[6];
 529   opad[7] = tmps[gid].opad[7];
 530
 531   for (u32 i = 0; i < 8; i += 8)
 532   {
 533     u32 dgst[8];
 534
 535     dgst[0] = tmps[gid].dgst[i + 0];
 536     dgst[1] = tmps[gid].dgst[i + 1];
 537     dgst[2] = tmps[gid].dgst[i + 2];
 538     dgst[3] = tmps[gid].dgst[i + 3];
 539     dgst[4] = tmps[gid].dgst[i + 4];
 540     dgst[5] = tmps[gid].dgst[i + 5];
 541     dgst[6] = tmps[gid].dgst[i + 6];
 542     dgst[7] = tmps[gid].dgst[i + 7];
 543
 544     u32 out[8];
 545
 546     out[0] = tmps[gid].out[i + 0];
 547     out[1] = tmps[gid].out[i + 1];
 548     out[2] = tmps[gid].out[i + 2];
 549     out[3] = tmps[gid].out[i + 3];
 550     out[4] = tmps[gid].out[i + 4];
 551     out[5] = tmps[gid].out[i + 5];
 552     out[6] = tmps[gid].out[i + 6];
 553     out[7] = tmps[gid].out[i + 7];
 554
 555     for (u32 j = 0; j < loop_cnt; j++)
 556     {
 557       u32 w0[4];
 558       u32 w1[4];
 559       u32 w2[4];
 560       u32 w3[4];
 561
 562       w0[0] = dgst[0];
 563       w0[1] = dgst[1];
 564       w0[2] = dgst[2];
 565       w0[3] = dgst[3];
 566       w1[0] = dgst[4];
 567       w1[1] = dgst[5];
 568       w1[2] = dgst[6];
 569       w1[3] = dgst[7];
 570       w2[0] = 0x80000000;
 571       w2[1] = 0;
 572       w2[2] = 0;
 573       w2[3] = 0;
 574       w3[0] = 0;
 575       w3[1] = 0;
 576       w3[2] = 0;
 577       w3[3] = (64 + 32) * 8;
 578
 579       hmac_sha256_run (w0, w1, w2, w3, ipad, opad, dgst);
 580
 581       out[0] ^= dgst[0];
 582       out[1] ^= dgst[1];
 583       out[2] ^= dgst[2];
 584       out[3] ^= dgst[3];
 585       out[4] ^= dgst[4];
 586       out[5] ^= dgst[5];
 587       out[6] ^= dgst[6];
 588       out[7] ^= dgst[7];
 589     }
 590
 591     tmps[gid].dgst[i + 0] = dgst[0];
 592     tmps[gid].dgst[i + 1] = dgst[1];
 593     tmps[gid].dgst[i + 2] = dgst[2];
 594     tmps[gid].dgst[i + 3] = dgst[3];
 595     tmps[gid].dgst[i + 4] = dgst[4];
 596     tmps[gid].dgst[i + 5] = dgst[5];
 597     tmps[gid].dgst[i + 6] = dgst[6];
 598     tmps[gid].dgst[i + 7] = dgst[7];
 599
 600     tmps[gid].out[i + 0] = out[0];
 601     tmps[gid].out[i + 1] = out[1];
 602     tmps[gid].out[i + 2] = out[2];
 603     tmps[gid].out[i + 3] = out[3];
 604     tmps[gid].out[i + 4] = out[4];
 605     tmps[gid].out[i + 5] = out[5];
 606     tmps[gid].out[i + 6] = out[6];
 607     tmps[gid].out[i + 7] = out[7];
 608   }
 609 }
 610
 611 extern "C" __global__ void __launch_bounds__ (256, 1) m12800_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)
 612 {
 613   /**
 614    * base
 615    */
 616
 617   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 618
 619   if (gid >= gid_max) return;
 620
 621   const u32 lid = threadIdx.x;
 622
 623   const u32x r0 = tmps[gid].out[DGST_R0];
 624   const u32x r1 = tmps[gid].out[DGST_R1];
 625   const u32x r2 = tmps[gid].out[DGST_R2];
 626   const u32x r3 = tmps[gid].out[DGST_R3];
 627
 628   #define il_pos 0
 629
 630   #include VECT_COMPARE_M
 631 }