nv/m06213.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _RIPEMD160_
   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__ static void ripemd160_transform (const u32x w[16], u32x dgst[5])
  33 {
  34   u32x a1 = dgst[0];
  35   u32x b1 = dgst[1];
  36   u32x c1 = dgst[2];
  37   u32x d1 = dgst[3];
  38   u32x e1 = dgst[4];
  39
  40   RIPEMD160_STEP (RIPEMD160_F , a1, b1, c1, d1, e1, w[ 0], RIPEMD160C00, RIPEMD160S00);
  41   RIPEMD160_STEP (RIPEMD160_F , e1, a1, b1, c1, d1, w[ 1], RIPEMD160C00, RIPEMD160S01);
  42   RIPEMD160_STEP (RIPEMD160_F , d1, e1, a1, b1, c1, w[ 2], RIPEMD160C00, RIPEMD160S02);
  43   RIPEMD160_STEP (RIPEMD160_F , c1, d1, e1, a1, b1, w[ 3], RIPEMD160C00, RIPEMD160S03);
  44   RIPEMD160_STEP (RIPEMD160_F , b1, c1, d1, e1, a1, w[ 4], RIPEMD160C00, RIPEMD160S04);
  45   RIPEMD160_STEP (RIPEMD160_F , a1, b1, c1, d1, e1, w[ 5], RIPEMD160C00, RIPEMD160S05);
  46   RIPEMD160_STEP (RIPEMD160_F , e1, a1, b1, c1, d1, w[ 6], RIPEMD160C00, RIPEMD160S06);
  47   RIPEMD160_STEP (RIPEMD160_F , d1, e1, a1, b1, c1, w[ 7], RIPEMD160C00, RIPEMD160S07);
  48   RIPEMD160_STEP (RIPEMD160_F , c1, d1, e1, a1, b1, w[ 8], RIPEMD160C00, RIPEMD160S08);
  49   RIPEMD160_STEP (RIPEMD160_F , b1, c1, d1, e1, a1, w[ 9], RIPEMD160C00, RIPEMD160S09);
  50   RIPEMD160_STEP (RIPEMD160_F , a1, b1, c1, d1, e1, w[10], RIPEMD160C00, RIPEMD160S0A);
  51   RIPEMD160_STEP (RIPEMD160_F , e1, a1, b1, c1, d1, w[11], RIPEMD160C00, RIPEMD160S0B);
  52   RIPEMD160_STEP (RIPEMD160_F , d1, e1, a1, b1, c1, w[12], RIPEMD160C00, RIPEMD160S0C);
  53   RIPEMD160_STEP (RIPEMD160_F , c1, d1, e1, a1, b1, w[13], RIPEMD160C00, RIPEMD160S0D);
  54   RIPEMD160_STEP (RIPEMD160_F , b1, c1, d1, e1, a1, w[14], RIPEMD160C00, RIPEMD160S0E);
  55   RIPEMD160_STEP (RIPEMD160_F , a1, b1, c1, d1, e1, w[15], RIPEMD160C00, RIPEMD160S0F);
  56
  57   RIPEMD160_STEP (RIPEMD160_Go, e1, a1, b1, c1, d1, w[ 7], RIPEMD160C10, RIPEMD160S10);
  58   RIPEMD160_STEP (RIPEMD160_Go, d1, e1, a1, b1, c1, w[ 4], RIPEMD160C10, RIPEMD160S11);
  59   RIPEMD160_STEP (RIPEMD160_Go, c1, d1, e1, a1, b1, w[13], RIPEMD160C10, RIPEMD160S12);
  60   RIPEMD160_STEP (RIPEMD160_Go, b1, c1, d1, e1, a1, w[ 1], RIPEMD160C10, RIPEMD160S13);
  61   RIPEMD160_STEP (RIPEMD160_Go, a1, b1, c1, d1, e1, w[10], RIPEMD160C10, RIPEMD160S14);
  62   RIPEMD160_STEP (RIPEMD160_Go, e1, a1, b1, c1, d1, w[ 6], RIPEMD160C10, RIPEMD160S15);
  63   RIPEMD160_STEP (RIPEMD160_Go, d1, e1, a1, b1, c1, w[15], RIPEMD160C10, RIPEMD160S16);
  64   RIPEMD160_STEP (RIPEMD160_Go, c1, d1, e1, a1, b1, w[ 3], RIPEMD160C10, RIPEMD160S17);
  65   RIPEMD160_STEP (RIPEMD160_Go, b1, c1, d1, e1, a1, w[12], RIPEMD160C10, RIPEMD160S18);
  66   RIPEMD160_STEP (RIPEMD160_Go, a1, b1, c1, d1, e1, w[ 0], RIPEMD160C10, RIPEMD160S19);
  67   RIPEMD160_STEP (RIPEMD160_Go, e1, a1, b1, c1, d1, w[ 9], RIPEMD160C10, RIPEMD160S1A);
  68   RIPEMD160_STEP (RIPEMD160_Go, d1, e1, a1, b1, c1, w[ 5], RIPEMD160C10, RIPEMD160S1B);
  69   RIPEMD160_STEP (RIPEMD160_Go, c1, d1, e1, a1, b1, w[ 2], RIPEMD160C10, RIPEMD160S1C);
  70   RIPEMD160_STEP (RIPEMD160_Go, b1, c1, d1, e1, a1, w[14], RIPEMD160C10, RIPEMD160S1D);
  71   RIPEMD160_STEP (RIPEMD160_Go, a1, b1, c1, d1, e1, w[11], RIPEMD160C10, RIPEMD160S1E);
  72   RIPEMD160_STEP (RIPEMD160_Go, e1, a1, b1, c1, d1, w[ 8], RIPEMD160C10, RIPEMD160S1F);
  73
  74   RIPEMD160_STEP (RIPEMD160_H , d1, e1, a1, b1, c1, w[ 3], RIPEMD160C20, RIPEMD160S20);
  75   RIPEMD160_STEP (RIPEMD160_H , c1, d1, e1, a1, b1, w[10], RIPEMD160C20, RIPEMD160S21);
  76   RIPEMD160_STEP (RIPEMD160_H , b1, c1, d1, e1, a1, w[14], RIPEMD160C20, RIPEMD160S22);
  77   RIPEMD160_STEP (RIPEMD160_H , a1, b1, c1, d1, e1, w[ 4], RIPEMD160C20, RIPEMD160S23);
  78   RIPEMD160_STEP (RIPEMD160_H , e1, a1, b1, c1, d1, w[ 9], RIPEMD160C20, RIPEMD160S24);
  79   RIPEMD160_STEP (RIPEMD160_H , d1, e1, a1, b1, c1, w[15], RIPEMD160C20, RIPEMD160S25);
  80   RIPEMD160_STEP (RIPEMD160_H , c1, d1, e1, a1, b1, w[ 8], RIPEMD160C20, RIPEMD160S26);
  81   RIPEMD160_STEP (RIPEMD160_H , b1, c1, d1, e1, a1, w[ 1], RIPEMD160C20, RIPEMD160S27);
  82   RIPEMD160_STEP (RIPEMD160_H , a1, b1, c1, d1, e1, w[ 2], RIPEMD160C20, RIPEMD160S28);
  83   RIPEMD160_STEP (RIPEMD160_H , e1, a1, b1, c1, d1, w[ 7], RIPEMD160C20, RIPEMD160S29);
  84   RIPEMD160_STEP (RIPEMD160_H , d1, e1, a1, b1, c1, w[ 0], RIPEMD160C20, RIPEMD160S2A);
  85   RIPEMD160_STEP (RIPEMD160_H , c1, d1, e1, a1, b1, w[ 6], RIPEMD160C20, RIPEMD160S2B);
  86   RIPEMD160_STEP (RIPEMD160_H , b1, c1, d1, e1, a1, w[13], RIPEMD160C20, RIPEMD160S2C);
  87   RIPEMD160_STEP (RIPEMD160_H , a1, b1, c1, d1, e1, w[11], RIPEMD160C20, RIPEMD160S2D);
  88   RIPEMD160_STEP (RIPEMD160_H , e1, a1, b1, c1, d1, w[ 5], RIPEMD160C20, RIPEMD160S2E);
  89   RIPEMD160_STEP (RIPEMD160_H , d1, e1, a1, b1, c1, w[12], RIPEMD160C20, RIPEMD160S2F);
  90
  91   RIPEMD160_STEP (RIPEMD160_Io, c1, d1, e1, a1, b1, w[ 1], RIPEMD160C30, RIPEMD160S30);
  92   RIPEMD160_STEP (RIPEMD160_Io, b1, c1, d1, e1, a1, w[ 9], RIPEMD160C30, RIPEMD160S31);
  93   RIPEMD160_STEP (RIPEMD160_Io, a1, b1, c1, d1, e1, w[11], RIPEMD160C30, RIPEMD160S32);
  94   RIPEMD160_STEP (RIPEMD160_Io, e1, a1, b1, c1, d1, w[10], RIPEMD160C30, RIPEMD160S33);
  95   RIPEMD160_STEP (RIPEMD160_Io, d1, e1, a1, b1, c1, w[ 0], RIPEMD160C30, RIPEMD160S34);
  96   RIPEMD160_STEP (RIPEMD160_Io, c1, d1, e1, a1, b1, w[ 8], RIPEMD160C30, RIPEMD160S35);
  97   RIPEMD160_STEP (RIPEMD160_Io, b1, c1, d1, e1, a1, w[12], RIPEMD160C30, RIPEMD160S36);
  98   RIPEMD160_STEP (RIPEMD160_Io, a1, b1, c1, d1, e1, w[ 4], RIPEMD160C30, RIPEMD160S37);
  99   RIPEMD160_STEP (RIPEMD160_Io, e1, a1, b1, c1, d1, w[13], RIPEMD160C30, RIPEMD160S38);
 100   RIPEMD160_STEP (RIPEMD160_Io, d1, e1, a1, b1, c1, w[ 3], RIPEMD160C30, RIPEMD160S39);
 101   RIPEMD160_STEP (RIPEMD160_Io, c1, d1, e1, a1, b1, w[ 7], RIPEMD160C30, RIPEMD160S3A);
 102   RIPEMD160_STEP (RIPEMD160_Io, b1, c1, d1, e1, a1, w[15], RIPEMD160C30, RIPEMD160S3B);
 103   RIPEMD160_STEP (RIPEMD160_Io, a1, b1, c1, d1, e1, w[14], RIPEMD160C30, RIPEMD160S3C);
 104   RIPEMD160_STEP (RIPEMD160_Io, e1, a1, b1, c1, d1, w[ 5], RIPEMD160C30, RIPEMD160S3D);
 105   RIPEMD160_STEP (RIPEMD160_Io, d1, e1, a1, b1, c1, w[ 6], RIPEMD160C30, RIPEMD160S3E);
 106   RIPEMD160_STEP (RIPEMD160_Io, c1, d1, e1, a1, b1, w[ 2], RIPEMD160C30, RIPEMD160S3F);
 107
 108   RIPEMD160_STEP (RIPEMD160_J , b1, c1, d1, e1, a1, w[ 4], RIPEMD160C40, RIPEMD160S40);
 109   RIPEMD160_STEP (RIPEMD160_J , a1, b1, c1, d1, e1, w[ 0], RIPEMD160C40, RIPEMD160S41);
 110   RIPEMD160_STEP (RIPEMD160_J , e1, a1, b1, c1, d1, w[ 5], RIPEMD160C40, RIPEMD160S42);
 111   RIPEMD160_STEP (RIPEMD160_J , d1, e1, a1, b1, c1, w[ 9], RIPEMD160C40, RIPEMD160S43);
 112   RIPEMD160_STEP (RIPEMD160_J , c1, d1, e1, a1, b1, w[ 7], RIPEMD160C40, RIPEMD160S44);
 113   RIPEMD160_STEP (RIPEMD160_J , b1, c1, d1, e1, a1, w[12], RIPEMD160C40, RIPEMD160S45);
 114   RIPEMD160_STEP (RIPEMD160_J , a1, b1, c1, d1, e1, w[ 2], RIPEMD160C40, RIPEMD160S46);
 115   RIPEMD160_STEP (RIPEMD160_J , e1, a1, b1, c1, d1, w[10], RIPEMD160C40, RIPEMD160S47);
 116   RIPEMD160_STEP (RIPEMD160_J , d1, e1, a1, b1, c1, w[14], RIPEMD160C40, RIPEMD160S48);
 117   RIPEMD160_STEP (RIPEMD160_J , c1, d1, e1, a1, b1, w[ 1], RIPEMD160C40, RIPEMD160S49);
 118   RIPEMD160_STEP (RIPEMD160_J , b1, c1, d1, e1, a1, w[ 3], RIPEMD160C40, RIPEMD160S4A);
 119   RIPEMD160_STEP (RIPEMD160_J , a1, b1, c1, d1, e1, w[ 8], RIPEMD160C40, RIPEMD160S4B);
 120   RIPEMD160_STEP (RIPEMD160_J , e1, a1, b1, c1, d1, w[11], RIPEMD160C40, RIPEMD160S4C);
 121   RIPEMD160_STEP (RIPEMD160_J , d1, e1, a1, b1, c1, w[ 6], RIPEMD160C40, RIPEMD160S4D);
 122   RIPEMD160_STEP (RIPEMD160_J , c1, d1, e1, a1, b1, w[15], RIPEMD160C40, RIPEMD160S4E);
 123   RIPEMD160_STEP (RIPEMD160_J , b1, c1, d1, e1, a1, w[13], RIPEMD160C40, RIPEMD160S4F);
 124
 125   u32x a2 = dgst[0];
 126   u32x b2 = dgst[1];
 127   u32x c2 = dgst[2];
 128   u32x d2 = dgst[3];
 129   u32x e2 = dgst[4];
 130
 131   //RIPEMD160_STEP_WORKAROUND_BUG (RIPEMD160_J , a2, b2, c2, d2, e2, w[ 5], RIPEMD160C50, RIPEMD160S50);
 132   RIPEMD160_STEP (RIPEMD160_J , a2, b2, c2, d2, e2, w[ 5], RIPEMD160C50, RIPEMD160S50);
 133   RIPEMD160_STEP (RIPEMD160_J , e2, a2, b2, c2, d2, w[14], RIPEMD160C50, RIPEMD160S51);
 134   RIPEMD160_STEP (RIPEMD160_J , d2, e2, a2, b2, c2, w[ 7], RIPEMD160C50, RIPEMD160S52);
 135   RIPEMD160_STEP (RIPEMD160_J , c2, d2, e2, a2, b2, w[ 0], RIPEMD160C50, RIPEMD160S53);
 136   RIPEMD160_STEP (RIPEMD160_J , b2, c2, d2, e2, a2, w[ 9], RIPEMD160C50, RIPEMD160S54);
 137   RIPEMD160_STEP (RIPEMD160_J , a2, b2, c2, d2, e2, w[ 2], RIPEMD160C50, RIPEMD160S55);
 138   RIPEMD160_STEP (RIPEMD160_J , e2, a2, b2, c2, d2, w[11], RIPEMD160C50, RIPEMD160S56);
 139   RIPEMD160_STEP (RIPEMD160_J , d2, e2, a2, b2, c2, w[ 4], RIPEMD160C50, RIPEMD160S57);
 140   RIPEMD160_STEP (RIPEMD160_J , c2, d2, e2, a2, b2, w[13], RIPEMD160C50, RIPEMD160S58);
 141   RIPEMD160_STEP (RIPEMD160_J , b2, c2, d2, e2, a2, w[ 6], RIPEMD160C50, RIPEMD160S59);
 142   RIPEMD160_STEP (RIPEMD160_J , a2, b2, c2, d2, e2, w[15], RIPEMD160C50, RIPEMD160S5A);
 143   RIPEMD160_STEP (RIPEMD160_J , e2, a2, b2, c2, d2, w[ 8], RIPEMD160C50, RIPEMD160S5B);
 144   RIPEMD160_STEP (RIPEMD160_J , d2, e2, a2, b2, c2, w[ 1], RIPEMD160C50, RIPEMD160S5C);
 145   RIPEMD160_STEP (RIPEMD160_J , c2, d2, e2, a2, b2, w[10], RIPEMD160C50, RIPEMD160S5D);
 146   RIPEMD160_STEP (RIPEMD160_J , b2, c2, d2, e2, a2, w[ 3], RIPEMD160C50, RIPEMD160S5E);
 147   RIPEMD160_STEP (RIPEMD160_J , a2, b2, c2, d2, e2, w[12], RIPEMD160C50, RIPEMD160S5F);
 148
 149   RIPEMD160_STEP (RIPEMD160_Io, e2, a2, b2, c2, d2, w[ 6], RIPEMD160C60, RIPEMD160S60);
 150   RIPEMD160_STEP (RIPEMD160_Io, d2, e2, a2, b2, c2, w[11], RIPEMD160C60, RIPEMD160S61);
 151   RIPEMD160_STEP (RIPEMD160_Io, c2, d2, e2, a2, b2, w[ 3], RIPEMD160C60, RIPEMD160S62);
 152   RIPEMD160_STEP (RIPEMD160_Io, b2, c2, d2, e2, a2, w[ 7], RIPEMD160C60, RIPEMD160S63);
 153   RIPEMD160_STEP (RIPEMD160_Io, a2, b2, c2, d2, e2, w[ 0], RIPEMD160C60, RIPEMD160S64);
 154   RIPEMD160_STEP (RIPEMD160_Io, e2, a2, b2, c2, d2, w[13], RIPEMD160C60, RIPEMD160S65);
 155   RIPEMD160_STEP (RIPEMD160_Io, d2, e2, a2, b2, c2, w[ 5], RIPEMD160C60, RIPEMD160S66);
 156   RIPEMD160_STEP (RIPEMD160_Io, c2, d2, e2, a2, b2, w[10], RIPEMD160C60, RIPEMD160S67);
 157   RIPEMD160_STEP (RIPEMD160_Io, b2, c2, d2, e2, a2, w[14], RIPEMD160C60, RIPEMD160S68);
 158   RIPEMD160_STEP (RIPEMD160_Io, a2, b2, c2, d2, e2, w[15], RIPEMD160C60, RIPEMD160S69);
 159   RIPEMD160_STEP (RIPEMD160_Io, e2, a2, b2, c2, d2, w[ 8], RIPEMD160C60, RIPEMD160S6A);
 160   RIPEMD160_STEP (RIPEMD160_Io, d2, e2, a2, b2, c2, w[12], RIPEMD160C60, RIPEMD160S6B);
 161   RIPEMD160_STEP (RIPEMD160_Io, c2, d2, e2, a2, b2, w[ 4], RIPEMD160C60, RIPEMD160S6C);
 162   RIPEMD160_STEP (RIPEMD160_Io, b2, c2, d2, e2, a2, w[ 9], RIPEMD160C60, RIPEMD160S6D);
 163   RIPEMD160_STEP (RIPEMD160_Io, a2, b2, c2, d2, e2, w[ 1], RIPEMD160C60, RIPEMD160S6E);
 164   RIPEMD160_STEP (RIPEMD160_Io, e2, a2, b2, c2, d2, w[ 2], RIPEMD160C60, RIPEMD160S6F);
 165
 166   RIPEMD160_STEP (RIPEMD160_H , d2, e2, a2, b2, c2, w[15], RIPEMD160C70, RIPEMD160S70);
 167   RIPEMD160_STEP (RIPEMD160_H , c2, d2, e2, a2, b2, w[ 5], RIPEMD160C70, RIPEMD160S71);
 168   RIPEMD160_STEP (RIPEMD160_H , b2, c2, d2, e2, a2, w[ 1], RIPEMD160C70, RIPEMD160S72);
 169   RIPEMD160_STEP (RIPEMD160_H , a2, b2, c2, d2, e2, w[ 3], RIPEMD160C70, RIPEMD160S73);
 170   RIPEMD160_STEP (RIPEMD160_H , e2, a2, b2, c2, d2, w[ 7], RIPEMD160C70, RIPEMD160S74);
 171   RIPEMD160_STEP (RIPEMD160_H , d2, e2, a2, b2, c2, w[14], RIPEMD160C70, RIPEMD160S75);
 172   RIPEMD160_STEP (RIPEMD160_H , c2, d2, e2, a2, b2, w[ 6], RIPEMD160C70, RIPEMD160S76);
 173   RIPEMD160_STEP (RIPEMD160_H , b2, c2, d2, e2, a2, w[ 9], RIPEMD160C70, RIPEMD160S77);
 174   RIPEMD160_STEP (RIPEMD160_H , a2, b2, c2, d2, e2, w[11], RIPEMD160C70, RIPEMD160S78);
 175   RIPEMD160_STEP (RIPEMD160_H , e2, a2, b2, c2, d2, w[ 8], RIPEMD160C70, RIPEMD160S79);
 176   RIPEMD160_STEP (RIPEMD160_H , d2, e2, a2, b2, c2, w[12], RIPEMD160C70, RIPEMD160S7A);
 177   RIPEMD160_STEP (RIPEMD160_H , c2, d2, e2, a2, b2, w[ 2], RIPEMD160C70, RIPEMD160S7B);
 178   RIPEMD160_STEP (RIPEMD160_H , b2, c2, d2, e2, a2, w[10], RIPEMD160C70, RIPEMD160S7C);
 179   RIPEMD160_STEP (RIPEMD160_H , a2, b2, c2, d2, e2, w[ 0], RIPEMD160C70, RIPEMD160S7D);
 180   RIPEMD160_STEP (RIPEMD160_H , e2, a2, b2, c2, d2, w[ 4], RIPEMD160C70, RIPEMD160S7E);
 181   RIPEMD160_STEP (RIPEMD160_H , d2, e2, a2, b2, c2, w[13], RIPEMD160C70, RIPEMD160S7F);
 182
 183   RIPEMD160_STEP (RIPEMD160_Go, c2, d2, e2, a2, b2, w[ 8], RIPEMD160C80, RIPEMD160S80);
 184   RIPEMD160_STEP (RIPEMD160_Go, b2, c2, d2, e2, a2, w[ 6], RIPEMD160C80, RIPEMD160S81);
 185   RIPEMD160_STEP (RIPEMD160_Go, a2, b2, c2, d2, e2, w[ 4], RIPEMD160C80, RIPEMD160S82);
 186   RIPEMD160_STEP (RIPEMD160_Go, e2, a2, b2, c2, d2, w[ 1], RIPEMD160C80, RIPEMD160S83);
 187   RIPEMD160_STEP (RIPEMD160_Go, d2, e2, a2, b2, c2, w[ 3], RIPEMD160C80, RIPEMD160S84);
 188   RIPEMD160_STEP (RIPEMD160_Go, c2, d2, e2, a2, b2, w[11], RIPEMD160C80, RIPEMD160S85);
 189   RIPEMD160_STEP (RIPEMD160_Go, b2, c2, d2, e2, a2, w[15], RIPEMD160C80, RIPEMD160S86);
 190   RIPEMD160_STEP (RIPEMD160_Go, a2, b2, c2, d2, e2, w[ 0], RIPEMD160C80, RIPEMD160S87);
 191   RIPEMD160_STEP (RIPEMD160_Go, e2, a2, b2, c2, d2, w[ 5], RIPEMD160C80, RIPEMD160S88);
 192   RIPEMD160_STEP (RIPEMD160_Go, d2, e2, a2, b2, c2, w[12], RIPEMD160C80, RIPEMD160S89);
 193   RIPEMD160_STEP (RIPEMD160_Go, c2, d2, e2, a2, b2, w[ 2], RIPEMD160C80, RIPEMD160S8A);
 194   RIPEMD160_STEP (RIPEMD160_Go, b2, c2, d2, e2, a2, w[13], RIPEMD160C80, RIPEMD160S8B);
 195   RIPEMD160_STEP (RIPEMD160_Go, a2, b2, c2, d2, e2, w[ 9], RIPEMD160C80, RIPEMD160S8C);
 196   RIPEMD160_STEP (RIPEMD160_Go, e2, a2, b2, c2, d2, w[ 7], RIPEMD160C80, RIPEMD160S8D);
 197   RIPEMD160_STEP (RIPEMD160_Go, d2, e2, a2, b2, c2, w[10], RIPEMD160C80, RIPEMD160S8E);
 198   RIPEMD160_STEP (RIPEMD160_Go, c2, d2, e2, a2, b2, w[14], RIPEMD160C80, RIPEMD160S8F);
 199
 200   RIPEMD160_STEP (RIPEMD160_F , b2, c2, d2, e2, a2, w[12], RIPEMD160C90, RIPEMD160S90);
 201   RIPEMD160_STEP (RIPEMD160_F , a2, b2, c2, d2, e2, w[15], RIPEMD160C90, RIPEMD160S91);
 202   RIPEMD160_STEP (RIPEMD160_F , e2, a2, b2, c2, d2, w[10], RIPEMD160C90, RIPEMD160S92);
 203   RIPEMD160_STEP (RIPEMD160_F , d2, e2, a2, b2, c2, w[ 4], RIPEMD160C90, RIPEMD160S93);
 204   RIPEMD160_STEP (RIPEMD160_F , c2, d2, e2, a2, b2, w[ 1], RIPEMD160C90, RIPEMD160S94);
 205   RIPEMD160_STEP (RIPEMD160_F , b2, c2, d2, e2, a2, w[ 5], RIPEMD160C90, RIPEMD160S95);
 206   RIPEMD160_STEP (RIPEMD160_F , a2, b2, c2, d2, e2, w[ 8], RIPEMD160C90, RIPEMD160S96);
 207   RIPEMD160_STEP (RIPEMD160_F , e2, a2, b2, c2, d2, w[ 7], RIPEMD160C90, RIPEMD160S97);
 208   RIPEMD160_STEP (RIPEMD160_F , d2, e2, a2, b2, c2, w[ 6], RIPEMD160C90, RIPEMD160S98);
 209   RIPEMD160_STEP (RIPEMD160_F , c2, d2, e2, a2, b2, w[ 2], RIPEMD160C90, RIPEMD160S99);
 210   RIPEMD160_STEP (RIPEMD160_F , b2, c2, d2, e2, a2, w[13], RIPEMD160C90, RIPEMD160S9A);
 211   RIPEMD160_STEP (RIPEMD160_F , a2, b2, c2, d2, e2, w[14], RIPEMD160C90, RIPEMD160S9B);
 212   RIPEMD160_STEP (RIPEMD160_F , e2, a2, b2, c2, d2, w[ 0], RIPEMD160C90, RIPEMD160S9C);
 213   RIPEMD160_STEP (RIPEMD160_F , d2, e2, a2, b2, c2, w[ 3], RIPEMD160C90, RIPEMD160S9D);
 214   RIPEMD160_STEP (RIPEMD160_F , c2, d2, e2, a2, b2, w[ 9], RIPEMD160C90, RIPEMD160S9E);
 215   RIPEMD160_STEP (RIPEMD160_F , b2, c2, d2, e2, a2, w[11], RIPEMD160C90, RIPEMD160S9F);
 216
 217   const u32x a = dgst[1] + c1 + d2;
 218   const u32x b = dgst[2] + d1 + e2;
 219   const u32x c = dgst[3] + e1 + a2;
 220   const u32x d = dgst[4] + a1 + b2;
 221   const u32x e = dgst[0] + b1 + c2;
 222
 223   dgst[0] = a;
 224   dgst[1] = b;
 225   dgst[2] = c;
 226   dgst[3] = d;
 227   dgst[4] = e;
 228 }
 229
 230 __device__ static void hmac_run2 (const u32x w1[16], const u32x w2[16], const u32x ipad[5], const u32x opad[5], u32x dgst[5])
 231 {
 232   dgst[0] = ipad[0];
 233   dgst[1] = ipad[1];
 234   dgst[2] = ipad[2];
 235   dgst[3] = ipad[3];
 236   dgst[4] = ipad[4];
 237
 238   ripemd160_transform (w1, dgst);
 239   ripemd160_transform (w2, dgst);
 240
 241   u32x w[16];
 242
 243   w[ 0] = dgst[0];
 244   w[ 1] = dgst[1];
 245   w[ 2] = dgst[2];
 246   w[ 3] = dgst[3];
 247   w[ 4] = dgst[4];
 248   w[ 5] = 0x80;
 249   w[ 6] = 0;
 250   w[ 7] = 0;
 251   w[ 8] = 0;
 252   w[ 9] = 0;
 253   w[10] = 0;
 254   w[11] = 0;
 255   w[12] = 0;
 256   w[13] = 0;
 257   w[14] = (64 + 20) * 8;
 258   w[15] = 0;
 259
 260   dgst[0] = opad[0];
 261   dgst[1] = opad[1];
 262   dgst[2] = opad[2];
 263   dgst[3] = opad[3];
 264   dgst[4] = opad[4];
 265
 266   ripemd160_transform (w, dgst);
 267 }
 268
 269 __device__ static void hmac_run (u32x w[16], const u32x ipad[5], const u32x opad[5], u32x dgst[5])
 270 {
 271   dgst[0] = ipad[0];
 272   dgst[1] = ipad[1];
 273   dgst[2] = ipad[2];
 274   dgst[3] = ipad[3];
 275   dgst[4] = ipad[4];
 276
 277   ripemd160_transform (w, dgst);
 278
 279   w[ 0] = dgst[0];
 280   w[ 1] = dgst[1];
 281   w[ 2] = dgst[2];
 282   w[ 3] = dgst[3];
 283   w[ 4] = dgst[4];
 284   w[ 5] = 0x80;
 285   w[ 6] = 0;
 286   w[ 7] = 0;
 287   w[ 8] = 0;
 288   w[ 9] = 0;
 289   w[10] = 0;
 290   w[11] = 0;
 291   w[12] = 0;
 292   w[13] = 0;
 293   w[14] = (64 + 20) * 8;
 294   w[15] = 0;
 295
 296   dgst[0] = opad[0];
 297   dgst[1] = opad[1];
 298   dgst[2] = opad[2];
 299   dgst[3] = opad[3];
 300   dgst[4] = opad[4];
 301
 302   ripemd160_transform (w, dgst);
 303 }
 304
 305 __device__ static void hmac_init (u32x w[16], u32x ipad[5], u32x opad[5])
 306 {
 307   w[ 0] ^= 0x36363636;
 308   w[ 1] ^= 0x36363636;
 309   w[ 2] ^= 0x36363636;
 310   w[ 3] ^= 0x36363636;
 311   w[ 4] ^= 0x36363636;
 312   w[ 5] ^= 0x36363636;
 313   w[ 6] ^= 0x36363636;
 314   w[ 7] ^= 0x36363636;
 315   w[ 8] ^= 0x36363636;
 316   w[ 9] ^= 0x36363636;
 317   w[10] ^= 0x36363636;
 318   w[11] ^= 0x36363636;
 319   w[12] ^= 0x36363636;
 320   w[13] ^= 0x36363636;
 321   w[14] ^= 0x36363636;
 322   w[15] ^= 0x36363636;
 323
 324   ipad[0] = RIPEMD160M_A;
 325   ipad[1] = RIPEMD160M_B;
 326   ipad[2] = RIPEMD160M_C;
 327   ipad[3] = RIPEMD160M_D;
 328   ipad[4] = RIPEMD160M_E;
 329
 330   ripemd160_transform (w, ipad);
 331
 332   w[ 0] ^= 0x6a6a6a6a;
 333   w[ 1] ^= 0x6a6a6a6a;
 334   w[ 2] ^= 0x6a6a6a6a;
 335   w[ 3] ^= 0x6a6a6a6a;
 336   w[ 4] ^= 0x6a6a6a6a;
 337   w[ 5] ^= 0x6a6a6a6a;
 338   w[ 6] ^= 0x6a6a6a6a;
 339   w[ 7] ^= 0x6a6a6a6a;
 340   w[ 8] ^= 0x6a6a6a6a;
 341   w[ 9] ^= 0x6a6a6a6a;
 342   w[10] ^= 0x6a6a6a6a;
 343   w[11] ^= 0x6a6a6a6a;
 344   w[12] ^= 0x6a6a6a6a;
 345   w[13] ^= 0x6a6a6a6a;
 346   w[14] ^= 0x6a6a6a6a;
 347   w[15] ^= 0x6a6a6a6a;
 348
 349   opad[0] = RIPEMD160M_A;
 350   opad[1] = RIPEMD160M_B;
 351   opad[2] = RIPEMD160M_C;
 352   opad[3] = RIPEMD160M_D;
 353   opad[4] = RIPEMD160M_E;
 354
 355   ripemd160_transform (w, opad);
 356 }
 357
 358 __device__ static u32 u8add (const u32 a, const u32 b)
 359 {
 360   const u32 a1 = (a >>  0) & 0xff;
 361   const u32 a2 = (a >>  8) & 0xff;
 362   const u32 a3 = (a >> 16) & 0xff;
 363   const u32 a4 = (a >> 24) & 0xff;
 364
 365   const u32 b1 = (b >>  0) & 0xff;
 366   const u32 b2 = (b >>  8) & 0xff;
 367   const u32 b3 = (b >> 16) & 0xff;
 368   const u32 b4 = (b >> 24) & 0xff;
 369
 370   const u32 r1 = (a1 + b1) & 0xff;
 371   const u32 r2 = (a2 + b2) & 0xff;
 372   const u32 r3 = (a3 + b3) & 0xff;
 373   const u32 r4 = (a4 + b4) & 0xff;
 374
 375   const u32 r = r1 <<  0
 376                | r2 <<  8
 377                | r3 << 16
 378                | r4 << 24;
 379
 380   return r;
 381 }
 382
 383 extern "C" __global__ void __launch_bounds__ (256, 1) m06213_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc_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)
 384 {
 385   /**
 386    * base
 387    */
 388
 389   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 390
 391   if (gid >= gid_max) return;
 392
 393   u32x w0[4];
 394
 395   w0[0] = pws[gid].i[ 0];
 396   w0[1] = pws[gid].i[ 1];
 397   w0[2] = pws[gid].i[ 2];
 398   w0[3] = pws[gid].i[ 3];
 399
 400   u32x w1[4];
 401
 402   w1[0] = pws[gid].i[ 4];
 403   w1[1] = pws[gid].i[ 5];
 404   w1[2] = pws[gid].i[ 6];
 405   w1[3] = pws[gid].i[ 7];
 406
 407   u32x w2[4];
 408
 409   w2[0] = pws[gid].i[ 8];
 410   w2[1] = pws[gid].i[ 9];
 411   w2[2] = pws[gid].i[10];
 412   w2[3] = pws[gid].i[11];
 413
 414   u32x w3[4];
 415
 416   w3[0] = pws[gid].i[12];
 417   w3[1] = pws[gid].i[13];
 418   w3[2] = pws[gid].i[14];
 419   w3[3] = pws[gid].i[15];
 420
 421   /**
 422    * keyfile
 423    */
 424
 425   w0[0] = u8add (w0[0], esalt_bufs[salt_pos].keyfile_buf[ 0]);
 426   w0[1] = u8add (w0[1], esalt_bufs[salt_pos].keyfile_buf[ 1]);
 427   w0[2] = u8add (w0[2], esalt_bufs[salt_pos].keyfile_buf[ 2]);
 428   w0[3] = u8add (w0[3], esalt_bufs[salt_pos].keyfile_buf[ 3]);
 429   w1[0] = u8add (w1[0], esalt_bufs[salt_pos].keyfile_buf[ 4]);
 430   w1[1] = u8add (w1[1], esalt_bufs[salt_pos].keyfile_buf[ 5]);
 431   w1[2] = u8add (w1[2], esalt_bufs[salt_pos].keyfile_buf[ 6]);
 432   w1[3] = u8add (w1[3], esalt_bufs[salt_pos].keyfile_buf[ 7]);
 433   w2[0] = u8add (w2[0], esalt_bufs[salt_pos].keyfile_buf[ 8]);
 434   w2[1] = u8add (w2[1], esalt_bufs[salt_pos].keyfile_buf[ 9]);
 435   w2[2] = u8add (w2[2], esalt_bufs[salt_pos].keyfile_buf[10]);
 436   w2[3] = u8add (w2[3], esalt_bufs[salt_pos].keyfile_buf[11]);
 437   w3[0] = u8add (w3[0], esalt_bufs[salt_pos].keyfile_buf[12]);
 438   w3[1] = u8add (w3[1], esalt_bufs[salt_pos].keyfile_buf[13]);
 439   w3[2] = u8add (w3[2], esalt_bufs[salt_pos].keyfile_buf[14]);
 440   w3[3] = u8add (w3[3], esalt_bufs[salt_pos].keyfile_buf[15]);
 441
 442   /**
 443    * salt
 444    */
 445
 446   u32x salt_buf1[16];
 447
 448   salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[ 0];
 449   salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[ 1];
 450   salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[ 2];
 451   salt_buf1[ 3] = esalt_bufs[salt_pos].salt_buf[ 3];
 452   salt_buf1[ 4] = esalt_bufs[salt_pos].salt_buf[ 4];
 453   salt_buf1[ 5] = esalt_bufs[salt_pos].salt_buf[ 5];
 454   salt_buf1[ 6] = esalt_bufs[salt_pos].salt_buf[ 6];
 455   salt_buf1[ 7] = esalt_bufs[salt_pos].salt_buf[ 7];
 456   salt_buf1[ 8] = esalt_bufs[salt_pos].salt_buf[ 8];
 457   salt_buf1[ 9] = esalt_bufs[salt_pos].salt_buf[ 9];
 458   salt_buf1[10] = esalt_bufs[salt_pos].salt_buf[10];
 459   salt_buf1[11] = esalt_bufs[salt_pos].salt_buf[11];
 460   salt_buf1[12] = esalt_bufs[salt_pos].salt_buf[12];
 461   salt_buf1[13] = esalt_bufs[salt_pos].salt_buf[13];
 462   salt_buf1[14] = esalt_bufs[salt_pos].salt_buf[14];
 463   salt_buf1[15] = esalt_bufs[salt_pos].salt_buf[15];
 464
 465   u32x salt_buf2[16];
 466
 467   salt_buf2[ 0] = 0;
 468   salt_buf2[ 1] = 0x80;
 469   salt_buf2[ 2] = 0;
 470   salt_buf2[ 3] = 0;
 471   salt_buf2[ 4] = 0;
 472   salt_buf2[ 5] = 0;
 473   salt_buf2[ 6] = 0;
 474   salt_buf2[ 7] = 0;
 475   salt_buf2[ 8] = 0;
 476   salt_buf2[ 9] = 0;
 477   salt_buf2[10] = 0;
 478   salt_buf2[11] = 0;
 479   salt_buf2[12] = 0;
 480   salt_buf2[13] = 0;
 481   salt_buf2[14] = (64 + 64 + 4) * 8;
 482   salt_buf2[15] = 0;
 483
 484   const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen;
 485
 486   u32x w[16];
 487
 488   w[ 0] = w0[0];
 489   w[ 1] = w0[1];
 490   w[ 2] = w0[2];
 491   w[ 3] = w0[3];
 492   w[ 4] = w1[0];
 493   w[ 5] = w1[1];
 494   w[ 6] = w1[2];
 495   w[ 7] = w1[3];
 496   w[ 8] = w2[0];
 497   w[ 9] = w2[1];
 498   w[10] = w2[2];
 499   w[11] = w2[3];
 500   w[12] = w3[0];
 501   w[13] = w3[1];
 502   w[14] = w3[2];
 503   w[15] = w3[3];
 504
 505   u32x ipad[5];
 506   u32x opad[5];
 507
 508   hmac_init (w, ipad, opad);
 509
 510   tmps[gid].ipad[0] = ipad[0];
 511   tmps[gid].ipad[1] = ipad[1];
 512   tmps[gid].ipad[2] = ipad[2];
 513   tmps[gid].ipad[3] = ipad[3];
 514   tmps[gid].ipad[4] = ipad[4];
 515
 516   tmps[gid].opad[0] = opad[0];
 517   tmps[gid].opad[1] = opad[1];
 518   tmps[gid].opad[2] = opad[2];
 519   tmps[gid].opad[3] = opad[3];
 520   tmps[gid].opad[4] = opad[4];
 521
 522   for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 5, j += 1)
 523   {
 524     salt_buf2[0] = swap_workaround (j);
 525
 526     u32x dgst[5];
 527
 528     hmac_run2 (salt_buf1, salt_buf2, ipad, opad, dgst);
 529
 530     tmps[gid].dgst[i + 0] = dgst[0];
 531     tmps[gid].dgst[i + 1] = dgst[1];
 532     tmps[gid].dgst[i + 2] = dgst[2];
 533     tmps[gid].dgst[i + 3] = dgst[3];
 534     tmps[gid].dgst[i + 4] = dgst[4];
 535
 536     tmps[gid].out[i + 0] = dgst[0];
 537     tmps[gid].out[i + 1] = dgst[1];
 538     tmps[gid].out[i + 2] = dgst[2];
 539     tmps[gid].out[i + 3] = dgst[3];
 540     tmps[gid].out[i + 4] = dgst[4];
 541   }
 542 }
 543
 544 extern "C" __global__ void __launch_bounds__ (256, 1) m06213_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc_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)
 545 {
 546   const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen;
 547
 548   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 549
 550   if (gid >= gid_max) return;
 551
 552   u32x ipad[5];
 553   u32x opad[5];
 554
 555   ipad[0] = tmps[gid].ipad[0];
 556   ipad[1] = tmps[gid].ipad[1];
 557   ipad[2] = tmps[gid].ipad[2];
 558   ipad[3] = tmps[gid].ipad[3];
 559   ipad[4] = tmps[gid].ipad[4];
 560
 561   opad[0] = tmps[gid].opad[0];
 562   opad[1] = tmps[gid].opad[1];
 563   opad[2] = tmps[gid].opad[2];
 564   opad[3] = tmps[gid].opad[3];
 565   opad[4] = tmps[gid].opad[4];
 566
 567   for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 5)
 568   {
 569     u32x dgst[5];
 570     u32x out[5];
 571
 572     dgst[0] = tmps[gid].dgst[i + 0];
 573     dgst[1] = tmps[gid].dgst[i + 1];
 574     dgst[2] = tmps[gid].dgst[i + 2];
 575     dgst[3] = tmps[gid].dgst[i + 3];
 576     dgst[4] = tmps[gid].dgst[i + 4];
 577
 578     out[0] = tmps[gid].out[i + 0];
 579     out[1] = tmps[gid].out[i + 1];
 580     out[2] = tmps[gid].out[i + 2];
 581     out[3] = tmps[gid].out[i + 3];
 582     out[4] = tmps[gid].out[i + 4];
 583
 584     for (u32 j = 0; j < loop_cnt; j++)
 585     {
 586       u32x w[16];
 587
 588       w[ 0] = dgst[0];
 589       w[ 1] = dgst[1];
 590       w[ 2] = dgst[2];
 591       w[ 3] = dgst[3];
 592       w[ 4] = dgst[4];
 593       w[ 5] = 0x80;
 594       w[ 6] = 0;
 595       w[ 7] = 0;
 596       w[ 8] = 0;
 597       w[ 9] = 0;
 598       w[10] = 0;
 599       w[11] = 0;
 600       w[12] = 0;
 601       w[13] = 0;
 602       w[14] = (64 + 20) * 8;
 603       w[15] = 0;
 604
 605       hmac_run (w, ipad, opad, dgst);
 606
 607       out[0] ^= dgst[0];
 608       out[1] ^= dgst[1];
 609       out[2] ^= dgst[2];
 610       out[3] ^= dgst[3];
 611       out[4] ^= dgst[4];
 612     }
 613
 614     tmps[gid].dgst[i + 0] = dgst[0];
 615     tmps[gid].dgst[i + 1] = dgst[1];
 616     tmps[gid].dgst[i + 2] = dgst[2];
 617     tmps[gid].dgst[i + 3] = dgst[3];
 618     tmps[gid].dgst[i + 4] = dgst[4];
 619
 620     tmps[gid].out[i + 0] = out[0];
 621     tmps[gid].out[i + 1] = out[1];
 622     tmps[gid].out[i + 2] = out[2];
 623     tmps[gid].out[i + 3] = out[3];
 624     tmps[gid].out[i + 4] = out[4];
 625   }
 626 }
 627
 628 extern "C" __global__ void __launch_bounds__ (256, 1) m06213_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc_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)
 629 {
 630   /**
 631    * base
 632    */
 633
 634   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 635   const u32 lid = threadIdx.x;
 636
 637   if (gid >= gid_max) return;
 638
 639   u32 ukey1[8];
 640
 641   ukey1[0] = tmps[gid].out[ 0];
 642   ukey1[1] = tmps[gid].out[ 1];
 643   ukey1[2] = tmps[gid].out[ 2];
 644   ukey1[3] = tmps[gid].out[ 3];
 645   ukey1[4] = tmps[gid].out[ 4];
 646   ukey1[5] = tmps[gid].out[ 5];
 647   ukey1[6] = tmps[gid].out[ 6];
 648   ukey1[7] = tmps[gid].out[ 7];
 649
 650   u32 ukey2[8];
 651
 652   ukey2[0] = tmps[gid].out[ 8];
 653   ukey2[1] = tmps[gid].out[ 9];
 654   ukey2[2] = tmps[gid].out[10];
 655   ukey2[3] = tmps[gid].out[11];
 656   ukey2[4] = tmps[gid].out[12];
 657   ukey2[5] = tmps[gid].out[13];
 658   ukey2[6] = tmps[gid].out[14];
 659   ukey2[7] = tmps[gid].out[15];
 660
 661   u32 data[4];
 662
 663   data[0] = esalt_bufs[0].data_buf[0];
 664   data[1] = esalt_bufs[0].data_buf[1];
 665   data[2] = esalt_bufs[0].data_buf[2];
 666   data[3] = esalt_bufs[0].data_buf[3];
 667
 668   u32 tmp[4];
 669
 670   {
 671     tmp[0] = data[0];
 672     tmp[1] = data[1];
 673     tmp[2] = data[2];
 674     tmp[3] = data[3];
 675
 676     aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
 677
 678     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 679     {
 680       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 681
 682       d_return_buf[lid] = 1;
 683     }
 684   }
 685
 686   {
 687     tmp[0] = data[0];
 688     tmp[1] = data[1];
 689     tmp[2] = data[2];
 690     tmp[3] = data[3];
 691
 692     serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
 693
 694     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 695     {
 696       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 697
 698       d_return_buf[lid] = 1;
 699     }
 700   }
 701
 702   {
 703     tmp[0] = data[0];
 704     tmp[1] = data[1];
 705     tmp[2] = data[2];
 706     tmp[3] = data[3];
 707
 708     twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
 709
 710     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 711     {
 712       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 713
 714       d_return_buf[lid] = 1;
 715     }
 716   }
 717
 718   u32 ukey3[8];
 719
 720   ukey3[0] = tmps[gid].out[16];
 721   ukey3[1] = tmps[gid].out[17];
 722   ukey3[2] = tmps[gid].out[18];
 723   ukey3[3] = tmps[gid].out[19];
 724   ukey3[4] = tmps[gid].out[20];
 725   ukey3[5] = tmps[gid].out[21];
 726   ukey3[6] = tmps[gid].out[22];
 727   ukey3[7] = tmps[gid].out[23];
 728
 729   u32 ukey4[8];
 730
 731   ukey4[0] = tmps[gid].out[24];
 732   ukey4[1] = tmps[gid].out[25];
 733   ukey4[2] = tmps[gid].out[26];
 734   ukey4[3] = tmps[gid].out[27];
 735   ukey4[4] = tmps[gid].out[28];
 736   ukey4[5] = tmps[gid].out[29];
 737   ukey4[6] = tmps[gid].out[30];
 738   ukey4[7] = tmps[gid].out[31];
 739
 740   {
 741     tmp[0] = data[0];
 742     tmp[1] = data[1];
 743     tmp[2] = data[2];
 744     tmp[3] = data[3];
 745
 746     aes256_decrypt_xts     (ukey2, ukey4, tmp, tmp);
 747     twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
 748
 749     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 750     {
 751       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 752
 753       d_return_buf[lid] = 1;
 754     }
 755   }
 756
 757   {
 758     tmp[0] = data[0];
 759     tmp[1] = data[1];
 760     tmp[2] = data[2];
 761     tmp[3] = data[3];
 762
 763     serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
 764     aes256_decrypt_xts     (ukey1, ukey3, tmp, tmp);
 765
 766     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 767     {
 768       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 769
 770       d_return_buf[lid] = 1;
 771     }
 772   }
 773
 774   {
 775     tmp[0] = data[0];
 776     tmp[1] = data[1];
 777     tmp[2] = data[2];
 778     tmp[3] = data[3];
 779
 780     twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
 781     serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
 782
 783     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 784     {
 785       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 786
 787       d_return_buf[lid] = 1;
 788     }
 789   }
 790
 791   u32 ukey5[8];
 792
 793   ukey5[0] = tmps[gid].out[32];
 794   ukey5[1] = tmps[gid].out[33];
 795   ukey5[2] = tmps[gid].out[34];
 796   ukey5[3] = tmps[gid].out[35];
 797   ukey5[4] = tmps[gid].out[36];
 798   ukey5[5] = tmps[gid].out[37];
 799   ukey5[6] = tmps[gid].out[38];
 800   ukey5[7] = tmps[gid].out[39];
 801
 802   u32 ukey6[8];
 803
 804   ukey6[0] = tmps[gid].out[40];
 805   ukey6[1] = tmps[gid].out[41];
 806   ukey6[2] = tmps[gid].out[42];
 807   ukey6[3] = tmps[gid].out[43];
 808   ukey6[4] = tmps[gid].out[44];
 809   ukey6[5] = tmps[gid].out[45];
 810   ukey6[6] = tmps[gid].out[46];
 811   ukey6[7] = tmps[gid].out[47];
 812
 813   {
 814     tmp[0] = data[0];
 815     tmp[1] = data[1];
 816     tmp[2] = data[2];
 817     tmp[3] = data[3];
 818
 819     aes256_decrypt_xts     (ukey3, ukey6, tmp, tmp);
 820     twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
 821     serpent256_decrypt_xts (ukey1, ukey4, tmp, tmp);
 822
 823     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 824     {
 825       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 826
 827       d_return_buf[lid] = 1;
 828     }
 829   }
 830
 831   {
 832     tmp[0] = data[0];
 833     tmp[1] = data[1];
 834     tmp[2] = data[2];
 835     tmp[3] = data[3];
 836
 837     serpent256_decrypt_xts (ukey3, ukey6, tmp, tmp);
 838     twofish256_decrypt_xts (ukey2, ukey5, tmp, tmp);
 839     aes256_decrypt_xts     (ukey1, ukey4, tmp, tmp);
 840
 841     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
 842     {
 843       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
 844
 845       d_return_buf[lid] = 1;
 846     }
 847   }
 848 }